Skip to content

Commit

Permalink
[pre-commit.ci] auto code formatting
Browse files Browse the repository at this point in the history
  • Loading branch information
@pre-commit-ci
pre-commit-ci[bot] committed Jul 2, 2024
1 parent 88f37e8 commit 9df96fd
Show file tree
Hide file tree
Showing 172 changed files with 1,070 additions and 1,247 deletions.
29 changes: 15 additions & 14 deletions cub/cub/agent/agent_merge_sort.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -346,20 +346,21 @@ struct AgentPartition
OffsetT partition_at = (cub::min)(keys2_end - keys1_beg, items_per_tile * local_tile_idx);

OffsetT partition_diag =
ping ? MergePath<KeyT>(
keys_ping + keys1_beg,
keys_ping + keys2_beg,
keys1_end - keys1_beg,
keys2_end - keys2_beg,
partition_at,
compare_op)
: MergePath<KeyT>(
keys_pong + keys1_beg,
keys_pong + keys2_beg,
keys1_end - keys1_beg,
keys2_end - keys2_beg,
partition_at,
compare_op);
ping
? MergePath<KeyT>(
keys_ping + keys1_beg,
keys_ping + keys2_beg,
keys1_end - keys1_beg,
keys2_end - keys2_beg,
partition_at,
compare_op)
: MergePath<KeyT>(
keys_pong + keys1_beg,
keys_pong + keys2_beg,
keys1_end - keys1_beg,
keys2_end - keys2_beg,
partition_at,
compare_op);

merge_partitions[partition_idx] = keys1_beg + partition_diag;
}
Expand Down
132 changes: 66 additions & 66 deletions cub/cub/block/block_radix_sort.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -674,9 +674,9 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -733,9 +733,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
Sort(keys, decomposer, 0, detail::radix::traits_t<KeyT>::default_end_bit(decomposer));
}
Expand Down Expand Up @@ -874,13 +874,13 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
{
SortBlocked(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>(), decomposer);
}
Expand Down Expand Up @@ -943,9 +943,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
Sort(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
Sort(keys, values, decomposer, 0, detail::radix::traits_t<KeyT>::default_end_bit(decomposer));
}
Expand Down Expand Up @@ -1064,9 +1064,9 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -1123,9 +1123,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -1271,13 +1271,13 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
{
SortBlocked(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>(), decomposer);
}
Expand Down Expand Up @@ -1340,9 +1340,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescending(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
SortBlocked(
keys,
Expand Down Expand Up @@ -1473,9 +1473,9 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -1533,9 +1533,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -1677,13 +1677,13 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
{
SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<false>(), Int2Type<KEYS_ONLY>(), decomposer);
}
Expand Down Expand Up @@ -1742,9 +1742,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
SortBlockedToStriped(
keys,
Expand Down Expand Up @@ -1871,9 +1871,9 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer, int begin_bit, int end_bit)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -1931,9 +1931,9 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
NullType values[ITEMS_PER_THREAD];

Expand Down Expand Up @@ -2075,14 +2075,14 @@ public:
//! comparison (e.g., `(sizeof(float) + sizeof(long long int)) * 8`)
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(
KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(
KeyT (&keys)[ITEMS_PER_THREAD],
ValueT (&values)[ITEMS_PER_THREAD],
DecomposerT decomposer,
int begin_bit,
int end_bit)
{
SortBlockedToStriped(keys, values, begin_bit, end_bit, Int2Type<true>(), Int2Type<KEYS_ONLY>(), decomposer);
}
Expand Down Expand Up @@ -2141,10 +2141,10 @@ public:
//! modify members of the key.
template <class DecomposerT>
_CCCL_DEVICE _CCCL_FORCEINLINE //
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(
KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
typename ::cuda::std::enable_if< //
!::cuda::std::is_convertible<DecomposerT, int>::value>::type
SortDescendingBlockedToStriped(
KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&values)[ITEMS_PER_THREAD], DecomposerT decomposer)
{
SortBlockedToStriped(
keys,
Expand Down
12 changes: 6 additions & 6 deletions cub/cub/block/radix_rank_sort_operations.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -583,8 +583,8 @@ private:
public:
template <class DecomposerT = detail::identity_decomposer_t>
static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE //
bit_ordered_type
In(bit_ordered_type key, DecomposerT decomposer = {})
bit_ordered_type
In(bit_ordered_type key, DecomposerT decomposer = {})
{
key = bit_ordered_conversion_policy::to_bit_ordered(decomposer, key);
_CCCL_IF_CONSTEXPR (IS_DESCENDING)
Expand All @@ -596,8 +596,8 @@ public:

template <class DecomposerT = detail::identity_decomposer_t>
static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE //
bit_ordered_type
Out(bit_ordered_type key, DecomposerT decomposer = {})
bit_ordered_type
Out(bit_ordered_type key, DecomposerT decomposer = {})
{
_CCCL_IF_CONSTEXPR (IS_DESCENDING)
{
Expand All @@ -609,8 +609,8 @@ public:

template <class DecomposerT = detail::identity_decomposer_t>
static _CCCL_HOST_DEVICE _CCCL_FORCEINLINE //
bit_ordered_type
DefaultKey(DecomposerT decomposer = {})
bit_ordered_type
DefaultKey(DecomposerT decomposer = {})
{
return IS_DESCENDING ? traits::min_raw_binary_key(decomposer) : traits::max_raw_binary_key(decomposer);
}
Expand Down
22 changes: 11 additions & 11 deletions cub/cub/detail/nvtx3.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1958,17 +1958,17 @@ NVTX3_INLINE_IF_REQUESTED namespace NVTX3_VERSION_NAMESPACE
*/
constexpr event_attributes() noexcept
: attributes_{
NVTX_VERSION, // version
sizeof(nvtxEventAttributes_t), // size
0, // category
NVTX_COLOR_UNKNOWN, // color type
0, // color value
NVTX_PAYLOAD_UNKNOWN, // payload type
0, // reserved 4B
{0}, // payload value (union) // NOTE(bgruber): added braces
NVTX_MESSAGE_UNKNOWN, // message type
{0} // message value (union) // NOTE(bgruber): added braces
}
NVTX_VERSION, // version
sizeof(nvtxEventAttributes_t), // size
0, // category
NVTX_COLOR_UNKNOWN, // color type
0, // color value
NVTX_PAYLOAD_UNKNOWN, // payload type
0, // reserved 4B
{0}, // payload value (union) // NOTE(bgruber): added braces
NVTX_MESSAGE_UNKNOWN, // message type
{0} // message value (union) // NOTE(bgruber): added braces
}
{}

/**
Expand Down
36 changes: 18 additions & 18 deletions cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -301,15 +301,15 @@ struct DeviceBatchMemcpyPolicy
{
static constexpr bool PREFER_POW2_BITS = true;
using AgentSmallBufferPolicyT = AgentBatchMemcpyPolicy<
BLOCK_THREADS,
BUFFERS_PER_THREAD,
TLEV_BYTES_PER_THREAD,
PREFER_POW2_BITS,
LARGE_BUFFER_BLOCK_THREADS * LARGE_BUFFER_BYTES_PER_THREAD,
WARP_LEVEL_THRESHOLD,
BLOCK_LEVEL_THRESHOLD,
buff_delay_constructor_t,
block_delay_constructor_t>;
BLOCK_THREADS,
BUFFERS_PER_THREAD,
TLEV_BYTES_PER_THREAD,
PREFER_POW2_BITS,
LARGE_BUFFER_BLOCK_THREADS * LARGE_BUFFER_BYTES_PER_THREAD,
WARP_LEVEL_THRESHOLD,
BLOCK_LEVEL_THRESHOLD,
buff_delay_constructor_t,
block_delay_constructor_t>;

using AgentLargeBufferPolicyT =
AgentBatchMemcpyLargeBuffersPolicy<LARGE_BUFFER_BLOCK_THREADS, LARGE_BUFFER_BYTES_PER_THREAD>;
Expand All @@ -320,15 +320,15 @@ struct DeviceBatchMemcpyPolicy
{
static constexpr bool PREFER_POW2_BITS = false;
using AgentSmallBufferPolicyT = AgentBatchMemcpyPolicy<
BLOCK_THREADS,
BUFFERS_PER_THREAD,
TLEV_BYTES_PER_THREAD,
PREFER_POW2_BITS,
LARGE_BUFFER_BLOCK_THREADS * LARGE_BUFFER_BYTES_PER_THREAD,
WARP_LEVEL_THRESHOLD,
BLOCK_LEVEL_THRESHOLD,
buff_delay_constructor_t,
block_delay_constructor_t>;
BLOCK_THREADS,
BUFFERS_PER_THREAD,
TLEV_BYTES_PER_THREAD,
PREFER_POW2_BITS,
LARGE_BUFFER_BLOCK_THREADS * LARGE_BUFFER_BYTES_PER_THREAD,
WARP_LEVEL_THRESHOLD,
BLOCK_LEVEL_THRESHOLD,
buff_delay_constructor_t,
block_delay_constructor_t>;

using AgentLargeBufferPolicyT =
AgentBatchMemcpyLargeBuffersPolicy<LARGE_BUFFER_BLOCK_THREADS, LARGE_BUFFER_BYTES_PER_THREAD>;
Expand Down
Loading

0 comments on commit 9df96fd

Please sign in to comment.