Skip to content

Commit

Permalink
Apply formatting to libcudacxx (#1610)
Browse files Browse the repository at this point in the history 
* Drop additional clang-format files

* Do not format autogenerated files

* Apply formatting to libcudacxx
  • Loading branch information
@miscco
miscco authored Apr 11, 2024
1 parent a923e71 commit baee3f5
Show file tree
Hide file tree
Showing 2,657 changed files with 114,301 additions and 99,528 deletions.
18 changes: 12 additions & 6 deletions .clang-format
View file
Original file line number Diff line number Diff line change
Expand Up @@ -107,18 +107,24 @@ IfMacros: [
IndentWrappedFunctionNames: false
IncludeBlocks: Regroup
IncludeCategories:
- Regex: '^<cub'
Priority: 1
- Regex: '^<(cuda/std/detail/__config|cub/config.cuh|thrust/detail/config.h)'
Priority: 0
SortPriority: 0
- Regex: '^<thrust'
- Regex: '^<cub/'
Priority: 2
SortPriority: 1
- Regex: '^<cuda'
- Regex: '^<thrust/'
Priority: 3
SortPriority: 2
- Regex: '^<[a-z_]*>$'
- Regex: '^<cuda/'
Priority: 4
SortPriority: 3
- Regex: '^<[a-z_]*>$'
Priority: 5
SortPriority: 4
- Regex: '^<cuda'
Priority: 0
SortPriority: 0
InsertBraces: true
IndentCaseLabels: true
InsertNewlineAtEOF: true
Expand Down Expand Up @@ -157,7 +163,7 @@ SpaceBeforeParens: ControlStatements
SpaceBeforeRangeBasedForLoopColon: true
SpaceInEmptyParentheses: false
SpacesBeforeTrailingComments: 1
SpacesInAngles: Leave
SpacesInAngles: Never
SpacesInCStyleCastParentheses: false
SpacesInParentheses: false
SpacesInSquareBrackets: false
Expand Down
1 change: 0 additions & 1 deletion .pre-commit-config.yaml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,6 @@ repos:
args: ["-fallback-style=none", "-style=file", "-i"]
exclude: |
(?x)^(
^libcudacxx/.*|
^thrust/.*
)
Expand Down
18 changes: 9 additions & 9 deletions cub/benchmarks/bench/copy/memcpy.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -101,15 +101,15 @@ struct offset_to_size_t
#if !TUNE_BASE
template <unsigned int MagicNs, unsigned int L2W, unsigned int DCID>
using delay_constructor_t =
nvbench::tl::get< DCID,
nvbench::type_list<cub::detail::no_delay_constructor_t<L2W>,
cub::detail::fixed_delay_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backoff_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backoff_jitter_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backoff_jitter_window_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backon_jitter_window_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backon_jitter_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backon_constructor_t<MagicNs, L2W>>>;
nvbench::tl::get<DCID,
nvbench::type_list<cub::detail::no_delay_constructor_t<L2W>,
cub::detail::fixed_delay_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backoff_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backoff_jitter_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backoff_jitter_window_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backon_jitter_window_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backon_jitter_constructor_t<MagicNs, L2W>,
cub::detail::exponential_backon_constructor_t<MagicNs, L2W>>>;

using buff_delay_constructor_t =
delay_constructor_t<TUNE_BUFF_MAGIC_NS, TUNE_BUFF_L2_WRITE_LATENCY_NS, TUNE_BUFF_DELAY_CONSTRUCTOR_ID>;
Expand Down
6 changes: 3 additions & 3 deletions cub/cub/agent/agent_batch_memcpy.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -556,9 +556,9 @@ private:
//---------------------------------------------------------------------
/// Internal load/store type. For byte-wise memcpy, a single-byte type
using AliasT =
typename ::cuda::std::conditional< IsMemcpy,
std::iterator_traits<char*>,
std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::type::value_type;
typename ::cuda::std::conditional<IsMemcpy,
std::iterator_traits<char*>,
std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::type::value_type;

/// Types of the input and output buffers
using InputBufferT = cub::detail::value_t<InputBufferIt>;
Expand Down
12 changes: 6 additions & 6 deletions cub/cub/agent/agent_histogram.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -225,9 +225,9 @@ struct AgentHistogram
// Wrap the native input pointer with CacheModifiedInputIterator
// or directly use the supplied input iterator type
using WrappedSampleIteratorT =
cub::detail::conditional_t< std::is_pointer<SampleIteratorT>::value,
CacheModifiedInputIterator<LOAD_MODIFIER, SampleT, OffsetT>,
SampleIteratorT>;
cub::detail::conditional_t<std::is_pointer<SampleIteratorT>::value,
CacheModifiedInputIterator<LOAD_MODIFIER, SampleT, OffsetT>,
SampleIteratorT>;

/// Pixel input iterator type (for applying cache modifier)
typedef CacheModifiedInputIterator<LOAD_MODIFIER, PixelT, OffsetT> WrappedPixelIteratorT;
Expand All @@ -236,13 +236,13 @@ struct AgentHistogram
typedef CacheModifiedInputIterator<LOAD_MODIFIER, VecT, OffsetT> WrappedVecsIteratorT;

/// Parameterized BlockLoad type for samples
typedef BlockLoad< SampleT, BLOCK_THREADS, SAMPLES_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadSampleT;
typedef BlockLoad<SampleT, BLOCK_THREADS, SAMPLES_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadSampleT;

/// Parameterized BlockLoad type for pixels
typedef BlockLoad< PixelT, BLOCK_THREADS, PIXELS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadPixelT;
typedef BlockLoad<PixelT, BLOCK_THREADS, PIXELS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadPixelT;

/// Parameterized BlockLoad type for vecs
typedef BlockLoad< VecT, BLOCK_THREADS, VECS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadVecT;
typedef BlockLoad<VecT, BLOCK_THREADS, VECS_PER_THREAD, AgentHistogramPolicyT::LOAD_ALGORITHM> BlockLoadVecT;

/// Shared memory type required by this thread block
struct _TempStorage
Expand Down
2 changes: 1 addition & 1 deletion cub/cub/agent/agent_radix_sort_downsweep.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -183,7 +183,7 @@ struct AgentRadixSortDownsweep

// Radix ranking type to use
using BlockRadixRankT =
cub::detail::block_radix_rank_t< RANK_ALGORITHM, BLOCK_THREADS, RADIX_BITS, IS_DESCENDING, SCAN_ALGORITHM>;
cub::detail::block_radix_rank_t<RANK_ALGORITHM, BLOCK_THREADS, RADIX_BITS, IS_DESCENDING, SCAN_ALGORITHM>;

// Digit extractor type
using fundamental_digit_extractor_t = BFEDigitExtractor<KeyT>;
Expand Down
2 changes: 1 addition & 1 deletion cub/cub/agent/agent_radix_sort_onesweep.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -81,7 +81,7 @@ template <int NOMINAL_BLOCK_THREADS_4B,
BlockScanAlgorithm _SCAN_ALGORITHM,
RadixSortStoreAlgorithm _STORE_ALGORITHM,
int _RADIX_BITS,
typename ScalingType = RegBoundScaling< NOMINAL_BLOCK_THREADS_4B, NOMINAL_ITEMS_PER_THREAD_4B, ComputeT> >
typename ScalingType = RegBoundScaling<NOMINAL_BLOCK_THREADS_4B, NOMINAL_ITEMS_PER_THREAD_4B, ComputeT>>
struct AgentRadixSortOnesweepPolicy : ScalingType
{
enum
Expand Down
6 changes: 3 additions & 3 deletions cub/cub/agent/agent_reduce.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -145,9 +145,9 @@ struct AgentReduce
// Wrap the native input pointer with CacheModifiedInputIterator
// or directly use the supplied input iterator type
using WrappedInputIteratorT =
cub::detail::conditional_t< std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentReducePolicy::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentReducePolicy::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;

/// Constants
static constexpr int BLOCK_THREADS = AgentReducePolicy::BLOCK_THREADS;
Expand Down
18 changes: 9 additions & 9 deletions cub/cub/agent/agent_reduce_by_key.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -225,27 +225,27 @@ struct AgentReduceByKey
// CacheModifiedValuesInputIterator or directly use the supplied input
// iterator type
using WrappedKeysInputIteratorT =
cub::detail::conditional_t< std::is_pointer<KeysInputIteratorT>::value,
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, KeyInputT, OffsetT>,
KeysInputIteratorT>;
cub::detail::conditional_t<std::is_pointer<KeysInputIteratorT>::value,
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, KeyInputT, OffsetT>,
KeysInputIteratorT>;

// Cache-modified Input iterator wrapper type (for applying cache modifier)
// for values Wrap the native input pointer with
// CacheModifiedValuesInputIterator or directly use the supplied input
// iterator type
using WrappedValuesInputIteratorT =
cub::detail::conditional_t< std::is_pointer<ValuesInputIteratorT>::value,
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
ValuesInputIteratorT>;
cub::detail::conditional_t<std::is_pointer<ValuesInputIteratorT>::value,
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
ValuesInputIteratorT>;

// Cache-modified Input iterator wrapper type (for applying cache modifier)
// for fixup values Wrap the native input pointer with
// CacheModifiedValuesInputIterator or directly use the supplied input
// iterator type
using WrappedFixupInputIteratorT =
cub::detail::conditional_t< std::is_pointer<AggregatesOutputIteratorT>::value,
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
AggregatesOutputIteratorT>;
cub::detail::conditional_t<std::is_pointer<AggregatesOutputIteratorT>::value,
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
AggregatesOutputIteratorT>;

// Reduce-value-by-segment scan operator
using ReduceBySegmentOpT = ReduceBySegmentOp<ReductionOpT>;
Expand Down
6 changes: 3 additions & 3 deletions cub/cub/agent/agent_rle.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -231,9 +231,9 @@ struct AgentRle
// Wrap the native input pointer with CacheModifiedVLengthnputIterator
// Directly use the supplied input iterator type
using WrappedInputIteratorT =
cub::detail::conditional_t< std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentRlePolicyT::LOAD_MODIFIER, T, OffsetT>,
InputIteratorT>;
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentRlePolicyT::LOAD_MODIFIER, T, OffsetT>,
InputIteratorT>;

// Parameterized BlockLoad type for data
using BlockLoadT =
Expand Down
6 changes: 3 additions & 3 deletions cub/cub/agent/agent_scan.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -157,9 +157,9 @@ struct AgentScan
// Wrap the native input pointer with CacheModifiedInputIterator
// or directly use the supplied input iterator type
using WrappedInputIteratorT =
cub::detail::conditional_t< std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentScanPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentScanPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;

// Constants
enum
Expand Down
12 changes: 6 additions & 6 deletions cub/cub/agent/agent_scan_by_key.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -152,14 +152,14 @@ struct AgentScanByKey
static constexpr int ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD;

using WrappedKeysInputIteratorT =
cub::detail::conditional_t< std::is_pointer<KeysInputIteratorT>::value,
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, KeyT, OffsetT>,
KeysInputIteratorT>;
cub::detail::conditional_t<std::is_pointer<KeysInputIteratorT>::value,
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, KeyT, OffsetT>,
KeysInputIteratorT>;

using WrappedValuesInputIteratorT =
cub::detail::conditional_t< std::is_pointer<ValuesInputIteratorT>::value,
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
ValuesInputIteratorT>;
cub::detail::conditional_t<std::is_pointer<ValuesInputIteratorT>::value,
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
ValuesInputIteratorT>;

using BlockLoadKeysT = BlockLoad<KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentScanByKeyPolicyT::LOAD_ALGORITHM>;

Expand Down
6 changes: 3 additions & 3 deletions cub/cub/agent/agent_segment_fixup.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -180,9 +180,9 @@ struct AgentSegmentFixup
// Wrap the native input pointer with CacheModifiedValuesInputIterator
// or directly use the supplied input iterator type
using WrappedFixupInputIteratorT =
cub::detail::conditional_t< std::is_pointer<AggregatesOutputIteratorT>::value,
CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER, ValueT, OffsetT>,
AggregatesOutputIteratorT>;
cub::detail::conditional_t<std::is_pointer<AggregatesOutputIteratorT>::value,
CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER, ValueT, OffsetT>,
AggregatesOutputIteratorT>;

// Reduce-value-by-segment scan operator
using ReduceBySegmentOpT = ReduceByKeyOp<cub::Sum>;
Expand Down
12 changes: 6 additions & 6 deletions cub/cub/agent/agent_select_if.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -219,17 +219,17 @@ struct AgentSelectIf
// Wrap the native input pointer with CacheModifiedValuesInputIterator
// or directly use the supplied input iterator type
using WrappedInputIteratorT =
cub::detail::conditional_t< ::cuda::std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;
cub::detail::conditional_t<::cuda::std::is_pointer<InputIteratorT>::value,
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;

// Cache-modified Input iterator wrapper type (for applying cache modifier) for values
// Wrap the native input pointer with CacheModifiedValuesInputIterator
// or directly use the supplied input iterator type
using WrappedFlagsInputIteratorT =
cub::detail::conditional_t< ::cuda::std::is_pointer<FlagsInputIteratorT>::value,
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, FlagT, OffsetT>,
FlagsInputIteratorT>;
cub::detail::conditional_t<::cuda::std::is_pointer<FlagsInputIteratorT>::value,
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, FlagT, OffsetT>,
FlagsInputIteratorT>;

// Parameterized BlockLoad type for input data
using BlockLoadT = BlockLoad<InputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSelectIfPolicyT::LOAD_ALGORITHM>;
Expand Down
20 changes: 10 additions & 10 deletions cub/cub/agent/agent_spmv_orig.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -228,17 +228,17 @@ struct AgentSpmv

/// Input iterator wrapper types (for applying cache modifiers)

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::ROW_OFFSETS_SEARCH_LOAD_MODIFIER, OffsetT, OffsetT>
typedef CacheModifiedInputIterator<AgentSpmvPolicyT::ROW_OFFSETS_SEARCH_LOAD_MODIFIER, OffsetT, OffsetT>
RowOffsetsSearchIteratorT;

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::ROW_OFFSETS_LOAD_MODIFIER, OffsetT, OffsetT> RowOffsetsIteratorT;
typedef CacheModifiedInputIterator<AgentSpmvPolicyT::ROW_OFFSETS_LOAD_MODIFIER, OffsetT, OffsetT> RowOffsetsIteratorT;

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::COLUMN_INDICES_LOAD_MODIFIER, OffsetT, OffsetT>
typedef CacheModifiedInputIterator<AgentSpmvPolicyT::COLUMN_INDICES_LOAD_MODIFIER, OffsetT, OffsetT>
ColumnIndicesIteratorT;

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::VALUES_LOAD_MODIFIER, ValueT, OffsetT> ValueIteratorT;
typedef CacheModifiedInputIterator<AgentSpmvPolicyT::VALUES_LOAD_MODIFIER, ValueT, OffsetT> ValueIteratorT;

typedef CacheModifiedInputIterator< AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER, ValueT, OffsetT>
typedef CacheModifiedInputIterator<AgentSpmvPolicyT::VECTOR_VALUES_LOAD_MODIFIER, ValueT, OffsetT>
VectorValueIteratorT;

// Tuple type for scanning (pairs accumulated segment-value with segment-index)
Expand All @@ -248,23 +248,23 @@ struct AgentSpmv
typedef ReduceByKeyOp<cub::Sum> ReduceBySegmentOpT;

// BlockReduce specialization
typedef BlockReduce< ValueT, BLOCK_THREADS, BLOCK_REDUCE_WARP_REDUCTIONS> BlockReduceT;
typedef BlockReduce<ValueT, BLOCK_THREADS, BLOCK_REDUCE_WARP_REDUCTIONS> BlockReduceT;

// BlockScan specialization
typedef BlockScan< KeyValuePairT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> BlockScanT;
typedef BlockScan<KeyValuePairT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> BlockScanT;

// BlockScan specialization
typedef BlockScan< ValueT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> BlockPrefixSumT;
typedef BlockScan<ValueT, BLOCK_THREADS, AgentSpmvPolicyT::SCAN_ALGORITHM> BlockPrefixSumT;

// BlockExchange specialization
typedef BlockExchange< ValueT, BLOCK_THREADS, ITEMS_PER_THREAD> BlockExchangeT;
typedef BlockExchange<ValueT, BLOCK_THREADS, ITEMS_PER_THREAD> BlockExchangeT;

/// Merge item type (either a non-zero value or a row-end offset)
union MergeItem
{
// Value type to pair with index type OffsetT
// (NullType if loading values directly during merge)
using MergeValueT = cub::detail::conditional_t< AgentSpmvPolicyT::DIRECT_LOAD_NONZEROS, NullType, ValueT>;
using MergeValueT = cub::detail::conditional_t<AgentSpmvPolicyT::DIRECT_LOAD_NONZEROS, NullType, ValueT>;

OffsetT row_end_offset;
MergeValueT nonzero;
Expand Down
6 changes: 3 additions & 3 deletions cub/cub/agent/agent_three_way_partition.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -197,9 +197,9 @@ struct AgentThreeWayPartition
static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD;

using WrappedInputIteratorT =
cub::detail::conditional_t< std::is_pointer<InputIteratorT>::value,
cub::CacheModifiedInputIterator<PolicyT::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
cub::CacheModifiedInputIterator<PolicyT::LOAD_MODIFIER, InputT, OffsetT>,
InputIteratorT>;

// Parameterized BlockLoad type for input data
using BlockLoadT = cub::BlockLoad<InputT, BLOCK_THREADS, ITEMS_PER_THREAD, PolicyT::LOAD_ALGORITHM>;
Expand Down
4 changes: 2 additions & 2 deletions cub/cub/agent/agent_unique_by_key.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -162,10 +162,10 @@ struct AgentUniqueByKey
ValueInputIteratorT>::type; // Directly use the supplied input iterator type

// Parameterized BlockLoad type for input data
using BlockLoadKeys = BlockLoad< KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentUniqueByKeyPolicyT::LOAD_ALGORITHM>;
using BlockLoadKeys = BlockLoad<KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentUniqueByKeyPolicyT::LOAD_ALGORITHM>;

// Parameterized BlockLoad type for flags
using BlockLoadValues = BlockLoad< ValueT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentUniqueByKeyPolicyT::LOAD_ALGORITHM>;
using BlockLoadValues = BlockLoad<ValueT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentUniqueByKeyPolicyT::LOAD_ALGORITHM>;

// Parameterized BlockDiscontinuity type for items
using BlockDiscontinuityKeys = cub::BlockDiscontinuity<KeyT, BLOCK_THREADS>;
Expand Down
Loading

0 comments on commit baee3f5

Please sign in to comment.