diff --git a/c/parallel/src/histogram.cu b/c/parallel/src/histogram.cu index d01e7cc63119..c80face8d5aa 100644 --- a/c/parallel/src/histogram.cu +++ b/c/parallel/src/histogram.cu @@ -56,7 +56,7 @@ struct histogram_kernel_source } template ", chained_policy_t, - privatized_smem_bins, + privatization_mode_t, num_channels, num_active_channels, samples_iterator_t, @@ -324,17 +324,26 @@ static_assert(device_histogram_policy()(detail::current_tuning_cc()) == {4}, "Ho fflush(stdout); #endif - // TODO: This is tricky because we need to know the input to set this to a - // value greater than 0 (see dispatch_histogram.cuh), but we don't have this - // information here. - const int privatized_smem_bins = - num_output_levels_val - 1 > cub::detail::histogram::max_privatized_smem_bins ? 0 : 256; - const bool is_byte_sample = d_samples.value_type.size == 1; + const int num_privatized_bins = + is_byte_sample ? cub::detail::histogram::byte_sample_privatized_levels - 1 : num_output_levels_val - 1; + const int counter_size_bytes = static_cast(d_output_histograms.value_type.size); + const auto privatization = + is_evenly_segmented + ? cub::detail::histogram::select_privatization_mode_for_counter_size( + active_policy, num_privatized_bins, counter_size_bytes) + : cub::detail::histogram::select_privatization_mode_for_counter_size( + active_policy, num_privatized_bins, counter_size_bytes); + const std::string_view privatization_mode_t = + privatization == cub::detail::histogram::privatization_mode::static_smem + ? "cub::detail::histogram::HistogramPrivatizedStaticSmem" + : privatization == cub::detail::histogram::privatization_mode::dynamic_smem + ? "cub::detail::histogram::HistogramPrivatizedDynamicSmem" + : "cub::detail::histogram::HistogramPrivatizedGmem"; std::string init_kernel_name = histogram::get_init_kernel_name(num_active_channels, counter_cpp, offset_cpp); std::string sweep_kernel_name = histogram::get_sweep_kernel_name( - privatized_smem_bins, + privatization_mode_t, num_channels, num_active_channels, d_samples, @@ -574,6 +583,7 @@ CUresult cccl_device_histogram_even_impl( indirect_arg_t, // LevelT OffsetT, // OffsetT cub::detail::histogram::policy_selector, // PolicySelector + void, // PrivatizedCounterT: C Parallel preserves the counter width after type erasure indirect_arg_t, // SampleT histogram::histogram_kernel_source, // KernelSource cub::detail::CudaDriverLauncherFactory // KernelLauncherFactory diff --git a/cub/benchmarks/bench/histogram/even.cu b/cub/benchmarks/bench/histogram/even.cu index a175139b1b99..457e58c1750e 100644 --- a/cub/benchmarks/bench/histogram/even.cu +++ b/cub/benchmarks/bench/histogram/even.cu @@ -9,7 +9,6 @@ // %RANGE% TUNE_THREADS tpb 128:1024:32 // %RANGE% TUNE_RLE_COMPRESS rle 0:1:1 // %RANGE% TUNE_WORK_STEALING ws 0:1:1 -// %RANGE% TUNE_MEM_PREFERENCE mem 0:2:1 // %RANGE% TUNE_LOAD ld 0:2:1 // %RANGE% TUNE_LOAD_ALGORITHM_ID laid 0:2:1 // %RANGE% TUNE_VEC_SIZE_POW vec 0:2:1 @@ -52,7 +51,7 @@ static void even(nvbench::state& state, nvbench::type_list{}) + cuda::execution::tune(histogram_tuning_policy_selector{}) #endif // !TUNE_BASE ); _CCCL_TRY_CUDA_API( diff --git a/cub/benchmarks/bench/histogram/histogram_common.cuh b/cub/benchmarks/bench/histogram/histogram_common.cuh index c87565264042..7e0735c12156 100644 --- a/cub/benchmarks/bench/histogram/histogram_common.cuh +++ b/cub/benchmarks/bench/histogram/histogram_common.cuh @@ -19,14 +19,6 @@ # define TUNE_VEC_SIZE (1 << TUNE_VEC_SIZE_POW) -# if TUNE_MEM_PREFERENCE == 0 -constexpr cub::BlockHistogramMemoryPreference MEM_PREFERENCE = cub::GMEM; -# elif TUNE_MEM_PREFERENCE == 1 -constexpr cub::BlockHistogramMemoryPreference MEM_PREFERENCE = cub::SMEM; -# else // TUNE_MEM_PREFERENCE == 2 -constexpr cub::BlockHistogramMemoryPreference MEM_PREFERENCE = cub::BLEND; -# endif // TUNE_MEM_PREFERENCE - # if TUNE_LOAD_ALGORITHM_ID == 0 # define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT # elif TUNE_LOAD_ALGORITHM_ID == 1 @@ -35,25 +27,38 @@ constexpr cub::BlockHistogramMemoryPreference MEM_PREFERENCE = cub::BLEND; # define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_STRIPED # endif // TUNE_LOAD_ALGORITHM_ID -template -struct bench_policy_selector +// Only generated tuning variants instantiate this selector. The `.base` target used for +// production-policy comparisons defines TUNE_BASE=1 and calls DeviceHistogram without a +// tuning environment, so it exercises the shipping selector unchanged. +// +// A generated tuning point supplies one candidate kernel configuration. Apply that same +// candidate to every privatization mode so the tuner measures the candidate independently +// of the runtime bin count selected by the production storage thresholds. +template +struct histogram_tuning_policy_selector { - _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const -> cub::HistogramPolicy + _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const -> cub::HistogramPolicy { constexpr cub::BlockLoadAlgorithm load_algorithm = (TUNE_LOAD_ALGORITHM == cub::BLOCK_LOAD_STRIPED) ? (NUM_CHANNELS == 1 ? cub::BLOCK_LOAD_STRIPED : cub::BLOCK_LOAD_DIRECT) : TUNE_LOAD_ALGORITHM; - return {TUNE_THREADS, - TUNE_ITEMS, - TUNE_VEC_SIZE, - load_algorithm, - TUNE_LOAD_MODIFIER, - TUNE_RLE_COMPRESS, - MEM_PREFERENCE, - TUNE_WORK_STEALING, - 2048}; // TODO(bgruber): make tunable + constexpr auto sweep = cub::HistogramPrivatizationPolicy{ + TUNE_THREADS, + TUNE_ITEMS, + TUNE_VEC_SIZE, + load_algorithm, + TUNE_LOAD_MODIFIER, + TUNE_RLE_COMPRESS, + TUNE_WORK_STEALING}; + auto policy = + cub::detail::histogram::policy_selector_from_types{}( + cc); + policy.gmem = sweep; + policy.static_smem = sweep; + policy.dynamic_smem = sweep; + return policy; } }; #endif // !TUNE_BASE diff --git a/cub/benchmarks/bench/histogram/multi/even.cu b/cub/benchmarks/bench/histogram/multi/even.cu index 4131206a8bd7..4bb60a72c26c 100644 --- a/cub/benchmarks/bench/histogram/multi/even.cu +++ b/cub/benchmarks/bench/histogram/multi/even.cu @@ -9,7 +9,6 @@ // %RANGE% TUNE_THREADS tpb 128:1024:32 // %RANGE% TUNE_RLE_COMPRESS rle 0:1:1 // %RANGE% TUNE_WORK_STEALING ws 0:1:1 -// %RANGE% TUNE_MEM_PREFERENCE mem 0:2:1 // %RANGE% TUNE_LOAD ld 0:2:1 // %RANGE% TUNE_LOAD_ALGORITHM_ID laid 0:2:1 // %RANGE% TUNE_VEC_SIZE_POW vec 0:2:1 @@ -65,7 +64,8 @@ static void even(nvbench::state& state, nvbench::type_list{}) + cuda::execution::tune( + histogram_tuning_policy_selector{}) #endif // !TUNE_BASE ); _CCCL_TRY_CUDA_API( diff --git a/cub/benchmarks/bench/histogram/multi/range.cu b/cub/benchmarks/bench/histogram/multi/range.cu index dcb402d7db20..7be2518dfca8 100644 --- a/cub/benchmarks/bench/histogram/multi/range.cu +++ b/cub/benchmarks/bench/histogram/multi/range.cu @@ -11,7 +11,6 @@ // %RANGE% TUNE_THREADS tpb 128:1024:32 // %RANGE% TUNE_RLE_COMPRESS rle 0:1:1 // %RANGE% TUNE_WORK_STEALING ws 0:1:1 -// %RANGE% TUNE_MEM_PREFERENCE mem 0:2:1 // %RANGE% TUNE_LOAD ld 0:2:1 // %RANGE% TUNE_LOAD_ALGORITHM_ID laid 0:2:1 // %RANGE% TUNE_VEC_SIZE_POW vec 0:2:1 @@ -65,7 +64,8 @@ static void range(nvbench::state& state, nvbench::type_list{}) + cuda::execution::tune( + histogram_tuning_policy_selector{}) #endif // !TUNE_BASE ); _CCCL_TRY_CUDA_API( diff --git a/cub/benchmarks/bench/histogram/range.cu b/cub/benchmarks/bench/histogram/range.cu index e1e808466395..938fa11dc4c6 100644 --- a/cub/benchmarks/bench/histogram/range.cu +++ b/cub/benchmarks/bench/histogram/range.cu @@ -11,7 +11,6 @@ // %RANGE% TUNE_THREADS tpb 128:1024:32 // %RANGE% TUNE_RLE_COMPRESS rle 0:1:1 // %RANGE% TUNE_WORK_STEALING ws 0:1:1 -// %RANGE% TUNE_MEM_PREFERENCE mem 0:2:1 // %RANGE% TUNE_LOAD ld 0:2:1 // %RANGE% TUNE_LOAD_ALGORITHM_ID laid 0:2:1 // %RANGE% TUNE_VEC_SIZE_POW vec 0:2:1 @@ -51,7 +50,7 @@ static void range(nvbench::state& state, nvbench::type_list{}) + cuda::execution::tune(histogram_tuning_policy_selector{}) #endif // !TUNE_BASE ); _CCCL_TRY_CUDA_API( diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index daed613be77d..48f2fface97f 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -21,11 +21,9 @@ #include #include #include +#include #include -#include -#include -#include #include #include #include @@ -33,65 +31,15 @@ CUB_NAMESPACE_BEGIN -enum BlockHistogramMemoryPreference -{ - GMEM, - SMEM, - BLEND -}; - -#if _CCCL_HOSTED() -namespace detail -{ -[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr const char* to_string(BlockHistogramMemoryPreference mempref) noexcept -{ - switch (mempref) - { - case GMEM: - return "GMEM"; - case SMEM: - return "SMEM"; - case BLEND: - return "BLEND"; - } - return ""; -} -} // namespace detail - -inline ::std::ostream& operator<<(::std::ostream& os, BlockHistogramMemoryPreference mempref) -{ - return os << CUB_NS_QUALIFIER::detail::to_string(mempref); -} -#endif // _CCCL_HOSTED() - -CUB_NAMESPACE_END - -#if __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) -template <::cuda::std::same_as CharT> -struct std::formatter : formatter -{ - template - auto format(const CUB_NS_QUALIFIER::BlockHistogramMemoryPreference& mempref, FmtCtx& ctx) const - { - return formatter::format(CUB_NS_QUALIFIER::detail::to_string(mempref), ctx); - } -}; -#endif // __cpp_lib_format >= 201907L && !defined(_CCCL_DOXYGEN_INVOKED) - -CUB_NAMESPACE_BEGIN - -namespace detail -{ -//! Parameterizable tuning policy type for AgentHistogram +//! Deprecated [Since 3.5] template -struct agent_histogram_policy +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") AgentHistogramPolicy { /// Threads per thread block static constexpr int BLOCK_THREADS = ThreadsPerBlock; @@ -101,9 +49,6 @@ struct agent_histogram_policy /// Whether to perform localized RLE to compress samples before histogramming static constexpr bool IS_RLE_COMPRESS = RleCompress; - /// Whether to prefer privatized shared-memory bins (versus privatized global-memory bins) - static constexpr BlockHistogramMemoryPreference MEM_PREFERENCE = MemoryPreference; - /// Whether to dequeue tiles from a global work queue static constexpr bool IS_WORK_STEALING = WorkStealing; @@ -117,30 +62,33 @@ struct agent_histogram_policy ///< Cache load modifier for reading input elements static constexpr CacheLoadModifier LOAD_MODIFIER = LoadModifier; }; -} // namespace detail - -//! Deprecated [Since 3.5] -template -using AgentHistogramPolicy - CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") = detail::agent_histogram_policy< - ThreadsPerBlock, - PixelsPerThread, - LoadAlgorithm, - LoadModifier, - RleCompress, - MemoryPreference, - WorkStealing, - VecSize>; namespace detail::histogram { +struct HistogramPrivatizedStaticSmem +{}; + +struct HistogramPrivatizedDynamicSmem +{}; + +struct HistogramPrivatizedGmem +{}; + +inline constexpr auto histogram_privatized_static_smem = HistogramPrivatizedStaticSmem{}; +inline constexpr auto histogram_privatized_dynamic_smem = HistogramPrivatizedDynamicSmem{}; +inline constexpr auto histogram_privatized_gmem = HistogramPrivatizedGmem{}; + +template +inline constexpr bool is_privatized_static_smem_v = + ::cuda::std::is_same_v; + +template +inline constexpr bool is_privatized_dynamic_smem_v = + ::cuda::std::is_same_v; + +template +inline constexpr bool is_privatized_gmem_v = ::cuda::std::is_same_v; + // Return a native pixel pointer (specialized for CacheModifiedInputIterator types) template _CCCL_DEVICE _CCCL_FORCEINLINE auto NativePointer(CacheModifiedInputIterator itr) @@ -158,12 +106,11 @@ _CCCL_DEVICE _CCCL_FORCEINLINE auto NativePointer(IteratorT itr) //! @brief AgentHistogram implements a stateful abstraction of CUDA thread blocks for participating //! in device-wide histogram . //! -//! @tparam AgentHistogramPolicyT -//! Parameterized AgentHistogramPolicy tuning policy type +//! @tparam PolicySelector +//! Selector that returns the active HistogramPolicy. //! -//! @tparam PrivatizedSmemBins -//! Number of privatized shared-memory histogram bins of any channel. Zero indicates privatized -//! counters to be maintained in device-accessible memory. +//! @tparam PrivatizationMode +//! Storage mode for the privatized histogram. //! //! @tparam NumChannels //! Number of channels interleaved in the input data. Supports up to four channels. @@ -175,7 +122,7 @@ _CCCL_DEVICE _CCCL_FORCEINLINE auto NativePointer(IteratorT itr) //! Random-access input iterator type for reading samples //! //! @tparam CounterT -//! Integer type for counting sample occurrences per histogram bin +//! Integer type for per-block privatized histogram bins //! //! @tparam PrivatizedDecodeOpT //! The transform operator type for determining privatized counter indices from samples, one for @@ -187,29 +134,48 @@ _CCCL_DEVICE _CCCL_FORCEINLINE auto NativePointer(IteratorT itr) //! //! @tparam OffsetT //! Signed integer type for global offsets -template + typename OffsetT, + typename OutputCounterT = CounterT> struct AgentHistogram { - static constexpr int vec_size = AgentHistogramPolicyT::VEC_SIZE; - static constexpr int threads_per_block = AgentHistogramPolicyT::BLOCK_THREADS; - static constexpr int pixels_per_thread = AgentHistogramPolicyT::PIXELS_PER_THREAD; - static constexpr int samples_per_thread = pixels_per_thread * NumChannels; + static_assert(sizeof(CounterT) <= sizeof(OutputCounterT), + "The output histogram counter must be at least as wide as the local counter"); + static_assert(is_privatized_static_smem_v || is_privatized_dynamic_smem_v + || is_privatized_gmem_v); + static constexpr bool uses_static_smem = is_privatized_static_smem_v; + static constexpr bool uses_dynamic_smem = is_privatized_dynamic_smem_v; + static constexpr bool uses_gmem = is_privatized_gmem_v; + static constexpr auto policy = current_policy(); + static constexpr auto privatization_policy = + uses_static_smem ? policy.static_smem + : uses_dynamic_smem + ? policy.dynamic_smem + : policy.gmem; + static constexpr int privatized_static_smem_bins = + uses_static_smem ? policy.max_privatized_static_smem_single_channel_bytes / int{sizeof(CounterT)} : 0; + static_assert(!uses_static_smem || privatized_static_smem_bins > 0, + "Static-SMEM privatization requires room for at least one bin"); + static constexpr int vec_size = privatization_policy.vec_size; + static constexpr int threads_per_block = privatization_policy.threads_per_block; + static constexpr int items_per_thread = privatization_policy.items_per_thread; + static constexpr int samples_per_thread = items_per_thread * NumChannels; static constexpr int vecs_per_thread = samples_per_thread / vec_size; - static constexpr int tile_pixels = pixels_per_thread * threads_per_block; + static constexpr int tile_pixels = items_per_thread * threads_per_block; static constexpr int tile_samples = samples_per_thread * threads_per_block; - static constexpr bool is_rle_compress = AgentHistogramPolicyT::IS_RLE_COMPRESS; - static constexpr bool is_work_stealing = AgentHistogramPolicyT::IS_WORK_STEALING; - static constexpr CacheLoadModifier load_modifier = AgentHistogramPolicyT::LOAD_MODIFIER; - static constexpr auto mem_preference = - (PrivatizedSmemBins > 0) ? BlockHistogramMemoryPreference{AgentHistogramPolicyT::MEM_PREFERENCE} : GMEM; + static constexpr bool is_rle_compress = privatization_policy.rle_compress; + static constexpr bool is_work_stealing = privatization_policy.work_stealing; + static constexpr CacheLoadModifier load_modifier = privatization_policy.load_modifier; using SampleT = it_value_t; using PixelT = typename CubVector::Type; @@ -225,15 +191,17 @@ struct AgentHistogram using WrappedPixelIteratorT = CacheModifiedInputIterator; using WrappedVecsIteratorT = CacheModifiedInputIterator; using BlockLoadSampleT = - BlockLoad; - using BlockLoadPixelT = - BlockLoad; - using BlockLoadVecT = BlockLoad; + BlockLoad; + using BlockLoadPixelT = BlockLoad; + using BlockLoadVecT = BlockLoad; struct _TempStorage { - // Smem needed for block-privatized smem histogram (with 1 word of padding) - CounterT histograms[NumActiveChannels][PrivatizedSmemBins + 1]; + // The one-element fallback keeps this type well-formed for modes that do not + // use the compile-time-sized histogram. Static-SMEM mode uses exactly the + // configured number of bins; out-of-range samples are rejected before the + // atomic update and therefore require no padding bin. + CounterT privatized_histogram[NumActiveChannels][privatized_static_smem_bins > 0 ? privatized_static_smem_bins : 1]; int tile_idx; union @@ -246,79 +214,64 @@ struct AgentHistogram using TempStorage = Uninitialized<_TempStorage>; - _TempStorage& temp_storage; + _TempStorage& static_smem_storage; WrappedSampleIteratorT d_wrapped_samples; // with cache modifier applied, if possible SampleT* d_native_samples; // possibly nullptr if unavailable - const int* num_output_bins; // one for each channel - const int* num_privatized_bins; // one for each channel - CounterT* d_privatized_histograms[NumActiveChannels]; // one for each channel - CounterT** d_output_histograms; // in global memory + const int* num_output_bins; // array of ints, one for each channel + const int* num_privatized_bins; // array of ints, one for each channel + CounterT* gmem_privatized_histograms[NumActiveChannels]; // one for each channel + CounterT* dynamic_smem_privatized_histograms[NumActiveChannels]; // dynamic shared-memory channel bases, when enabled + OutputCounterT** output_histogram; // final output, in global memory const OutputDecodeOpT* output_decode_op; // determines output bin-id from privatized counter index, one for each // channel - const PrivatizedDecodeOpT* privatized_decode_op; // determines privatized counter index from sample, one for each - // channel - bool prefer_smem; // for privatized counterss - - template - _CCCL_DEVICE _CCCL_FORCEINLINE void ZeroBinCounters(TwoDimSubscriptableCounterT& privatized_histograms) + PrivatizedDecodeOpT* privatized_decode_op; // determines privatized counter index from sample, one for each channel + _CCCL_DEVICE _CCCL_FORCEINLINE CounterT* PrivatizedHistogram(int channel) { - _CCCL_PRAGMA_UNROLL_FULL() - for (int ch = 0; ch < NumActiveChannels; ++ch) + if constexpr (uses_dynamic_smem) { - for (int bin = static_cast(threadIdx.x); bin < num_privatized_bins[ch]; bin += threads_per_block) - { - privatized_histograms[ch][bin] = 0; - } + return dynamic_smem_privatized_histograms[channel]; + } + else if constexpr (uses_static_smem) + { + return static_smem_storage.privatized_histogram[channel]; + } + else + { + return gmem_privatized_histograms[channel]; } - - // TODO(bgruber): do we also need the __syncthreads() when prefer_smem is false? - // Barrier to make sure all threads are done updating counters - __syncthreads(); } - // Update final output histograms from privatized histograms - template - _CCCL_DEVICE _CCCL_FORCEINLINE void StoreOutput(TwoDimSubscriptableCounterT& privatized_histograms) + _CCCL_DEVICE _CCCL_FORCEINLINE void ZeroBinCounters() { - // Barrier to make sure all threads are done updating counters - __syncthreads(); - - // Apply privatized bin counts to output bin counts _CCCL_PRAGMA_UNROLL_FULL() for (int ch = 0; ch < NumActiveChannels; ++ch) { - const int channel_bins = num_privatized_bins[ch]; - for (int bin = static_cast(threadIdx.x); bin < channel_bins; bin += threads_per_block) + CounterT* privatized_histogram = PrivatizedHistogram(ch); + for (int bin = static_cast(threadIdx.x); bin < num_privatized_bins[ch]; bin += threads_per_block) { - int output_bin = -1; - const CounterT count = privatized_histograms[ch][bin]; - const bool is_valid = count > 0; - output_decode_op[ch].template BinSelect(static_cast(bin), output_bin, is_valid); - - if (output_bin >= 0) - { - atomicAdd(&d_output_histograms[ch][output_bin], count); - } + privatized_histogram[bin] = 0; } } + + // Barrier to make sure all threads are done updating counters + __syncthreads(); } // Accumulate pixels. Specialized for RLE compression. - template _CCCL_DEVICE _CCCL_FORCEINLINE void AccumulatePixels( - SampleT samples[pixels_per_thread][NumChannels], - bool is_valid[pixels_per_thread], - TwoDimSubscriptableCounterT& privatized_histograms, + SampleT samples[items_per_thread][NumChannels], + bool is_valid[items_per_thread], ::cuda::std::true_type is_rle_compress) { _CCCL_PRAGMA_UNROLL_FULL() for (int ch = 0; ch < NumActiveChannels; ++ch) { + CounterT* privatized_histogram = PrivatizedHistogram(ch); // Bin pixels - int bins[pixels_per_thread]; + int bins[items_per_thread]; _CCCL_PRAGMA_UNROLL_FULL() - for (int pixel = 0; pixel < pixels_per_thread; ++pixel) + for (int pixel = 0; pixel < items_per_thread; ++pixel) { bins[pixel] = -1; privatized_decode_op[ch].template BinSelect(samples[pixel][ch], bins[pixel], is_valid[pixel]); @@ -327,15 +280,13 @@ struct AgentHistogram CounterT accumulator = 1; _CCCL_PRAGMA_UNROLL_FULL() - for (int pixel = 0; pixel < pixels_per_thread - 1; ++pixel) + for (int pixel = 0; pixel < items_per_thread - 1; ++pixel) { if (bins[pixel] != bins[pixel + 1]) { if (bins[pixel] >= 0) { - NV_IF_ELSE_TARGET(NV_PROVIDES_SM_60, - (atomicAdd_block(privatized_histograms[ch] + bins[pixel], accumulator);), - (atomicAdd(privatized_histograms[ch] + bins[pixel], accumulator);)); + atomicAdd_block(privatized_histogram + bins[pixel], accumulator); } accumulator = 0; @@ -344,36 +295,31 @@ struct AgentHistogram } // Last pixel - if (bins[pixels_per_thread - 1] >= 0) + if (bins[items_per_thread - 1] >= 0) { - NV_IF_ELSE_TARGET(NV_PROVIDES_SM_60, - (atomicAdd_block(privatized_histograms[ch] + bins[pixels_per_thread - 1], accumulator);), - (atomicAdd(privatized_histograms[ch] + bins[pixels_per_thread - 1], accumulator);)); + atomicAdd_block(privatized_histogram + bins[items_per_thread - 1], accumulator); } } } // Accumulate pixels. Specialized for individual accumulation of each pixel. - template _CCCL_DEVICE _CCCL_FORCEINLINE void AccumulatePixels( - SampleT samples[pixels_per_thread][NumChannels], - bool is_valid[pixels_per_thread], - TwoDimSubscriptableCounterT& privatized_histograms, + SampleT samples[items_per_thread][NumChannels], + bool is_valid[items_per_thread], ::cuda::std::false_type is_rle_compress) { _CCCL_PRAGMA_UNROLL_FULL() - for (int pixel = 0; pixel < pixels_per_thread; ++pixel) + for (int pixel = 0; pixel < items_per_thread; ++pixel) { _CCCL_PRAGMA_UNROLL_FULL() for (int ch = 0; ch < NumActiveChannels; ++ch) { - int bin = -1; + CounterT* privatized_histogram = PrivatizedHistogram(ch); + int bin = -1; privatized_decode_op[ch].template BinSelect(samples[pixel][ch], bin, is_valid[pixel]); if (bin >= 0) { - NV_IF_ELSE_TARGET(NV_PROVIDES_SM_60, - (atomicAdd_block(privatized_histograms[ch] + bin, 1);), - (atomicAdd(privatized_histograms[ch] + bin, 1);)); + atomicAdd_block(privatized_histogram + bin, 1); } } } @@ -381,27 +327,27 @@ struct AgentHistogram // Load full, aligned tile using pixel iterator _CCCL_DEVICE _CCCL_FORCEINLINE void - LoadFullAlignedTile(OffsetT block_offset, SampleT (&samples)[pixels_per_thread][NumChannels]) + LoadFullAlignedTile(OffsetT block_offset, SampleT (&samples)[items_per_thread][NumChannels]) { if constexpr (NumActiveChannels == 1) { using AliasedVecs = VecT[vecs_per_thread]; WrappedVecsIteratorT d_wrapped_vecs(reinterpret_cast(d_native_samples + block_offset)); // Load using a wrapped vec iterator - BlockLoadVecT{temp_storage.vec_load}.Load(d_wrapped_vecs, reinterpret_cast(samples)); + BlockLoadVecT{static_smem_storage.vec_load}.Load(d_wrapped_vecs, reinterpret_cast(samples)); } else { - using AliasedPixels = PixelT[pixels_per_thread]; + using AliasedPixels = PixelT[items_per_thread]; WrappedPixelIteratorT d_wrapped_pixels(reinterpret_cast(d_native_samples + block_offset)); // Load using a wrapped pixel iterator - BlockLoadPixelT{temp_storage.pixel_load}.Load(d_wrapped_pixels, reinterpret_cast(samples)); + BlockLoadPixelT{static_smem_storage.pixel_load}.Load(d_wrapped_pixels, reinterpret_cast(samples)); } } template _CCCL_DEVICE _CCCL_FORCEINLINE void - LoadTile(OffsetT block_offset, int valid_samples, SampleT (&samples)[pixels_per_thread][NumChannels]) + LoadTile(OffsetT block_offset, int valid_samples, SampleT (&samples)[items_per_thread][NumChannels]) { if constexpr (IsFullTile) { @@ -413,7 +359,7 @@ struct AgentHistogram { // Load using sample iterator using AliasedSamples = SampleT[samples_per_thread]; - BlockLoadSampleT{temp_storage.sample_load}.Load( + BlockLoadSampleT{static_smem_storage.sample_load}.Load( d_wrapped_samples + block_offset, reinterpret_cast(samples)); } } @@ -422,28 +368,28 @@ struct AgentHistogram if constexpr (IsAligned) { // Load partially-full, aligned tile using the pixel iterator - using AliasedPixels = PixelT[pixels_per_thread]; + using AliasedPixels = PixelT[items_per_thread]; WrappedPixelIteratorT d_wrapped_pixels((PixelT*) (d_native_samples + block_offset)); int valid_pixels = valid_samples / NumChannels; // Load using a wrapped pixel iterator - BlockLoadPixelT{temp_storage.pixel_load}.Load( + BlockLoadPixelT{static_smem_storage.pixel_load}.Load( d_wrapped_pixels, reinterpret_cast(samples), valid_pixels); } else { using AliasedSamples = SampleT[samples_per_thread]; - BlockLoadSampleT{temp_storage.sample_load}.Load( + BlockLoadSampleT{static_smem_storage.sample_load}.Load( d_wrapped_samples + block_offset, reinterpret_cast(samples), valid_samples); } } } template - _CCCL_DEVICE _CCCL_FORCEINLINE void MarkValid(bool (&is_valid)[pixels_per_thread], int valid_samples) + _CCCL_DEVICE _CCCL_FORCEINLINE void MarkValid(bool (&is_valid)[items_per_thread], int valid_samples) { _CCCL_PRAGMA_UNROLL_FULL() - for (int pixel = 0; pixel < pixels_per_thread; ++pixel) + for (int pixel = 0; pixel < items_per_thread; ++pixel) { if constexpr (IsStriped) { @@ -451,7 +397,7 @@ struct AgentHistogram } else { - is_valid[pixel] = IsFullTile || (((threadIdx.x * pixels_per_thread + pixel) * NumChannels) < valid_samples); + is_valid[pixel] = IsFullTile || (((threadIdx.x * items_per_thread + pixel) * NumChannels) < valid_samples); } } } @@ -466,20 +412,13 @@ struct AgentHistogram template _CCCL_DEVICE _CCCL_FORCEINLINE void ConsumeTile(OffsetT block_offset, int valid_samples) { - SampleT samples[pixels_per_thread][NumChannels]; - bool is_valid[pixels_per_thread]; + SampleT samples[items_per_thread][NumChannels]; + bool is_valid[items_per_thread]; LoadTile(block_offset, valid_samples, samples); - MarkValid(is_valid, valid_samples); + MarkValid(is_valid, valid_samples); - if (prefer_smem) - { - AccumulatePixels(samples, is_valid, temp_storage.histograms, ::cuda::std::bool_constant{}); - } - else - { - AccumulatePixels(samples, is_valid, d_privatized_histograms, ::cuda::std::bool_constant{}); - } + AccumulatePixels(samples, is_valid, ::cuda::std::bool_constant{}); } //! @brief Consume row tiles. Specialized for work-stealing from queue @@ -533,12 +472,12 @@ struct AgentHistogram // Get next tile if (threadIdx.x == 0) { - temp_storage.tile_idx = tile_queue.Drain(1) + num_even_share_tiles; + static_smem_storage.tile_idx = tile_queue.Drain(1) + num_even_share_tiles; } __syncthreads(); - tile_idx = temp_storage.tile_idx; + tile_idx = static_smem_storage.tile_idx; } } @@ -586,8 +525,8 @@ struct AgentHistogram //! @brief Constructor //! - //! @param temp_storage - //! Reference to temp_storage + //! @param static_smem_storage + //! Shared storage used by the agent //! //! @param d_samples //! Input data to reduce @@ -598,47 +537,58 @@ struct AgentHistogram //! @param num_privatized_bins //! The number bins per privatized histogram //! - //! @param d_output_histograms + //! @param output_histogram //! Reference to final output histograms //! - //! @param d_privatized_histograms - //! Reference to privatized histograms + //! @param gmem_privatized_histograms + //! Global-memory privatized histograms, or `nullptr` entries for a shared-memory mode //! //! @param output_decode_op //! The transform operator for determining output bin-ids from privatized counter indices, one for each channel //! //! @param privatized_decode_op //! The transform operator for determining privatized counter indices from samples, one for each channel + //! + //! @param dynamic_smem_privatized_histograms + //! Base of the runtime-sized shared-memory histogram, or `nullptr` for a static-SMEM or global-memory mode _CCCL_DEVICE _CCCL_FORCEINLINE AgentHistogram( - TempStorage& temp_storage, + TempStorage& static_smem_storage, SampleIteratorT d_samples, const int* num_output_bins, const int* num_privatized_bins, - CounterT** d_output_histograms, - CounterT** d_privatized_histograms, + OutputCounterT** output_histogram, + CounterT** gmem_privatized_histograms, const OutputDecodeOpT* output_decode_op, - const PrivatizedDecodeOpT* privatized_decode_op) - : temp_storage(temp_storage.Alias()) + PrivatizedDecodeOpT* privatized_decode_op, + CounterT* dynamic_smem_privatized_histograms) + : static_smem_storage(static_smem_storage.Alias()) , d_wrapped_samples(d_samples) , d_native_samples(NativePointer(d_wrapped_samples)) , num_output_bins(num_output_bins) , num_privatized_bins(num_privatized_bins) - , d_output_histograms(d_output_histograms) + , output_histogram(output_histogram) , output_decode_op(output_decode_op) , privatized_decode_op(privatized_decode_op) - , prefer_smem((mem_preference == SMEM) ? true : // prefer smem privatized histograms - (mem_preference == GMEM) ? false - : // prefer gmem privatized histograms - blockIdx.x & 1) // prefer blended privatized histograms { - const int blockId = static_cast((blockIdx.y * gridDim.x) + blockIdx.x); - - // TODO(bgruber): d_privatized_histograms seems only used when !prefer_smem, can we skip it if prefer_smem? - // Initialize the locations of this block's privatized histograms - for (int ch = 0; ch < NumActiveChannels; ++ch) + if constexpr (uses_dynamic_smem) + { + _CCCL_ASSERT(dynamic_smem_privatized_histograms != nullptr, "Dynamic-SMEM mode requires a shared-memory base"); + CounterT* channel_histogram = dynamic_smem_privatized_histograms; + _CCCL_PRAGMA_UNROLL_FULL() + for (int ch = 0; ch < NumActiveChannels; ++ch) + { + this->dynamic_smem_privatized_histograms[ch] = channel_histogram; + channel_histogram += num_privatized_bins[ch]; + } + } + else if constexpr (uses_gmem) { - const auto offset = static_cast<::cuda::std::int64_t>(blockId) * num_privatized_bins[ch]; - this->d_privatized_histograms[ch] = d_privatized_histograms[ch] + offset; + const int block_id = static_cast((blockIdx.y * gridDim.x) + blockIdx.x); + for (int ch = 0; ch < NumActiveChannels; ++ch) + { + const auto offset = static_cast<::cuda::std::int64_t>(block_id) * num_privatized_bins[ch]; + this->gmem_privatized_histograms[ch] = gmem_privatized_histograms[ch] + offset; + } } } @@ -697,26 +647,33 @@ struct AgentHistogram //! Initialize privatized bin counters. Specialized for privatized shared-memory counters _CCCL_DEVICE _CCCL_FORCEINLINE void InitBinCounters() { - if (prefer_smem) - { - ZeroBinCounters(temp_storage.histograms); - } - else - { - ZeroBinCounters(d_privatized_histograms); - } + ZeroBinCounters(); } //! Store privatized histogram to device-accessible memory. Specialized for privatized shared-memory counters _CCCL_DEVICE _CCCL_FORCEINLINE void StoreOutput() { - if (prefer_smem) - { - StoreOutput(temp_storage.histograms); - } - else + // Barrier to make sure all threads are done updating counters + __syncthreads(); + + // Apply privatized bin counts to output bin counts + _CCCL_PRAGMA_UNROLL_FULL() + for (int ch = 0; ch < NumActiveChannels; ++ch) { - StoreOutput(d_privatized_histograms); + CounterT* privatized_histogram = PrivatizedHistogram(ch); + const int channel_bins = num_privatized_bins[ch]; + for (int bin = static_cast(threadIdx.x); bin < channel_bins; bin += threads_per_block) + { + int output_bin = -1; + const CounterT count = privatized_histogram[bin]; + const bool is_valid = count > 0; + output_decode_op[ch].template BinSelect(static_cast(bin), output_bin, is_valid); + + if (output_bin >= 0) + { + atomicAdd(&output_histogram[ch][output_bin], static_cast(count)); + } + } } } }; diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index 673be75f6fd7..9d1402cbb77e 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -43,6 +43,7 @@ #include #include #include +#include #include #include #include @@ -53,8 +54,23 @@ CUB_NAMESPACE_BEGIN namespace detail::histogram { -// Maximum number of bins per channel for which we will use a privatized smem strategy -static constexpr int max_privatized_smem_bins = 256; +inline constexpr int byte_sample_privatized_levels = + static_cast(::cuda::std::numeric_limits::max()) + 2; + +template +struct local_counter +{ + using type = OutputCounterT; +}; + +template +struct local_counter> +{ + using type = typename PolicySelector::local_counter_type; +}; + +template +using local_counter_t = typename local_counter::type; template + typename SampleT, + typename OutputCounterT = CounterT> struct DeviceHistogramKernelSource { + static_assert(sizeof(CounterT) <= sizeof(OutputCounterT), + "The output histogram counter must be at least as wide as the local counter"); + using TransformsT = detail::histogram::Transforms; template _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramInitKernel() { - return &DeviceHistogramInitKernel; + return &DeviceHistogramInitKernel; } /// Returns the default histogram sweep kernel that receives pre-initialized decode operators from the host. - template + template _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramSweepKernel() { return &DeviceHistogramSweepKernel< PolicyT, - PRIVATIZED_SMEM_BINS, + PrivatizationMode, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, - OffsetT>; + OffsetT, + OutputCounterT>; } /// Returns the device-init histogram sweep kernel that initializes decode operators from level arrays in the kernel. template ; + IsEven, + OutputCounterT>; } CUB_RUNTIME_FUNCTION static constexpr size_t CounterSize() @@ -159,7 +181,7 @@ struct DeviceHistogramKernelSource template ; + ::cuda::compute_capability cc{}; if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) { @@ -217,7 +241,7 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( { return kernel_source.template HistogramSweepKernelDeviceInit< PolicySelector, - PRIVATIZED_SMEM_BINS, + PrivatizationMode, FirstLevelArrayT, SecondLevelArrayT, IsEven, @@ -228,12 +252,80 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( using output_decode_op_t = typename FirstLevelArrayT::value_type; using privatized_decode_op_t = typename SecondLevelArrayT::value_type; return kernel_source - .template HistogramSweepKernel(); + .template HistogramSweepKernel(); } }(); - const int threads_per_block = active_policy.threads_per_block; - const int pixels_per_thread = active_policy.pixels_per_thread; + const HistogramPrivatizationPolicy sweep = + is_privatized_static_smem_v ? active_policy.static_smem + : is_privatized_dynamic_smem_v + ? active_policy.dynamic_smem + : active_policy.gmem; + const int threads_per_block = sweep.threads_per_block; + const int items_per_thread = sweep.items_per_thread; + + int dynamic_smem_bytes = 0; + if constexpr (is_privatized_dynamic_smem_v) + { + for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) + { + dynamic_smem_bytes += (num_privatized_levels[channel] - 1) * static_cast(kernel_source.CounterSize()); + } + int max_dynamic_smem_bytes{}; + NV_IF_ELSE_TARGET( + NV_IS_HOST, + ({ + if (const auto error = + CubDebug(launcher_factory.max_dynamic_smem_size_for(max_dynamic_smem_bytes, sweep_kernel))) + { + return error; + } + }), + ({ + int max_shared_smem_bytes{}; + if (const auto error = CubDebug(launcher_factory.MaxSharedMemory(max_shared_smem_bytes))) + { + return error; + } + ::cudaFuncAttributes sweep_kernel_attributes{}; + if (const auto error = CubDebug(::cudaFuncGetAttributes(&sweep_kernel_attributes, sweep_kernel))) + { + return error; + } + max_dynamic_smem_bytes = max_shared_smem_bytes - static_cast(sweep_kernel_attributes.sharedSizeBytes); + })) + if (dynamic_smem_bytes > max_dynamic_smem_bytes) + { + return detail::histogram:: + dispatch( + d_temp_storage, + temp_storage_bytes, + d_samples, + d_output_histograms, + num_privatized_levels, + num_output_levels, + first_level_array, + second_level_array, + max_num_output_bins, + num_row_pixels, + num_rows, + row_stride_samples, + stream, + policy_selector, + kernel_source, + launcher_factory); + } + NV_IF_TARGET(NV_IS_HOST, ({ + const int opt_in_dynamic_smem_bytes = + (::cuda::std::min) (dynamic_smem_limit_bytes(active_policy), + max_dynamic_smem_bytes); + if (const auto error = CubDebug( + launcher_factory.set_max_dynamic_smem_size_for(sweep_kernel, opt_in_dynamic_smem_bytes))) + { + return error; + } + })) + } // Get SM count int sm_count; @@ -244,8 +336,8 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( // Get SM occupancy for sweep_kernel int histogram_sweep_sm_occupancy; - if (const auto error = - CubDebug(launcher_factory.MaxSmOccupancy(histogram_sweep_sm_occupancy, sweep_kernel, threads_per_block))) + if (const auto error = CubDebug(launcher_factory.MaxSmOccupancy( + histogram_sweep_sm_occupancy, sweep_kernel, threads_per_block, dynamic_smem_bytes))) { return error; } @@ -262,7 +354,7 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( } // Get grid dimensions, trying to keep total blocks ~histogram_sweep_occupancy - int pixels_per_tile = threads_per_block * pixels_per_thread; + int pixels_per_tile = threads_per_block * items_per_thread; int tiles_per_row = static_cast(::cuda::ceil_div(num_row_pixels, pixels_per_tile)); int blocks_per_row = ::cuda::std::min(histogram_sweep_occupancy, tiles_per_row); int blocks_per_col = @@ -281,10 +373,13 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( void* allocations[NUM_ALLOCATIONS] = {}; size_t allocation_sizes[NUM_ALLOCATIONS]; + constexpr bool requires_global_privatization = is_privatized_gmem_v; for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) { allocation_sizes[CHANNEL] = - size_t(num_thread_blocks) * (num_privatized_levels[CHANNEL] - 1) * kernel_source.CounterSize(); + requires_global_privatization + ? size_t(num_thread_blocks) * (num_privatized_levels[CHANNEL] - 1) * kernel_source.CounterSize() + : 0; } allocation_sizes[NUM_ALLOCATIONS - 1] = GridQueue::AllocationSize(); @@ -307,11 +402,11 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( GridQueue tile_queue(allocations[NUM_ALLOCATIONS - 1]); // Wrap arrays so we can pass them by-value to the kernel - ::cuda::std::array d_privatized_histograms_wrapper; + ::cuda::std::array d_privatized_histograms_wrapper; ::cuda::std::array num_privatized_bins_wrapper; ::cuda::std::array num_output_bins_wrapper; - auto* typed_allocations = reinterpret_cast(allocations); + auto* typed_allocations = reinterpret_cast(allocations); ::cuda::std::copy(typed_allocations, typed_allocations + NUM_ACTIVE_CHANNELS, d_privatized_histograms_wrapper.begin()); auto minus_one = ::cuda::proclaim_return_type([](int levels) { @@ -321,22 +416,21 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( num_privatized_levels.begin(), num_privatized_levels.end(), num_privatized_bins_wrapper.begin(), minus_one); ::cuda::std::transform(num_output_levels.begin(), num_output_levels.end(), num_output_bins_wrapper.begin(), minus_one); - constexpr int histogram_init_threads_per_block = 256; - int histogram_init_grid_dims = - (max_num_output_bins + histogram_init_threads_per_block - 1) / histogram_init_threads_per_block; + const int histogram_init_grid_dims = + (max_num_output_bins + active_policy.init_threads_per_block - 1) / active_policy.init_threads_per_block; // Log DeviceHistogramInitKernel configuration #ifdef CUB_DEBUG_LOG _CubLog("Invoking DeviceHistogramInitKernel<<<%d, %d, 0, %lld>>>()\n", histogram_init_grid_dims, - histogram_init_threads_per_block, + active_policy.init_threads_per_block, (long long) stream); #endif // CUB_DEBUG_LOG // Invoke histogram_init_kernel if (const auto error = CubDebug( launcher_factory(histogram_init_grid_dims, - histogram_init_threads_per_block, + active_policy.init_threads_per_block, 0, stream, /* dependent_launch */ cc >= ::cuda::compute_capability{9, 0}) @@ -353,20 +447,24 @@ CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE auto dispatch( // Log histogram_sweep_kernel configuration #ifdef CUB_DEBUG_LOG - _CubLog("Invoking histogram_sweep_kernel<<<{%d, %d, %d}, %d, 0, %lld>>>(), %d pixels " + _CubLog("Invoking histogram_sweep_kernel<<<{%d, %d, %d}, %d, %d, %lld>>>(), %d pixels " "per thread, %d SM occupancy\n", sweep_grid_dims.x, sweep_grid_dims.y, sweep_grid_dims.z, threads_per_block, + dynamic_smem_bytes, (long long) stream, - pixels_per_thread, + items_per_thread, histogram_sweep_sm_occupancy); #endif // CUB_DEBUG_LOG if (const auto error = CubDebug( - launcher_factory( - sweep_grid_dims, threads_per_block, 0, stream, /* dependent_launch */ cc >= ::cuda::compute_capability{9, 0}) + launcher_factory(sweep_grid_dims, + threads_per_block, + dynamic_smem_bytes, + stream, + /* dependent_launch */ cc >= ::cuda::compute_capability{9, 0}) .doit(sweep_kernel, d_samples, num_output_bins_wrapper, @@ -465,7 +563,8 @@ template < typename LevelT, typename OffsetT, typename PolicySelector, - typename SampleT = it_value_t, /// The sample value type of the input iterator + typename PrivatizedCounterT = CounterT, + typename SampleT = it_value_t, /// The sample value type of the input iterator typename KernelSource = DeviceHistogramKernelSource, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY, @@ -512,15 +611,61 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t __dispatch_even_device } int max_num_output_bins = max_levels - 1; - if (max_num_output_bins > detail::histogram::max_privatized_smem_bins) + ::cuda::compute_capability cc{}; + if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) { - // Dispatch shared-privatized approach - constexpr int PRIVATIZED_SMEM_BINS = 0; + return error; + } + const HistogramPolicy active_policy = policy_selector(cc); + const auto privatization = [&] { + if constexpr (::cuda::std::is_void_v) + { + return select_privatization_mode_for_counter_size( + active_policy, max_num_output_bins, static_cast(kernel_source.CounterSize())); + } + else + { + return select_privatization_mode( + active_policy, max_num_output_bins); + } + }(); + if (privatization == privatization_mode::dynamic_smem) + { + if (const auto error = CubDebug( + (detail::histogram::dispatch( + d_temp_storage, + temp_storage_bytes, + d_samples, + d_output_histograms, + num_output_levels, + num_output_levels, + upper_level, + lower_level, + max_num_output_bins, + num_row_pixels, + num_rows, + row_stride_samples, + stream, + policy_selector, + kernel_source, + launcher_factory)))) + { + return error; + } + } + else if (privatization == privatization_mode::gmem) + { + // Dispatch global-memory-privatized approach if (const auto error = CubDebug( (detail::histogram::dispatch( @@ -547,12 +692,10 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t __dispatch_even_device else { // Dispatch shared-privatized approach - constexpr int PRIVATIZED_SMEM_BINS = detail::histogram::max_privatized_smem_bins; - if (const auto error = CubDebug( (detail::histogram::dispatch( @@ -637,7 +780,8 @@ template < typename LevelT, typename OffsetT, typename PolicySelector, - typename SampleT = it_value_t, /// The sample value type of the input iterator + typename PrivatizedCounterT = CounterT, + typename SampleT = it_value_t, /// The sample value type of the input iterator typename KernelSource = DeviceHistogramKernelSource, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY, @@ -665,7 +809,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t __dispatch_even_device for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) { - num_privatized_levels[channel] = 257; + num_privatized_levels[channel] = byte_sample_privatized_levels; int num_levels = num_output_levels[channel]; if (kernel_source.MayOverflow(num_levels - 1, upper_level, lower_level, channel)) @@ -687,31 +831,58 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t __dispatch_even_device } int max_num_output_bins = max_levels - 1; - constexpr int PRIVATIZED_SMEM_BINS = 256; + ::cuda::compute_capability cc{}; + if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) + { + return error; + } + const HistogramPolicy active_policy = policy_selector(cc); + constexpr int num_privatized_bins = byte_sample_privatized_levels - 1; + const auto privatization = [&] { + if constexpr (::cuda::std::is_void_v) + { + return select_privatization_mode_for_counter_size( + active_policy, num_privatized_bins, static_cast(kernel_source.CounterSize())); + } + else + { + return select_privatization_mode( + active_policy, num_privatized_bins); + } + }(); - if (const auto error = CubDebug( - (detail::histogram::dispatch( - d_temp_storage, - temp_storage_bytes, - d_samples, - d_output_histograms, - num_privatized_levels, - num_output_levels, - upper_level, - lower_level, - max_num_output_bins, - num_row_pixels, - num_rows, - row_stride_samples, - stream, - policy_selector, - kernel_source, - launcher_factory)))) + const auto dispatch_with = [&](auto mode) { + using privatization_mode_t = decltype(mode); + return detail::histogram::dispatch( + d_temp_storage, + temp_storage_bytes, + d_samples, + d_output_histograms, + num_privatized_levels, + num_output_levels, + upper_level, + lower_level, + max_num_output_bins, + num_row_pixels, + num_rows, + row_stride_samples, + stream, + policy_selector, + kernel_source, + launcher_factory); + }; + + const auto error = + privatization == privatization_mode::static_smem ? dispatch_with(HistogramPrivatizedStaticSmem{}) + : privatization == privatization_mode::dynamic_smem + ? dispatch_with(HistogramPrivatizedDynamicSmem{}) + : dispatch_with(HistogramPrivatizedGmem{}); + if (CubDebug(error)) { return error; } @@ -720,50 +891,62 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t __dispatch_even_device } // TODO(bgruber): drop in CCCL 4.0 -template -_CCCL_HOST_DEVICE_API constexpr auto convert_pdl_trigger(int) +template +_CCCL_HOST_DEVICE_API constexpr auto convert_pdl_trigger_bytes(int) -> decltype(ActivePolicy::init_kernel_pdl_trigger_max_bins) { - return ActivePolicy::init_kernel_pdl_trigger_max_bins; + return ActivePolicy::init_kernel_pdl_trigger_max_bins * int{sizeof(CounterT)}; } // TODO(bgruber): drop in CCCL 4.0 -template -_CCCL_HOST_DEVICE_API constexpr auto convert_pdl_trigger(long) +template +_CCCL_HOST_DEVICE_API constexpr auto convert_pdl_trigger_bytes(long) { return 0; } // TODO(bgruber): drop in CCCL 4.0 -template -_CCCL_HOST_DEVICE_API constexpr auto convert_policy() -> HistogramPolicy +template +_CCCL_HOST_DEVICE_API constexpr auto convert_legacy_policy() -> HistogramPolicy { - using ap = typename ActivePolicy::AgentHistogramPolicyT; - return HistogramPolicy{ - ap::BLOCK_THREADS, - ap::PIXELS_PER_THREAD, - ap::VEC_SIZE, - ap::LOAD_ALGORITHM, - ap::LOAD_MODIFIER, - ap::IS_RLE_COMPRESS, - ap::MEM_PREFERENCE, - ap::IS_WORK_STEALING, - convert_pdl_trigger(0)}; + using sweep = typename ActivePolicy::AgentHistogramPolicyT; + const auto kernel_config = HistogramPrivatizationPolicy{ + sweep::BLOCK_THREADS, + sweep::PIXELS_PER_THREAD, + sweep::VEC_SIZE, + sweep::LOAD_ALGORITHM, + sweep::LOAD_MODIFIER, + sweep::IS_RLE_COMPRESS, + sweep::IS_WORK_STEALING}; + return { + kernel_config, + kernel_config, + kernel_config, + 256, + 256 * int{sizeof(CounterT)}, + 0, + 0, + 0, + 0, + 0, + 0, + 0, + convert_pdl_trigger_bytes(0)}; } // TODO(bgruber): drop in CCCL 4.0 -template -struct policy_selector_from_max_policy +template +struct policy_selector_from_hub { private: - struct extract_policy_dispatch_t + struct dispatch_t { HistogramPolicy& policy; - template + template _CCCL_HOST_DEVICE_API constexpr cudaError_t Invoke() { - policy = convert_policy(); + policy = convert_legacy_policy(); return cudaSuccess; } }; @@ -774,27 +957,45 @@ public: NV_IF_ELSE_TARGET(NV_IS_HOST, ({ HistogramPolicy policy{}; - extract_policy_dispatch_t dispatch{policy}; + dispatch_t dispatch{policy}; _CCCL_VERIFY(MaxPolicy::Invoke(cc.get() * 10, dispatch) == cudaSuccess, ""); return policy; }), - ({ return convert_policy(); })); + ({ return convert_legacy_policy(); })); } }; -template < - int NUM_CHANNELS, - int NUM_ACTIVE_CHANNELS, - typename SampleIteratorT, - typename CounterT, - typename LevelT, - typename OffsetT, - bool IsByteSample, - typename PolicySelector, - typename SampleT = it_value_t, /// The sample value type of the input iterator - typename KernelSource = - DeviceHistogramKernelSource, - typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> +template +struct max_policy_from_hub +{ + using type = typename PolicyHub::MaxPolicy; +}; + +template <> +struct max_policy_from_hub +{ + using type = void; +}; + +template , /// The sample value type of the input iterator + typename KernelSource = DeviceHistogramKernelSource< + NUM_CHANNELS, + NUM_ACTIVE_CHANNELS, + SampleIteratorT, + local_counter_t, + LevelT, + OffsetT, + SampleT, + CounterT>, + typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( void* d_temp_storage, size_t& temp_storage_bytes, @@ -811,6 +1012,8 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( KernelSource kernel_source = {}, KernelLauncherFactory launcher_factory = {}) { + using LocalCounterT = local_counter_t; + if constexpr (IsByteSample) { using TransformsT = Transforms; @@ -828,7 +1031,7 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) { - num_privatized_levels[channel] = 257; + num_privatized_levels[channel] = byte_sample_privatized_levels; output_decode_op[channel].Init(d_levels[channel], num_output_levels[channel]); if (num_output_levels[channel] > max_levels) @@ -838,20 +1041,108 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( } int max_num_output_bins = max_levels - 1; - constexpr int PRIVATIZED_SMEM_BINS = 256; + ::cuda::compute_capability cc{}; + if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) + { + return error; + } + const HistogramPolicy active_policy = policy_selector(cc); + constexpr int num_privatized_bins = byte_sample_privatized_levels - 1; + const auto privatization = + select_privatization_mode(active_policy, num_privatized_bins); + + const auto dispatch_with = [&](auto mode) { + using privatization_mode_t = decltype(mode); + return detail::histogram::dispatch( + d_temp_storage, + temp_storage_bytes, + d_samples, + d_output_histograms, + num_privatized_levels, + num_output_levels, + output_decode_op, + privatized_decode_op, + max_num_output_bins, + num_row_pixels, + num_rows, + row_stride_samples, + stream, + policy_selector, + kernel_source, + launcher_factory); + }; + + const auto error = + privatization == privatization_mode::static_smem ? dispatch_with(HistogramPrivatizedStaticSmem{}) + : privatization == privatization_mode::dynamic_smem + ? dispatch_with(HistogramPrivatizedDynamicSmem{}) + : dispatch_with(HistogramPrivatizedGmem{}); + if (CubDebug(error)) + { + return error; + } + } + else + { + ::cuda::compute_capability cc{}; + if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) + { + return error; + } + const HistogramPolicy active_policy = policy_selector(cc); - if (const auto error = CubDebug( - (detail::histogram::dispatch( + using TransformsT = Transforms; + + // Use the pass-thru transform op for converting privatized bins to output bins + using OutputDecodeOpT = typename TransformsT::PassThruTransform; + + ::cuda::std::array output_decode_op{}; + int max_levels = num_output_levels[0]; + + for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) + { + if (num_output_levels[channel] > max_levels) + { + max_levels = num_output_levels[channel]; + } + } + int max_num_output_bins = max_levels - 1; + const auto privatization = + select_privatization_mode(active_policy, max_num_output_bins); + + constexpr bool supports_cached_search = + ::cuda::std::is_integral_v || ::cuda::std::is_floating_point_v; + if constexpr (supports_cached_search) + { + if (privatization == privatization_mode::dynamic_smem + || (privatization == privatization_mode::gmem + && use_cached_search_for_gmem_range(active_policy, max_num_output_bins))) + { + using PrivatizedDecodeOpT = typename TransformsT::template CachedSearchTransform; + ::cuda::std::array privatized_decode_op{}; + for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) + { + privatized_decode_op[channel].Init(d_levels[channel], num_output_levels[channel]); + } + + const auto dispatch_with = [&](auto mode) { + using privatization_mode_t = decltype(mode); + return detail::histogram::dispatch( d_temp_storage, temp_storage_bytes, d_samples, d_output_histograms, - num_privatized_levels, + num_output_levels, num_output_levels, output_decode_op, privatized_decode_op, @@ -862,45 +1153,30 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( stream, policy_selector, kernel_source, - launcher_factory)))) - { - return error; + launcher_factory); + }; + + return CubDebug(privatization == privatization_mode::dynamic_smem + ? dispatch_with(HistogramPrivatizedDynamicSmem{}) + : dispatch_with(HistogramPrivatizedGmem{})); + } } - } - else - { - using TransformsT = Transforms; - // Use the search transform op for converting samples to privatized bins using PrivatizedDecodeOpT = typename TransformsT::template SearchTransform; - - // Use the pass-thru transform op for converting privatized bins to output bins - using OutputDecodeOpT = typename TransformsT::PassThruTransform; - ::cuda::std::array privatized_decode_op{}; - ::cuda::std::array output_decode_op{}; - int max_levels = num_output_levels[0]; - for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) { privatized_decode_op[channel].Init(d_levels[channel], num_output_levels[channel]); - if (num_output_levels[channel] > max_levels) - { - max_levels = num_output_levels[channel]; - } } - int max_num_output_bins = max_levels - 1; // Dispatch - if (max_num_output_bins > max_privatized_smem_bins) + if (privatization != privatization_mode::static_smem) { // Too many bins to keep in shared memory. - constexpr int PRIVATIZED_SMEM_BINS = 0; - if (const auto error = CubDebug( (detail::histogram::dispatch( @@ -927,12 +1203,10 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( else { // Dispatch shared-privatized approach - constexpr int PRIVATIZED_SMEM_BINS = max_privatized_smem_bins; - if (const auto error = CubDebug( (detail::histogram::dispatch( @@ -961,19 +1235,25 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_range( return cudaSuccess; } -template < - int NUM_CHANNELS, - int NUM_ACTIVE_CHANNELS, - typename SampleIteratorT, - typename CounterT, - typename LevelT, - typename OffsetT, - bool IsByteSample, - typename PolicySelector, - typename SampleT = it_value_t, /// The sample value type of the input iterator - typename KernelSource = - DeviceHistogramKernelSource, - typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> +template , /// The sample value type of the input iterator + typename KernelSource = DeviceHistogramKernelSource< + NUM_CHANNELS, + NUM_ACTIVE_CHANNELS, + SampleIteratorT, + local_counter_t, + LevelT, + OffsetT, + SampleT, + CounterT>, + typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even( void* d_temp_storage, size_t& temp_storage_bytes, @@ -991,6 +1271,8 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even( KernelSource kernel_source = {}, KernelLauncherFactory launcher_factory = {}) { + using LocalCounterT = local_counter_t; + if constexpr (IsByteSample) { using TransformsT = Transforms; @@ -1010,7 +1292,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even( for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) { - num_privatized_levels[channel] = 257; + num_privatized_levels[channel] = byte_sample_privatized_levels; int num_levels = num_output_levels[channel]; if (kernel_source.MayOverflow(static_cast(num_levels - 1), upper_level, lower_level, channel)) @@ -1031,31 +1313,48 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even( } int max_num_output_bins = max_levels - 1; - constexpr int PRIVATIZED_SMEM_BINS = 256; - - if (const auto error = CubDebug( - (detail::histogram::dispatch( - d_temp_storage, - temp_storage_bytes, - d_samples, - d_output_histograms, - num_privatized_levels, - num_output_levels, - output_decode_op, - privatized_decode_op, - max_num_output_bins, - num_row_pixels, - num_rows, - row_stride_samples, - stream, - policy_selector, - kernel_source, - launcher_factory)))) + ::cuda::compute_capability cc{}; + if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) + { + return error; + } + const HistogramPolicy active_policy = policy_selector(cc); + constexpr int num_privatized_bins = byte_sample_privatized_levels - 1; + const auto privatization = + select_privatization_mode(active_policy, num_privatized_bins); + + const auto dispatch_with = [&](auto mode) { + using privatization_mode_t = decltype(mode); + return detail::histogram::dispatch( + d_temp_storage, + temp_storage_bytes, + d_samples, + d_output_histograms, + num_privatized_levels, + num_output_levels, + output_decode_op, + privatized_decode_op, + max_num_output_bins, + num_row_pixels, + num_rows, + row_stride_samples, + stream, + policy_selector, + kernel_source, + launcher_factory); + }; + + const auto error = + privatization == privatization_mode::static_smem ? dispatch_with(HistogramPrivatizedStaticSmem{}) + : privatization == privatization_mode::dynamic_smem + ? dispatch_with(HistogramPrivatizedDynamicSmem{}) + : dispatch_with(HistogramPrivatizedGmem{}); + if (CubDebug(error)) { return error; } @@ -1097,14 +1396,47 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even( } int max_num_output_bins = max_levels - 1; - if (max_num_output_bins > max_privatized_smem_bins) + ::cuda::compute_capability cc{}; + if (const auto error = CubDebug(launcher_factory.PtxComputeCap(cc))) { - constexpr int PRIVATIZED_SMEM_BINS = 0; + return error; + } + const HistogramPolicy active_policy = policy_selector(cc); + + const auto privatization = + select_privatization_mode(active_policy, max_num_output_bins); + if (privatization == privatization_mode::dynamic_smem) + { + return CubDebug((detail::histogram::dispatch( + d_temp_storage, + temp_storage_bytes, + d_samples, + d_output_histograms, + num_output_levels, + num_output_levels, + output_decode_op, + privatized_decode_op, + max_num_output_bins, + num_row_pixels, + num_rows, + row_stride_samples, + stream, + policy_selector, + kernel_source, + launcher_factory))); + } + if (privatization == privatization_mode::gmem) + { if (const auto error = CubDebug( (detail::histogram::dispatch( @@ -1130,12 +1462,10 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even( } else { - constexpr int PRIVATIZED_SMEM_BINS = max_privatized_smem_bins; - if (const auto error = CubDebug( (detail::histogram::dispatch( @@ -1269,12 +1599,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") Dispatc * @param stream * CUDA stream to launch kernels within. Default is stream0. */ - template , - /* fallback_policy_hub */ - detail::histogram::policy_hub, - PolicyHub>::MaxPolicy, - bool IsByteSample> + template ::type, bool IsByteSample> CUB_RUNTIME_FUNCTION static cudaError_t DispatchRange( void* d_temp_storage, size_t& temp_storage_bytes, @@ -1289,8 +1614,14 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") Dispatc ::cuda::std::bool_constant is_byte_sample, KernelSource kernel_source = {}, KernelLauncherFactory launcher_factory = {}, - [[maybe_unused]] MaxPolicyT max_policy = {}) + [[maybe_unused]] ::cuda::std::_If<::cuda::std::is_void_v, ::cuda::std::nullptr_t, MaxPolicyT> + max_policy = {}) { + static constexpr bool uses_default_policy = ::cuda::std::is_void_v; + using policy_selector_t = ::cuda::std::_If< + uses_default_policy, + detail::histogram::policy_selector_from_types, + detail::histogram::policy_selector_from_hub>; return detail::histogram::dispatch_range( d_temp_storage, temp_storage_bytes, @@ -1303,7 +1634,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") Dispatc row_stride_samples, stream, is_byte_sample, - detail::histogram::policy_selector_from_max_policy{}, + policy_selector_t{}, kernel_source, launcher_factory); } @@ -1355,12 +1686,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") Dispatc * CUDA stream to launch kernels within. Default is stream0. * */ - template , - /* fallback_policy_hub */ - detail::histogram::policy_hub, - PolicyHub>::MaxPolicy, - bool IsByteSample> + template ::type, bool IsByteSample> CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t DispatchEven( void* d_temp_storage, size_t& temp_storage_bytes, @@ -1376,8 +1702,14 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") Dispatc ::cuda::std::bool_constant is_byte_sample, KernelSource kernel_source = {}, KernelLauncherFactory launcher_factory = {}, - [[maybe_unused]] MaxPolicyT max_policy = {}) + [[maybe_unused]] ::cuda::std::_If<::cuda::std::is_void_v, ::cuda::std::nullptr_t, MaxPolicyT> + max_policy = {}) { + static constexpr bool uses_default_policy = ::cuda::std::is_void_v; + using policy_selector_t = ::cuda::std::_If< + uses_default_policy, + detail::histogram::policy_selector_from_types, + detail::histogram::policy_selector_from_hub>; return detail::histogram::dispatch_even( d_temp_storage, temp_storage_bytes, @@ -1391,7 +1723,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") Dispatc row_stride_samples, stream, is_byte_sample, - detail::histogram::policy_selector_from_max_policy{}, + policy_selector_t{}, kernel_source, launcher_factory); } diff --git a/cub/cub/device/dispatch/kernels/kernel_histogram.cuh b/cub/cub/device/dispatch/kernels/kernel_histogram.cuh index cc64ddea1862..78b833598ee9 100644 --- a/cub/cub/device/dispatch/kernels/kernel_histogram.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_histogram.cuh @@ -20,24 +20,106 @@ #include #include +#include +#include CUB_NAMESPACE_BEGIN namespace detail::histogram { -template +template struct Transforms { //--------------------------------------------------------------------- // Transform functors for converting samples to bin-ids //--------------------------------------------------------------------- - // Searches for bin given a list of bin-boundary levels + //! @brief Finds a RANGE bin with binary search. + //! + //! Uses `UpperBound` without interpolation or per-thread state. template struct SearchTransform { LevelIteratorT d_levels; // Pointer to levels array int num_output_levels; // Number of levels in array + _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void Init(LevelIteratorT d_levels_, int num_output_levels_) + { + d_levels = d_levels_; + num_output_levels = num_output_levels_; + } + + _CCCL_DEVICE _CCCL_FORCEINLINE void Precompute() {} + + template + _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(SampleT sample, int& bin, bool valid) const + { + using WrappedLevelIteratorT = + ::cuda::std::_If<::cuda::std::is_pointer_v, + CacheModifiedInputIterator, + LevelIteratorT>; + WrappedLevelIteratorT wrapped_levels(d_levels); + + const int num_bins = num_output_levels - 1; + if (valid) + { + bin = UpperBound(wrapped_levels, num_output_levels, static_cast(sample)) - 1; + if (bin >= num_bins) + { + bin = -1; + } + } + } + }; + + //! @brief Finds a RANGE bin with piecewise-linear interpolation and a per-thread bracket cache. + //! + //! This transform is used by the runtime-sized shared-memory kernel and by + //! selected global-memory kernels. It precomputes interpolation parameters + //! once per thread, verifies each interpolated guess against the level array, + //! and falls back to binary search for irregular levels. `BinSelectState` + //! remembers the most recently resolved bracket so consecutive samples in + //! that bracket require no level loads. + template + struct CachedSearchTransform + { + //! @brief Computes a non-negative interpolation distance without signed overflow. + template + [[nodiscard]] _CCCL_HOST_DEVICE_API static constexpr auto interpolation_difference(T lhs, T rhs) + { + if constexpr (::cuda::std::is_integral_v) + { + using unsigned_t = ::cuda::std::make_unsigned_t; + return static_cast(lhs) - static_cast(rhs); + } + else + { + return lhs - rhs; + } + } + + struct BinSelectState + { + LevelT lo; // cached d_levels[bin] + LevelT hi; // cached d_levels[bin + 1] + int bin = -1; // cached bin; < 0 means empty + }; + + BinSelectState most_recent_bin; + + LevelIteratorT d_levels; // Pointer to levels array + int num_output_levels; // Number of levels in array + // Interpolation state shared by all samples processed by a thread. + float inv_scale; // num_bins / (float)(last - first); valid iff have_precompute + LevelT first; // cached d_levels[0] + LevelT last; // cached d_levels[num_bins] + bool have_precompute; // whether the fields above are valid + + // Piecewise-linear interpolation state split at the midpoint level. + LevelT mid; // cached d_levels[mid_bin] + float inv_scale_lo; // mid_bin / (float)(mid - first) + float inv_scale_hi; // (num_bins - mid_bin) / (float)(last - mid) + int mid_bin; // split bin index (num_bins / 2) + //! @brief Initializer //! //! @param d_levels_ Pointer to levels array @@ -46,30 +128,190 @@ struct Transforms { this->d_levels = d_levels_; this->num_output_levels = num_output_levels_; + this->have_precompute = false; + this->inv_scale = 0.0f; + this->first = LevelT{}; + this->last = LevelT{}; + this->mid = LevelT{}; + this->inv_scale_lo = 0.0f; + this->inv_scale_hi = 0.0f; + this->mid_bin = 0; + this->most_recent_bin = BinSelectState{}; } - // Method for converting samples to bin-ids - template - _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(_SampleT sample, int& bin, bool valid) const + //! @brief Precomputes interpolation slopes from the device level array. + _CCCL_DEVICE _CCCL_FORCEINLINE void Precompute() { - /// Level iterator wrapper type - // Wrap the native input pointer with CacheModifiedInputIterator - // or Directly use the supplied input iterator type + const int num_bins = num_output_levels - 1; using WrappedLevelIteratorT = ::cuda::std::_If<::cuda::std::is_pointer_v, - CacheModifiedInputIterator, + CacheModifiedInputIterator, LevelIteratorT>; - WrappedLevelIteratorT wrapped_levels(d_levels); + const LevelT first = wrapped_levels[0]; + const LevelT last = wrapped_levels[num_bins]; + if (!(first < last)) + { + have_precompute = false; + return; + } + + this->first = first; + this->last = last; + inv_scale = static_cast(num_bins) / static_cast(interpolation_difference(last, first)); + have_precompute = true; + + // Use a single secant if the midpoint does not split the level range. + this->mid_bin = 0; + inv_scale_lo = inv_scale; + inv_scale_hi = inv_scale; + mid = first; + const int split = num_bins >> 1; + if (split > 0 && split < num_bins) + { + const LevelT split_level = wrapped_levels[split]; + if ((first < split_level) && (split_level < last)) + { + mid = split_level; + mid_bin = split; + inv_scale_lo = static_cast(split) / static_cast(interpolation_difference(split_level, first)); + inv_scale_hi = + static_cast(num_bins - split) / static_cast(interpolation_difference(last, split_level)); + } + } + } + + //! @brief Implements cached/interpolated bin selection. + //! + //! A cached-bracket hit returns immediately. Otherwise, this computes and + //! verifies an interpolated guess, checks one adjacent bracket, and finally + //! falls back to `UpperBound` for arbitrary level distributions. + template + _CCCL_DEVICE _CCCL_FORCEINLINE void BinSelect(SampleT sample, int& bin, bool valid) + { + using WrappedLevelIteratorT = + ::cuda::std::_If<::cuda::std::is_pointer_v, + CacheModifiedInputIterator, + LevelIteratorT>; + WrappedLevelIteratorT wrapped_levels(d_levels); const int num_bins = num_output_levels - 1; - if (valid) + if (!valid) { - bin = UpperBound(wrapped_levels, num_output_levels, static_cast(sample)) - 1; + return; + } + + const LevelT s = static_cast(sample); + + if (most_recent_bin.bin >= 0 && !(s < most_recent_bin.lo) && (s < most_recent_bin.hi)) + { + bin = most_recent_bin.bin; + return; + } + + const LevelT first_level = have_precompute ? first : wrapped_levels[0]; + const LevelT last_level = have_precompute ? last : wrapped_levels[num_bins]; + + if (!(first_level < last_level)) + { + bin = UpperBound(wrapped_levels, num_output_levels, s) - 1; if (bin >= num_bins) { bin = -1; } + return; + } + + if (s < first_level || !(s < last_level)) + { + bin = -1; + return; + } + + const auto delta = interpolation_difference(s, first_level); + int guess; + if (have_precompute) + { + if (mid_bin > 0) + { + if (s < mid) + { + guess = static_cast(static_cast(delta) * inv_scale_lo); + } + else + { + const auto delta_hi = interpolation_difference(s, mid); + guess = mid_bin + static_cast(static_cast(delta_hi) * inv_scale_hi); + } + } + else + { + guess = static_cast(static_cast(delta) * inv_scale); + } + } + else + { + const auto range = interpolation_difference(last_level, first_level); + guess = + static_cast((static_cast(delta) * static_cast(num_bins)) / static_cast(range)); + } + if (guess < 0) + { + guess = 0; + } + else if (guess > num_bins - 1) + { + guess = num_bins - 1; + } + + const LevelT lvl_lo = wrapped_levels[guess]; + const LevelT lvl_hi = wrapped_levels[guess + 1]; + + if (!(s < lvl_lo) && (s < lvl_hi)) + { + bin = guess; + most_recent_bin = BinSelectState{lvl_lo, lvl_hi, guess}; + return; + } + + if (s < lvl_lo) + { + const int g2 = guess - 1; + if (g2 >= 0) + { + const LevelT lvl2_lo = wrapped_levels[g2]; + if (!(s < lvl2_lo)) + { + bin = g2; + most_recent_bin = BinSelectState{lvl2_lo, lvl_lo, g2}; + return; + } + } + } + else + { + const int g2 = guess + 1; + if (g2 <= num_bins - 1) + { + const LevelT lvl2_hi = wrapped_levels[g2 + 1]; + if (s < lvl2_hi) + { + bin = g2; + most_recent_bin = BinSelectState{lvl_hi, lvl2_hi, g2}; + return; + } + } + } + + bin = UpperBound(wrapped_levels, num_output_levels, s) - 1; + if (bin >= num_bins) + { + bin = -1; + return; + } + if (bin >= 0) + { + most_recent_bin = BinSelectState{wrapped_levels[bin], wrapped_levels[bin + 1], bin}; } } }; @@ -77,12 +319,12 @@ struct Transforms // Scales samples to evenly-spaced bins struct ScaleTransform { - using CommonT = ::cuda::std::common_type_t; + using CommonT = ::cuda::std::common_type_t; static_assert(::cuda::std::is_convertible_v, - "The common type of `LevelT` and `SampleT` must be " + "The common type of `LevelT` and `InputSampleT` must be " "convertible to `int`."); static_assert(::cuda::is_trivially_copyable_v, - "The common type of `LevelT` and `SampleT` must be " + "The common type of `LevelT` and `InputSampleT` must be " "trivially copyable."); // An arithmetic type that's used for bin computation of integral types, guaranteed to not @@ -92,7 +334,7 @@ struct Transforms // multiplication result. // If CommonT used to be a 128-bit wide integral type already, we use CommonT's arithmetic using IntArithmeticT = ::cuda::std::_If< // - sizeof(SampleT) + sizeof(CommonT) <= sizeof(uint32_t), // + sizeof(InputSampleT) + sizeof(CommonT) <= sizeof(uint32_t), // uint32_t, // #if _CCCL_HAS_INT128() ::cuda::std::_If< // @@ -265,9 +507,11 @@ struct Transforms m_scale = this->ComputeScale(num_levels, m_max, m_min); } + _CCCL_DEVICE _CCCL_FORCEINLINE void Precompute() {} + // Method for converting samples to bin-ids template - _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(SampleT sample, int& bin, bool valid) const + _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(InputSampleT sample, int& bin, bool valid) const { const CommonT common_sample = static_cast(sample); @@ -297,9 +541,11 @@ struct Transforms _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void Init(T, int) {} + _CCCL_DEVICE _CCCL_FORCEINLINE void Precompute() {} + // Method for converting samples to bin-ids - template - _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(_SampleT sample, int& bin, bool valid) const + template + _CCCL_HOST_DEVICE _CCCL_FORCEINLINE void BinSelect(SampleT sample, int& bin, bool valid) const { if (valid) { @@ -348,14 +594,18 @@ _CCCL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( _CCCL_PDL_GRID_DEPENDENCY_SYNC(); // TODO(bgruber): if we had the guarantee that there would be no pending // writes/reads to the temp storage, we could omit the sync here - // we trigger the sweep kernel only if we have a small number of remaining writes in this kernel - NV_IF_TARGET(NV_PROVIDES_SM_90, ({ - if (::cuda::std::reduce(num_output_bins_wrapper.begin(), num_output_bins_wrapper.end()) - <= policy.init_kernel_pdl_trigger_max_bins) - { - _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); - } - })); + // Trigger the sweep only when the remaining output writes occupy at most the tuned byte threshold. + NV_IF_TARGET( + NV_PROVIDES_SM_90, ({ + const auto output_histogram_bytes = + ::cuda::std::reduce(num_output_bins_wrapper.begin(), num_output_bins_wrapper.end(), ::cuda::std::uint64_t{0}) + * sizeof(CounterT); + if (output_histogram_bytes + <= static_cast<::cuda::std::uint64_t>(policy.max_output_histogram_bytes_for_init_kernel_pdl)) + { + _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); + } + })); if ((threadIdx.x == 0) && (blockIdx.x == 0)) { @@ -374,6 +624,36 @@ _CCCL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( } } +template +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto histogram_privatization_policy() -> HistogramPrivatizationPolicy +{ + if constexpr (is_privatized_static_smem_v) + { + return current_policy().static_smem; + } + else if constexpr (is_privatized_dynamic_smem_v) + { + return current_policy().dynamic_smem; + } + else + { + return current_policy().gmem; + } +} + +template +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr int histogram_min_blocks_per_sm() +{ + if constexpr (is_privatized_static_smem_v) + { + return current_policy().static_smem_min_blocks_per_sm; + } + else + { + return 0; + } +} + //! Histogram privatized sweep kernel entry point (multi-block). //! Computes privatized histograms, one per thread block. //! This kernel receives pre-initialized decode operators from the host. @@ -381,8 +661,8 @@ _CCCL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( //! @tparam PolicySelector //! Selects the tuning policy //! -//! @tparam PrivatizedSmemBins -//! Maximum number of histogram bins per channel (e.g., up to 256) +//! @tparam PrivatizationMode +//! Storage mode for the privatized histogram //! //! @tparam NumChannels //! Number of channels interleaved in the input data (may be greater than the number of channels @@ -446,67 +726,72 @@ _CCCL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( //! @param tile_queue //! Drain queue descriptor for dynamically mapping tile data onto thread blocks template + typename OffsetT, + typename OutputCounterT = CounterT> #if _CCCL_HAS_CONCEPTS() requires histogram_policy_selector #endif // _CCCL_HAS_CONCEPTS() -__launch_bounds__(int(current_policy().threads_per_block)) +__launch_bounds__(int(histogram_privatization_policy().threads_per_block), + int(histogram_min_blocks_per_sm())) _CCCL_KERNEL_ATTRIBUTES void DeviceHistogramSweepKernel( const SampleIteratorT d_samples, const ::cuda::std::array num_output_bins_wrapper, const ::cuda::std::array num_privatized_bins_wrapper, - ::cuda::std::array d_output_histograms_wrapper, + ::cuda::std::array d_output_histograms_wrapper, ::cuda::std::array d_privatized_histograms_wrapper, - const ::cuda::std::array output_decode_op_wrapper, - const ::cuda::std::array privatized_decode_op_wrapper, + ::cuda::std::array output_decode_op_wrapper, + ::cuda::std::array privatized_decode_op_wrapper, const OffsetT num_row_pixels, const OffsetT num_rows, const OffsetT row_stride_samples, const int tiles_per_row, GridQueue tile_queue) { - static constexpr HistogramPolicy hp = current_policy(); - - // Thread block type for compositing input tiles - using AgentHistogramPolicyT = agent_histogram_policy< - hp.threads_per_block, - hp.pixels_per_thread, - hp.load_algorithm, - hp.load_modifier, - hp.rle_compress, - hp.mem_preference, - hp.use_work_stealing, - hp.vec_size>; using AgentHistogramT = - AgentHistogram; + OffsetT, + OutputCounterT>; + + __shared__ typename AgentHistogramT::TempStorage static_smem; + extern __shared__ __align__(16) unsigned char dynamic_smem[]; - // Shared memory for AgentHistogram - __shared__ typename AgentHistogramT::TempStorage temp_storage; + CounterT* dynamic_smem_privatized_histograms = nullptr; + if constexpr (is_privatized_dynamic_smem_v) + { + dynamic_smem_privatized_histograms = reinterpret_cast(dynamic_smem); + } + + _CCCL_PRAGMA_UNROLL_FULL() + for (int channel = 0; channel < NumActiveChannels; ++channel) + { + output_decode_op_wrapper[channel].Precompute(); + privatized_decode_op_wrapper[channel].Precompute(); + } AgentHistogramT agent( - temp_storage, + static_smem, d_samples, num_output_bins_wrapper.data(), num_privatized_bins_wrapper.data(), d_output_histograms_wrapper.data(), d_privatized_histograms_wrapper.data(), output_decode_op_wrapper.data(), - privatized_decode_op_wrapper.data()); + privatized_decode_op_wrapper.data(), + dynamic_smem_privatized_histograms); // Initialize counters agent.InitBinCounters(); @@ -525,8 +810,8 @@ __launch_bounds__(int(current_policy().threads_per_block)) //! @tparam PolicySelector //! Selects the tuning policy //! -//! @tparam PrivatizedSmemBins -//! Maximum number of histogram bins per channel (e.g., up to 256) +//! @tparam PrivatizationMode +//! Storage mode for the privatized histogram //! //! @tparam NumChannels //! Number of channels interleaved in the input data (may be greater than the number of channels @@ -602,7 +887,7 @@ __launch_bounds__(int(current_policy().threads_per_block)) //! @param tile_queue //! Drain queue descriptor for dynamically mapping tile data onto thread blocks template + bool IsEven, + typename OutputCounterT = CounterT> #if _CCCL_HAS_CONCEPTS() requires histogram_policy_selector #endif // _CCCL_HAS_CONCEPTS() -__launch_bounds__(int(current_policy().threads_per_block)) +__launch_bounds__(int(histogram_privatization_policy().threads_per_block), + int(histogram_min_blocks_per_sm())) _CCCL_KERNEL_ATTRIBUTES void DeviceHistogramSweepDeviceInitKernel( const SampleIteratorT d_samples, ::cuda::std::array num_output_bins_wrapper, ::cuda::std::array num_privatized_bins_wrapper, - ::cuda::std::array d_output_histograms_wrapper, + ::cuda::std::array d_output_histograms_wrapper, ::cuda::std::array d_privatized_histograms_wrapper, const FirstLevelArrayT first_level_array, const SecondLevelArrayT second_level_array, @@ -632,8 +919,6 @@ __launch_bounds__(int(current_policy().threads_per_block)) const int tiles_per_row, const GridQueue tile_queue) { - static constexpr HistogramPolicy hp = current_policy(); - OutputDecodeOpT output_decode_op[NumActiveChannels]; PrivatizedDecodeOpT privatized_decode_op[NumActiveChannels]; if constexpr (IsEven) @@ -648,6 +933,7 @@ __launch_bounds__(int(current_policy().threads_per_block)) output_decode_op[channel].Init(num_levels, upper_level, lower_level); } } + else { _CCCL_PRAGMA_UNROLL_FULL() @@ -660,39 +946,44 @@ __launch_bounds__(int(current_policy().threads_per_block)) } } - // Thread block type for compositing input tiles - using AgentHistogramPolicyT = agent_histogram_policy< - hp.threads_per_block, - hp.pixels_per_thread, - hp.load_algorithm, - hp.load_modifier, - hp.rle_compress, - hp.mem_preference, - hp.use_work_stealing, - hp.vec_size>; + _CCCL_PRAGMA_UNROLL_FULL() + for (int channel = 0; channel < NumActiveChannels; ++channel) + { + output_decode_op[channel].Precompute(); + privatized_decode_op[channel].Precompute(); + } + using AgentHistogramT = - AgentHistogram; + OffsetT, + OutputCounterT>; + + __shared__ typename AgentHistogramT::TempStorage static_smem; + extern __shared__ __align__(16) unsigned char dynamic_smem[]; - // Shared memory for AgentHistogram - __shared__ typename AgentHistogramT::TempStorage temp_storage; + CounterT* dynamic_smem_privatized_histograms = nullptr; + if constexpr (is_privatized_dynamic_smem_v) + { + dynamic_smem_privatized_histograms = reinterpret_cast(dynamic_smem); + } AgentHistogramT agent( - temp_storage, + static_smem, d_samples, num_output_bins_wrapper.data(), num_privatized_bins_wrapper.data(), d_output_histograms_wrapper.data(), d_privatized_histograms_wrapper.data(), output_decode_op, - privatized_decode_op); + privatized_decode_op, + dynamic_smem_privatized_histograms); // Initialize counters agent.InitBinCounters(); diff --git a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh index 2084666a5af3..2af37486d959 100644 --- a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh @@ -22,32 +22,74 @@ #include #include #include +#include CUB_NAMESPACE_BEGIN -//! The tuning policy for all algorithms in @ref DeviceHistogram. -struct HistogramPolicy +//! The tuning policy for one DeviceHistogram privatization technique. +struct HistogramPrivatizationPolicy { int threads_per_block; //!< Number of threads in a CUDA block - int pixels_per_thread; //!< Number of pixels processed per thread + int items_per_thread; //!< Number of items processed per thread int vec_size; //!< Vectorization size for loading samples - BlockLoadAlgorithm load_algorithm; //!< The @ref BlockLoadAlgorithm used for loading samples from global memory - CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading samples from global memory - bool rle_compress; //!< Whether to perform localized RLE to compress samples before histogramming - BlockHistogramMemoryPreference mem_preference; //!< Whether to prefer privatized shared-memory or global-memory bins, - //!< or a mix of both - bool use_work_stealing; //!< Whether to dequeue tiles from a global work queue - int init_kernel_pdl_trigger_max_bins; //!< Maximum number of bins for the init kernel to trigger the histogram kernel - //!< early using PDL + BlockLoadAlgorithm load_algorithm; //!< Algorithm used for loading samples + CacheLoadModifier load_modifier; //!< Cache modifier used for loading samples + bool rle_compress; //!< Whether to locally run-length encode samples + bool work_stealing; //!< Whether blocks dequeue tiles from a global work queue [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool - operator==(const HistogramPolicy& lhs, const HistogramPolicy& rhs) noexcept + operator==(const HistogramPrivatizationPolicy& lhs, const HistogramPrivatizationPolicy& rhs) noexcept { - return lhs.threads_per_block == rhs.threads_per_block && lhs.pixels_per_thread == rhs.pixels_per_thread + return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread && lhs.vec_size == rhs.vec_size && lhs.load_algorithm == rhs.load_algorithm && lhs.load_modifier == rhs.load_modifier && lhs.rle_compress == rhs.rle_compress - && lhs.mem_preference == rhs.mem_preference && lhs.use_work_stealing == rhs.use_work_stealing - && lhs.init_kernel_pdl_trigger_max_bins == rhs.init_kernel_pdl_trigger_max_bins; + && lhs.work_stealing == rhs.work_stealing; + } + +#if _CCCL_HOSTED() + friend ::std::ostream& operator<<(::std::ostream& os, const HistogramPrivatizationPolicy& p) + { + return os + << "HistogramPrivatizationPolicy { .threads_per_block = " << p.threads_per_block + << ", .items_per_thread = " << p.items_per_thread << ", .vec_size = " << p.vec_size + << ", .load_algorithm = " << p.load_algorithm << ", .load_modifier = " << p.load_modifier + << ", .rle_compress = " << p.rle_compress << ", .work_stealing = " << p.work_stealing << " }"; + } +#endif // _CCCL_HOSTED() +}; + +//! The tuning policy for all DeviceHistogram sweep passes. +struct HistogramPolicy +{ + HistogramPrivatizationPolicy gmem; //!< Policy for global-memory privatization + HistogramPrivatizationPolicy static_smem; //!< Policy for compile-time-sized shared-memory privatization + HistogramPrivatizationPolicy dynamic_smem; //!< Policy for runtime-sized shared-memory privatization + int init_threads_per_block; //!< Number of threads in a histogram initialization block + int max_privatized_static_smem_single_channel_bytes; //!< Single-channel compile-time-sized SMEM limit + int max_privatized_dynamic_smem_single_channel_bytes; //!< Single-channel runtime-sized SMEM limit + int static_smem_min_blocks_per_sm; //!< Minimum blocks per SM requested by the static-SMEM launch bounds + int max_privatized_dynamic_smem_multi_channel_range_bytes; //!< Multi-channel HistogramRange SMEM limit + int max_privatized_dynamic_smem_2_channel_even_bytes; //!< Two-channel HistogramEven SMEM limit + int max_privatized_dynamic_smem_3_channel_even_bytes; //!< Three-channel HistogramEven SMEM limit + int max_privatized_dynamic_smem_4_channel_even_bytes; //!< Four-channel HistogramEven SMEM limit + int min_cached_search_gmem_range_bins; //!< Minimum RANGE bin count for cached search with GMEM privatization + int max_output_histogram_bytes_for_init_kernel_pdl; //!< Largest output allocation for init-kernel PDL + + [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool + operator==(const HistogramPolicy& lhs, const HistogramPolicy& rhs) noexcept + { + return lhs.gmem == rhs.gmem && lhs.static_smem == rhs.static_smem && lhs.dynamic_smem == rhs.dynamic_smem + && lhs.init_threads_per_block == rhs.init_threads_per_block + && lhs.max_privatized_static_smem_single_channel_bytes == rhs.max_privatized_static_smem_single_channel_bytes + && lhs.max_privatized_dynamic_smem_single_channel_bytes == rhs.max_privatized_dynamic_smem_single_channel_bytes + && lhs.static_smem_min_blocks_per_sm == rhs.static_smem_min_blocks_per_sm + && lhs.max_privatized_dynamic_smem_multi_channel_range_bytes + == rhs.max_privatized_dynamic_smem_multi_channel_range_bytes + && lhs.max_privatized_dynamic_smem_2_channel_even_bytes == rhs.max_privatized_dynamic_smem_2_channel_even_bytes + && lhs.max_privatized_dynamic_smem_3_channel_even_bytes == rhs.max_privatized_dynamic_smem_3_channel_even_bytes + && lhs.max_privatized_dynamic_smem_4_channel_even_bytes == rhs.max_privatized_dynamic_smem_4_channel_even_bytes + && lhs.min_cached_search_gmem_range_bins == rhs.min_cached_search_gmem_range_bins + && lhs.max_output_histogram_bytes_for_init_kernel_pdl == rhs.max_output_histogram_bytes_for_init_kernel_pdl; } [[nodiscard]] _CCCL_HOST_DEVICE_API friend constexpr bool @@ -60,218 +102,124 @@ struct HistogramPolicy friend ::std::ostream& operator<<(::std::ostream& os, const HistogramPolicy& p) { return os - << "HistogramPolicy { .threads_per_block = " << p.threads_per_block << ", .pixels_per_thread = " - << p.pixels_per_thread << ", .vec_size = " << p.vec_size << ", .load_algorithm = " << p.load_algorithm - << ", .load_modifier = " << p.load_modifier << ", .rle_compress = " << p.rle_compress - << ", .mem_preference = " << p.mem_preference << ", .use_work_stealing = " << p.use_work_stealing - << ", .init_kernel_pdl_trigger_max_bins = " << p.init_kernel_pdl_trigger_max_bins << " }"; + << "HistogramPolicy { .gmem = " << p.gmem << ", .static_smem = " << p.static_smem + << ", .dynamic_smem = " << p.dynamic_smem << ", .init_threads_per_block = " << p.init_threads_per_block + << ", .max_privatized_static_smem_single_channel_bytes = " << p.max_privatized_static_smem_single_channel_bytes + << ", .max_privatized_dynamic_smem_single_channel_bytes = " + << p.max_privatized_dynamic_smem_single_channel_bytes << ", .static_smem_min_blocks_per_sm = " + << p.static_smem_min_blocks_per_sm << ", .max_privatized_dynamic_smem_multi_channel_range_bytes = " + << p.max_privatized_dynamic_smem_multi_channel_range_bytes + << ", .max_privatized_dynamic_smem_2_channel_even_bytes = " + << p.max_privatized_dynamic_smem_2_channel_even_bytes + << ", .max_privatized_dynamic_smem_3_channel_even_bytes = " + << p.max_privatized_dynamic_smem_3_channel_even_bytes + << ", .max_privatized_dynamic_smem_4_channel_even_bytes = " + << p.max_privatized_dynamic_smem_4_channel_even_bytes + << ", .min_cached_search_gmem_range_bins = " << p.min_cached_search_gmem_range_bins + << ", .max_output_histogram_bytes_for_init_kernel_pdl = " << p.max_output_histogram_bytes_for_init_kernel_pdl + << " }"; } -#endif // _CCCL_HOSTED() +#endif }; namespace detail::histogram { -// TODO(bgruber): drop in CCCL 4.0 -enum class primitive_sample -{ - no, - yes -}; - -// TODO(bgruber): drop in CCCL 4.0 -enum class sample_size -{ - _1, - _2, - _4, - _8, - unknown -}; - -// TODO(bgruber): drop in CCCL 4.0 -enum class counter_size +enum class privatization_mode { - _4, - unknown + gmem, + static_smem, + dynamic_smem }; -// TODO(bgruber): drop in CCCL 4.0 -template -_CCCL_HOST_DEVICE_API constexpr primitive_sample is_primitive_sample() -{ - return is_primitive::value ? primitive_sample::yes : primitive_sample::no; -} - -// TODO(bgruber): drop in CCCL 4.0 -template -_CCCL_HOST_DEVICE_API constexpr counter_size classify_counter_size() +template +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr int max_privatized_smem_bins(int max_privatized_smem_bytes) { - return sizeof(CounterT) == 4 ? counter_size::_4 : counter_size::unknown; + static_assert(NumActiveChannels > 0); + if (max_privatized_smem_bytes <= 0) + { + return 0; + } + return max_privatized_smem_bytes / int{sizeof(CounterT)} / NumActiveChannels; } -// TODO(bgruber): drop in CCCL 4.0 -template -_CCCL_HOST_DEVICE_API constexpr sample_size classify_sample_size() +template +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr int dynamic_smem_limit_bytes(const HistogramPolicy& policy) { - return sizeof(SampleT) == 1 ? sample_size::_1 : sizeof(SampleT) == 2 ? sample_size::_2 : sample_size::unknown; + int dynamic_smem_max_bytes = policy.max_privatized_dynamic_smem_single_channel_bytes; + if constexpr (NumActiveChannels > 1) + { + if constexpr (IsEven) + { + dynamic_smem_max_bytes = + NumActiveChannels == 2 ? policy.max_privatized_dynamic_smem_2_channel_even_bytes + : NumActiveChannels == 3 ? policy.max_privatized_dynamic_smem_3_channel_even_bytes + : NumActiveChannels == 4 + ? policy.max_privatized_dynamic_smem_4_channel_even_bytes + : 0; + } + else + { + dynamic_smem_max_bytes = policy.max_privatized_dynamic_smem_multi_channel_range_bytes; + } + } + return dynamic_smem_max_bytes; } -// TODO(bgruber): drop in CCCL 4.0 -template (), - sample_size SampleSize = classify_sample_size()> -struct sm90_tuning; - -template -struct sm90_tuning -{ - static constexpr int threads = 768; - static constexpr int items = 12; - - static constexpr CacheLoadModifier load_modifier = LOAD_LDG; - static constexpr BlockHistogramMemoryPreference mem_preference = SMEM; - - static constexpr BlockLoadAlgorithm load_algorithm = BLOCK_LOAD_DIRECT; - - static constexpr bool rle_compress = false; - static constexpr bool use_work_stealing = false; -}; - -template -struct sm90_tuning +template +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto +select_privatization_mode(const HistogramPolicy& policy, int num_bins) -> privatization_mode { - static constexpr int threads = 960; - static constexpr int items = 10; - - static constexpr CacheLoadModifier load_modifier = LOAD_DEFAULT; - static constexpr BlockHistogramMemoryPreference mem_preference = SMEM; - - static constexpr BlockLoadAlgorithm load_algorithm = BLOCK_LOAD_DIRECT; - - static constexpr bool rle_compress = true; - static constexpr bool use_work_stealing = false; -}; - -// TODO(bgruber): drop in CCCL 4.0 -template (), - sample_size SampleSize = classify_sample_size()> -struct sm100_tuning; - -// even -template -struct sm100_tuning -{ - // ipt_12.tpb_928.rle_0.ws_0.mem_1.ld_2.laid_0.vec_2 1.033332 0.940517 1.031835 1.195876 - static constexpr int items = 12; - static constexpr int threads = 928; - static constexpr bool rle_compress = false; - static constexpr bool use_work_stealing = false; - static constexpr BlockHistogramMemoryPreference mem_preference = SMEM; - static constexpr CacheLoadModifier load_modifier = LOAD_CA; - static constexpr BlockLoadAlgorithm load_algorithm = BLOCK_LOAD_DIRECT; - static constexpr int vec_size = 1 << 2; -}; + if (num_bins <= 0) + { + return privatization_mode::gmem; + } -// sample_size 2/4/8 showed no benefit over SM90 during verification benchmarks + const int static_smem_max_bins = + max_privatized_smem_bins(policy.max_privatized_static_smem_single_channel_bytes); + const int dynamic_smem_max_bytes = dynamic_smem_limit_bytes(policy); + const int dynamic_smem_max_bins = max_privatized_smem_bins(dynamic_smem_max_bytes); + if (num_bins <= static_smem_max_bins) + { + return privatization_mode::static_smem; + } + if (num_bins <= dynamic_smem_max_bins) + { + return privatization_mode::dynamic_smem; + } + return privatization_mode::gmem; +} -// range -template -struct sm100_tuning +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr bool +use_cached_search_for_gmem_range(const HistogramPolicy& policy, int num_bins) { - // ipt_12.tpb_448.rle_0.ws_0.mem_1.ld_1.laid_0.vec_2 1.078987 0.985542 1.085118 1.175637 - static constexpr int items = 12; - static constexpr int threads = 448; - static constexpr bool rle_compress = false; - static constexpr bool use_work_stealing = false; - static constexpr BlockHistogramMemoryPreference mem_preference = SMEM; - static constexpr CacheLoadModifier load_modifier = LOAD_LDG; - static constexpr BlockLoadAlgorithm load_algorithm = BLOCK_LOAD_DIRECT; - static constexpr int vec_size = 1 << 2; -}; - -// sample_size 2/4/8 showed no benefit over SM90 during verification benchmarks - -// multi.even and multi.range: none of the found tunings surpassed the SM90 tuning during verification benchmarks + return policy.min_cached_search_gmem_range_bins > 0 && num_bins >= policy.min_cached_search_gmem_range_bins; +} -// TODO(bgruber): drop in CCCL 4.0 -template -struct policy_hub +// The C Parallel API erases CounterT before host dispatch, so its bridge must select from the +// preserved runtime counter width. Typed CUB dispatch uses the overload above. +template +[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto +select_privatization_mode_for_counter_size(const HistogramPolicy& policy, int num_bins, int counter_size_bytes) + -> privatization_mode { - // TODO(bgruber): move inside t_scale in C++14 - static constexpr int v_scale = (sizeof(SampleT) + sizeof(int) - 1) / sizeof(int); - - _CCCL_HOST_DEVICE_API static constexpr int t_scale(int nominalItemsPerThread) + if (num_bins <= 0 || counter_size_bytes <= 0) { - return (::cuda::std::max) (nominalItemsPerThread / NumActiveChannels / v_scale, 1); + return privatization_mode::gmem; } - // SM50 - struct Policy500 : detail::chained_policy<500, Policy500, Policy500> - { - // TODO This might be worth it to separate usual histogram and the multi one - using AgentHistogramPolicyT = - agent_histogram_policy<384, t_scale(16), BLOCK_LOAD_DIRECT, LOAD_LDG, true, SMEM, false>; - }; - - // SM90 - struct Policy900 : detail::chained_policy<900, Policy900, Policy500> + const int static_smem_max_bins = policy.max_privatized_static_smem_single_channel_bytes / counter_size_bytes; + const int dynamic_smem_max_bins = + dynamic_smem_limit_bytes(policy) / counter_size_bytes / NumActiveChannels; + if (num_bins <= static_smem_max_bins) { - // Use values from tuning if a specialization exists, otherwise pick Policy500 - template - _CCCL_HOST_DEVICE_API static auto select_agent_policy(int) - -> agent_histogram_policy; - - template - _CCCL_HOST_DEVICE_API static auto select_agent_policy(long) -> typename Policy500::AgentHistogramPolicyT; - - using AgentHistogramPolicyT = - decltype(select_agent_policy< - sm90_tuning()>>(0)); - - static constexpr int init_kernel_pdl_trigger_max_bins = 2048; - }; - - struct Policy1000 : detail::chained_policy<1000, Policy1000, Policy900> + return privatization_mode::static_smem; + } + if (num_bins <= dynamic_smem_max_bins) { - // Use values from tuning if a specialization exists, otherwise pick Policy900 - template - _CCCL_HOST_DEVICE_API static auto select_agent_policy(int) -> agent_histogram_policy< - Tuning::threads, - Tuning::items, - Tuning::load_algorithm, - Tuning::load_modifier, - Tuning::rle_compress, - Tuning::mem_preference, - Tuning::use_work_stealing, - Tuning::vec_size>; - - template - _CCCL_HOST_DEVICE_API static auto select_agent_policy(long) -> typename Policy900::AgentHistogramPolicyT; - - using AgentHistogramPolicyT = - decltype(select_agent_policy< - sm100_tuning()>>( - 0)); - - static constexpr int init_kernel_pdl_trigger_max_bins = 2048; - }; - - using MaxPolicy = Policy1000; -}; + return privatization_mode::dynamic_smem; + } + return privatization_mode::gmem; +} #if _CCCL_HAS_CONCEPTS() template @@ -280,9 +228,10 @@ concept histogram_policy_selector = policy_selector; struct policy_selector { - bool sample_is_primitive; + bool sample_is_primitive; //!< Whether the sample opts into CUB's primitive-type tuning category + // Kept separately from sample_size_bytes to preserve the serialized C Parallel selector layout. int sample_size; - int counter_size; + int counter_size_bytes; int sample_size_bytes; int num_channels; int num_active_channels; @@ -298,43 +247,182 @@ private: public: [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const -> HistogramPolicy { + // SM100 and SM120 use the autoresearch launch shapes. Their dynamic-SMEM budgets differ because SM100 permits + // 227 KiB per block while SM120 permits 99 KiB per block. if (cc >= ::cuda::compute_capability{10, 0}) { - if (num_channels == 1 && num_active_channels == 1 && counter_size == 4 && sample_is_primitive && sample_size == 1) + const bool is_sm120_or_newer = cc >= ::cuda::compute_capability{12, 0}; + const bool single_channel = num_channels == 1 && num_active_channels == 1; + auto gmem = HistogramPrivatizationPolicy{384, t_scale(16), 4, BLOCK_LOAD_DIRECT, LOAD_LDG, true, false}; + + // Single-channel primitive samples with 32-bit counters use their per-sample-width tuning. + if (single_channel && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} && sample_is_primitive) { - if (is_even) + // Eight-bit EVEN and RANGE histograms retain the dedicated SM100 tunings already in main. + if (sample_size_bytes == 1) { - // ipt_12.tpb_928.rle_0.ws_0.mem_1.ld_2.laid_0.vec_2 1.033332 0.940517 1.031835 1.195876 - return HistogramPolicy{928, 12, 1 << 2, BLOCK_LOAD_DIRECT, LOAD_CA, false, SMEM, false, 2048}; + gmem = is_even ? HistogramPrivatizationPolicy{928, 12, 4, BLOCK_LOAD_DIRECT, LOAD_CA, false, false} + : HistogramPrivatizationPolicy{448, 12, 4, BLOCK_LOAD_DIRECT, LOAD_LDG, false, false}; } - else + // Sixteen-bit samples retain the SM90 tuning because autoresearch did not improve it. + else if (sample_size_bytes == 2) { - // ipt_12.tpb_448.rle_0.ws_0.mem_1.ld_1.laid_0.vec_2 1.078987 0.985542 1.085118 1.175637 - return HistogramPolicy{448, 12, 1 << 2, BLOCK_LOAD_DIRECT, LOAD_LDG, false, SMEM, false, 2048}; + gmem = HistogramPrivatizationPolicy{960, 10, 4, BLOCK_LOAD_DIRECT, LOAD_DEFAULT, true, false}; } + // Thirty-two-bit samples use the best sweep shape measured by autoresearch. + else if (sample_size_bytes == 4) + { + gmem = HistogramPrivatizationPolicy{768, 12, 4, BLOCK_LOAD_DIRECT, LOAD_LDG, true, false}; + } + // Sixty-four-bit samples use the best sweep shape measured by autoresearch. + else if (sample_size_bytes == 8) + { + gmem = HistogramPrivatizationPolicy{768, 6, 4, BLOCK_LOAD_DIRECT, LOAD_LDG, true, false}; + } + } + + auto static_smem = gmem; + const bool range_multi_static = + !is_even && num_channels > 1 && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} && sample_is_primitive; + const bool range_u32_static = + !is_even && single_channel && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} && sample_is_primitive + && sample_size_bytes == 4; + const bool range_u64_static = + !is_even && single_channel && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} && sample_is_primitive + && sample_size_bytes == 8; + // Multi-channel and 64-bit-sample RANGE favor narrower blocks in the static-SMEM tier. + if (range_multi_static || range_u64_static) + { + static_smem.threads_per_block = 384; + } + // Thirty-two-bit-sample RANGE retains the wider block that won in the static-SMEM tier. + else if (range_u32_static) + { + static_smem.threads_per_block = 768; + } + // Sixty-four-bit-sample RANGE recovers the higher static-tier items-per-thread count. + if (range_u64_static) + { + static_smem.items_per_thread = t_scale(16); } - // sample_size 2/4/8 showed no benefit over SM90 during verification benchmarks - // multi.even and multi.range: none of the found tunings surpassed the SM90 tuning during verification benchmarks + // All storage thresholds are byte budgets. Dispatch derives the corresponding + // bin limits from the local counter width and active channel count. + constexpr int max_privatized_static_smem_bytes = 1024; + constexpr int max_privatized_dynamic_smem_sm100_bytes = 228352; + constexpr int max_privatized_dynamic_smem_sm120_bytes = 99 * 1024; + constexpr int max_privatized_dynamic_smem_range_bytes_per_channel = 8192; + constexpr int max_privatized_dynamic_smem_even_bytes_per_channel = 32768; + constexpr int min_cached_search_gmem_single_channel_range_bins = 1; + constexpr int min_cached_search_gmem_multi_channel_range_bins = 16384; + constexpr int init_threads_per_block = 256; + constexpr int max_output_histogram_bytes_for_init_kernel_pdl = 8192; + const int max_privatized_dynamic_smem_single_channel_bytes = + is_sm120_or_newer ? max_privatized_dynamic_smem_sm120_bytes : max_privatized_dynamic_smem_sm100_bytes; + + const bool supports_dynamic_smem = + counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} && sample_is_primitive; + const bool has_single_channel_dynamic_smem = + supports_dynamic_smem && single_channel + && (sample_size_bytes == 1 || sample_size_bytes == 4 || sample_size_bytes == 8); + const bool has_multi_channel_dynamic_smem = supports_dynamic_smem && num_channels > 1; + int dynamic_smem_single_channel_bytes = + has_single_channel_dynamic_smem ? max_privatized_dynamic_smem_single_channel_bytes : 0; + int dynamic_smem_multi_channel_range_bytes = + has_single_channel_dynamic_smem || (has_multi_channel_dynamic_smem && !is_even) + ? max_privatized_dynamic_smem_range_bytes_per_channel * num_active_channels + : 0; + int dynamic_smem_multi_channel_even_bytes = 0; + if (has_multi_channel_dynamic_smem && is_even) + { + dynamic_smem_multi_channel_even_bytes = + max_privatized_dynamic_smem_even_bytes_per_channel * num_active_channels; + if (is_sm120_or_newer) + { + dynamic_smem_multi_channel_even_bytes = + (::cuda::std::min) (dynamic_smem_multi_channel_even_bytes, max_privatized_dynamic_smem_sm120_bytes); + } + } + const int init_kernel_pdl_trigger_bytes = + single_channel && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} && sample_is_primitive + && (sample_size_bytes == 1 || sample_size_bytes == 2 || sample_size_bytes == 4 || sample_size_bytes == 8) + ? max_output_histogram_bytes_for_init_kernel_pdl + : 0; + const int min_cached_search_gmem_range_bins = + !is_even && sample_is_primitive && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} + && (sample_size_bytes == 4 || sample_size_bytes == 8) + ? (single_channel ? min_cached_search_gmem_single_channel_range_bins + : min_cached_search_gmem_multi_channel_range_bins) + : 0; + + return HistogramPolicy{ + gmem, + static_smem, + gmem, + init_threads_per_block, + max_privatized_static_smem_bytes, + dynamic_smem_single_channel_bytes, + range_multi_static || range_u64_static ? 3 : 0, + dynamic_smem_multi_channel_range_bytes, + has_multi_channel_dynamic_smem && is_even && num_active_channels == 2 + ? dynamic_smem_multi_channel_even_bytes + : 0, + has_multi_channel_dynamic_smem && is_even && num_active_channels == 3 + ? dynamic_smem_multi_channel_even_bytes + : 0, + has_multi_channel_dynamic_smem && is_even && num_active_channels == 4 + ? dynamic_smem_multi_channel_even_bytes + : 0, + min_cached_search_gmem_range_bins, + init_kernel_pdl_trigger_bytes}; } + // SM90 uses its established single-channel 8-bit and 16-bit specializations. if (cc >= ::cuda::compute_capability{9, 0}) { - if (num_channels == 1 && num_active_channels == 1 && counter_size == 4 && sample_is_primitive) + auto sweep = HistogramPrivatizationPolicy{384, t_scale(16), 4, BLOCK_LOAD_DIRECT, LOAD_LDG, true, false}; + // Single-channel primitive samples with 32-bit counters use the established SM90 specializations. + if (num_channels == 1 && num_active_channels == 1 && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} + && sample_is_primitive) { - if (sample_size == 1) + // Eight-bit samples use the tuned SM90 sweep. + if (sample_size_bytes == 1) { - return HistogramPolicy{768, 12, 1 << 2, BLOCK_LOAD_DIRECT, LOAD_LDG, false, SMEM, false, 2048}; + sweep = HistogramPrivatizationPolicy{768, 12, 4, BLOCK_LOAD_DIRECT, LOAD_LDG, false, false}; } - else if (sample_size == 2) + // Sixteen-bit samples use the tuned SM90 sweep. + else if (sample_size_bytes == 2) { - return HistogramPolicy{960, 10, 1 << 2, BLOCK_LOAD_DIRECT, LOAD_DEFAULT, true, SMEM, false, 2048}; + sweep = HistogramPrivatizationPolicy{960, 10, 4, BLOCK_LOAD_DIRECT, LOAD_DEFAULT, true, false}; } } + constexpr int max_privatized_static_smem_bytes = 1024; + constexpr int init_threads_per_block = 256; + constexpr int max_output_histogram_bytes_for_init_kernel_pdl = 8192; + const int init_kernel_pdl_trigger_bytes = + num_channels == 1 && num_active_channels == 1 && counter_size_bytes == int{sizeof(::cuda::std::uint32_t)} + && sample_is_primitive && (sample_size_bytes == 1 || sample_size_bytes == 2) + ? max_output_histogram_bytes_for_init_kernel_pdl + : 0; + return HistogramPolicy{ + sweep, + sweep, + sweep, + init_threads_per_block, + max_privatized_static_smem_bytes, + 0, + 0, + 0, + 0, + 0, + 0, + 0, + init_kernel_pdl_trigger_bytes}; } - // fallback from SM50 - return HistogramPolicy{384, t_scale(16), 4, BLOCK_LOAD_DIRECT, LOAD_LDG, true, SMEM, false, 0}; + // Architectures before SM90 use the longstanding generic histogram tuning. + const auto sweep = HistogramPrivatizationPolicy{384, t_scale(16), 4, BLOCK_LOAD_DIRECT, LOAD_LDG, true, false}; + return HistogramPolicy{sweep, sweep, sweep, 256, 1024, 0, 0, 0, 0, 0, 0, 0, 0}; } }; @@ -347,15 +435,14 @@ struct policy_selector_from_types { [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const -> HistogramPolicy { - constexpr auto policies = policy_selector{ + return policy_selector{ is_primitive_v, int{sizeof(SampleT)}, int{sizeof(CounterT)}, int{sizeof(SampleT)}, NumChannels, NumActiveChannels, - IsEven}; - return policies(cc); + IsEven}(cc); } }; } // namespace detail::histogram diff --git a/cub/test/catch2_test_device_histogram.cu b/cub/test/catch2_test_device_histogram.cu index 6cdf13cb33a5..86e6e335e770 100644 --- a/cub/test/catch2_test_device_histogram.cu +++ b/cub/test/catch2_test_device_histogram.cu @@ -574,6 +574,25 @@ CUB_TEST_LIST("DeviceHistogram::Histogram* channel configs", test_even_and_range(256, 256 + 1, 128, 32); } +CUB_TEST("DeviceHistogram::Histogram* dynamic shared-memory privatization", "[histogram][device]", CUB_SMALL) +{ + int current_device{}; + REQUIRE(cudaSuccess == cudaGetDevice(¤t_device)); + + cuda::compute_capability cc{}; + REQUIRE(cudaSuccess == cub::detail::ptx_compute_cap(cc, current_device)); + if (cc < cuda::compute_capability{10, 0}) + { + SKIP("The runtime-sized shared-memory histogram policy is currently tuned for SM100"); + } + + using counter_t = unsigned int; + const int num_levels = GENERATE(1025, 4097, 8193); + + test_even_and_range(num_levels - 1, num_levels, 4096, 4); + test_even_and_range(num_levels - 1, num_levels, 4096, 4); +} + // Testing only HistogramEven is fine, because HistogramRange shares the loading logic and the different binning // implementations are not affected by the iterator. CUB_TEST("DeviceHistogram::HistogramEven sample iterator", "[histogram_even][device]", CUB_SMALL) @@ -657,6 +676,52 @@ CUB_TEST("DeviceHistogram::HistogramRange levels/samples aliasing", "[histogram_ } } +CUB_TEST("DeviceHistogram::HistogramRange interpolation avoids signed overflow", "[histogram_range][device]", CUB_SMALL) +{ + int current_device{}; + REQUIRE(cudaSuccess == cudaGetDevice(¤t_device)); + + cuda::compute_capability cc{}; + REQUIRE(cudaSuccess == cub::detail::ptx_compute_cap(cc, current_device)); + if (cc < cuda::compute_capability{10, 0}) + { + SKIP("The runtime-sized shared-memory histogram policy is currently tuned for SM100"); + } + + using sample_t = int; + constexpr int num_bins = 1024; + + c2h::host_vector h_levels(num_bins + 1); + constexpr auto lo = static_cast(cs::numeric_limits::lowest()); + constexpr auto hi = static_cast(cs::numeric_limits::max()); + constexpr auto range = hi - lo; + for (int i = 0; i <= num_bins; ++i) + { + h_levels[i] = static_cast(lo + (range * i) / num_bins); + } + + const c2h::host_vector h_samples{ + cs::numeric_limits::lowest(), -1, 0, 1, cs::numeric_limits::max() - 1}; + c2h::device_vector d_levels = h_levels; + c2h::device_vector d_samples = h_samples; + c2h::device_vector d_histogram(num_bins, 0); + + histogram_range( + thrust::raw_pointer_cast(d_samples.data()), + thrust::raw_pointer_cast(d_histogram.data()), + num_bins + 1, + thrust::raw_pointer_cast(d_levels.data()), + static_cast(d_samples.size())); + + c2h::host_vector expected(num_bins, 0); + for (const sample_t sample : h_samples) + { + const auto upper = std::upper_bound(h_levels.begin(), h_levels.end(), sample); + ++expected[static_cast(std::distance(h_levels.begin(), upper) - 1)]; + } + REQUIRE(d_histogram == expected); +} + // Limit this large-memory reproducer to the host launch path. #if TEST_LAUNCH == 0 CUB_TEST("DeviceHistogram::MultiHistogramEven large privatized offsets", "[histogram_even][device]", CUB_LARGE) diff --git a/cub/test/catch2_test_device_histogram_custom_policy_hub.cu b/cub/test/catch2_test_device_histogram_custom_policy_hub.cu index 9e976687d335..494c0828a45c 100644 --- a/cub/test/catch2_test_device_histogram_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_histogram_custom_policy_hub.cu @@ -20,16 +20,48 @@ using namespace cub; template struct my_policy_hub { - // simplified from Policy500 of the CUB histogram tunings - struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy> + struct Policy500 : cub::detail::chained_policy<500, Policy500, Policy500> { - using AgentHistogramPolicyT = AgentHistogramPolicy<384, 16, BLOCK_LOAD_DIRECT, LOAD_LDG, true, SMEM, false>; + using AgentHistogramPolicyT = AgentHistogramPolicy<384, 16, BLOCK_LOAD_DIRECT, LOAD_LDG, true, false>; + static constexpr int init_kernel_pdl_trigger_max_bins = 0; + }; + + struct Policy900 : cub::detail::chained_policy<900, Policy900, Policy500> + { + using AgentHistogramPolicyT = AgentHistogramPolicy<256, 8, BLOCK_LOAD_DIRECT, LOAD_DEFAULT, false, false>; static constexpr int init_kernel_pdl_trigger_max_bins = 2048; }; + + using MaxPolicy = Policy900; }; CUB_TEST("DispatchHistogram::DispatchEven: custom policy hub", "[histogram][device]", CUB_SMALL) { + using custom_max_policy_t = typename my_policy_hub::MaxPolicy; + const auto custom_sm75_policy = + cub::detail::histogram::policy_selector_from_hub{}(cuda::compute_capability{7, 5}); + const auto custom_sm90_policy = + cub::detail::histogram::policy_selector_from_hub{}(cuda::compute_capability{9, 0}); + const auto custom_wide_counter_policy = + cub::detail::histogram::policy_selector_from_hub{}( + cuda::compute_capability{7, 5}); + const auto custom_wide_counter_sm90_policy = + cub::detail::histogram::policy_selector_from_hub{}( + cuda::compute_capability{9, 0}); + REQUIRE(custom_sm75_policy.gmem.threads_per_block == 384); + REQUIRE(custom_sm75_policy.gmem.items_per_thread == 16); + REQUIRE(custom_sm90_policy.gmem.threads_per_block == 256); + REQUIRE(custom_sm90_policy.gmem.items_per_thread == 8); + REQUIRE(custom_sm75_policy.max_privatized_static_smem_single_channel_bytes == 256 * sizeof(unsigned int)); + REQUIRE( + custom_wide_counter_policy.max_privatized_static_smem_single_channel_bytes == 256 * sizeof(unsigned long long)); + REQUIRE(custom_sm90_policy.max_privatized_dynamic_smem_single_channel_bytes == 0); + REQUIRE(custom_sm90_policy.min_cached_search_gmem_range_bins == 0); + REQUIRE(custom_sm75_policy.max_output_histogram_bytes_for_init_kernel_pdl == 0); + REQUIRE(custom_sm90_policy.max_output_histogram_bytes_for_init_kernel_pdl == 8192); + REQUIRE(custom_wide_counter_policy.max_output_histogram_bytes_for_init_kernel_pdl == 0); + REQUIRE(custom_wide_counter_sm90_policy.max_output_histogram_bytes_for_init_kernel_pdl == 16384); + using sample_t = cuda::std::uint8_t; using counter_t = int; using level_t = int; diff --git a/cub/test/catch2_test_device_histogram_env.cu b/cub/test/catch2_test_device_histogram_env.cu index 6712bc678055..1fdb04e8fe78 100644 --- a/cub/test/catch2_test_device_histogram_env.cu +++ b/cub/test/catch2_test_device_histogram_env.cu @@ -739,6 +739,68 @@ CUB_TEST("DeviceHistogram::MultiHistogramEven uses environment", "[histogram][de REQUIRE(d_histogram_b == expected_b); } +CUB_TEST("DeviceHistogram::MultiHistogramEven handles the device-launch dynamic-SMEM boundary", + "[histogram][device]", + CUB_SMALL) +{ + int current_device{}; + REQUIRE(cudaSuccess == cudaGetDevice(¤t_device)); + + cuda::compute_capability cc{}; + REQUIRE(cudaSuccess == cub::detail::ptx_compute_cap(cc, current_device)); + if (cc < cuda::compute_capability{10, 0}) + { + SKIP("The runtime-sized shared-memory histogram policy is currently tuned for SM100"); + } + + [[maybe_unused]] constexpr int num_channels = 4; + constexpr int num_active_channels = 3; + // The direct-load kernel has no static shared-memory footprint, so 4,096 + // three-channel counters exactly fill the B200's 48 KiB device-launch limit. + constexpr int num_bins = 4096; + constexpr int num_levels = num_bins + 1; + auto d_samples = c2h::device_vector{0, 1, 2, 3}; + + cuda::std::array levels{num_levels, num_levels, num_levels}; + cuda::std::array lower_levels{0, 0, 0}; + cuda::std::array upper_levels{num_bins, num_bins, num_bins}; + + auto d_histogram_r = c2h::device_vector(num_bins, 0); + auto d_histogram_g = c2h::device_vector(num_bins, 0); + auto d_histogram_b = c2h::device_vector(num_bins, 0); + cuda::std::array d_histograms = { + thrust::raw_pointer_cast(d_histogram_r.data()), + thrust::raw_pointer_cast(d_histogram_g.data()), + thrust::raw_pointer_cast(d_histogram_b.data())}; + + size_t expected_bytes_allocated{}; + REQUIRE( + cudaSuccess + == cub::DeviceHistogram::MultiHistogramEven( + nullptr, + expected_bytes_allocated, + thrust::raw_pointer_cast(d_samples.data()), + d_histograms, + levels, + lower_levels, + upper_levels, + 1)); + + auto env = stdexec::env{expected_allocation_size(expected_bytes_allocated)}; + multi_histogram_even( + thrust::raw_pointer_cast(d_samples.data()), d_histograms, levels, lower_levels, upper_levels, 1, env); + + auto expected_r = c2h::device_vector(num_bins, 0); + auto expected_g = c2h::device_vector(num_bins, 0); + auto expected_b = c2h::device_vector(num_bins, 0); + expected_r[0] = 1; + expected_g[1] = 1; + expected_b[2] = 1; + REQUIRE(d_histogram_r == expected_r); + REQUIRE(d_histogram_g == expected_g); + REQUIRE(d_histogram_b == expected_b); +} + CUB_TEST_CASE("DeviceHistogram::MultiHistogramEven uses custom stream", "[histogram][device]", CUB_SMALL) { [[maybe_unused]] constexpr int NUM_CHANNELS = 4; @@ -1626,10 +1688,79 @@ struct histogram_tuning { _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::HistogramPolicy { - return {BlockThreads, 1, 1, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT, false, cub::SMEM, false, 0}; + constexpr auto sweep = + cub::HistogramPrivatizationPolicy{BlockThreads, 1, 1, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT, false, false}; + return {sweep, sweep, sweep, 256, 256 * sizeof(unsigned int), 0, 0, 0, 0, 0, 0, 0, 0}; } }; +template +struct histogram_tuning_with_local_counter : histogram_tuning +{ + using local_counter_type = LocalCounterT; +}; + +struct mixed_counter_histogram_tuning +{ + using local_counter_type = unsigned int; + + _CCCL_API constexpr auto operator()(cuda::compute_capability) const -> cub::HistogramPolicy + { + constexpr auto sweep = + cub::HistogramPrivatizationPolicy{128, 4, 1, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT, false, false}; + return { + sweep, + sweep, + sweep, + 256, + 512 * sizeof(unsigned int), + 228352, + 0, + 2048 * sizeof(unsigned int), + 28544 * sizeof(unsigned int) * 2, + 19029 * sizeof(unsigned int) * 3, + 8192 * sizeof(unsigned int) * 4, + 0, + 0}; + } +}; + +static_assert( + cuda::std::is_same_v< + cub::detail::histogram::local_counter_t, unsigned long long>, + unsigned int>); +static_assert(cuda::std::is_same_v, unsigned long long>, + unsigned long long>); + +CUB_TEST("DeviceHistogram supports narrower local counters than output counters", "[histogram][device]", CUB_SMALL) +{ + int current_device{}; + REQUIRE(cudaSuccess == cudaGetDevice(¤t_device)); + + cuda::compute_capability cc{}; + REQUIRE(cudaSuccess == cub::detail::ptx_compute_cap(cc, current_device)); + if (cc < cuda::compute_capability{10, 0}) + { + SKIP("The runtime-sized shared-memory histogram policy is currently tuned for SM100"); + } + + constexpr int num_samples = 4096; + constexpr int num_levels = num_samples + 1; + auto d_histogram = c2h::device_vector(num_samples, 0); + auto env = cuda::execution::tune(mixed_counter_histogram_tuning{}); + + histogram_even( + cuda::counting_iterator(0), + thrust::raw_pointer_cast(d_histogram.data()), + num_levels, + 0u, + static_cast(num_samples), + num_samples, + env); + + REQUIRE(d_histogram == c2h::host_vector(num_samples, 1)); +} + using block_sizes = c2h::type_list, cuda::std::integral_constant>; @@ -1750,20 +1881,54 @@ CUB_TEST("Test HistogramPolicy properties", "[histogram][device]", CUB_SMALL) // aggregate init constexpr auto p1 = cub::HistogramPolicy{ - 128, 7, 4, cub::BLOCK_LOAD_DIRECT, cub::CacheLoadModifier::LOAD_LDG, false, cub::SMEM, false, 2048}; + {128, 7, 4, cub::BLOCK_LOAD_DIRECT, cub::CacheLoadModifier::LOAD_LDG, false, false}, + {96, 3, 4, cub::BLOCK_LOAD_DIRECT, cub::CacheLoadModifier::LOAD_LDG, false, false}, + {128, 7, 4, cub::BLOCK_LOAD_DIRECT, cub::CacheLoadModifier::LOAD_LDG, false, false}, + 256, + 2052, + 12345, + 2, + 1024, + 4096, + 8192, + 16384, + 32768, + 2048}; # if _CCCL_STD_VER >= 2020 // designated init constexpr auto p2 = cub::HistogramPolicy{ - .threads_per_block = 128, - .pixels_per_thread = 7, - .vec_size = 4, - .load_algorithm = cub::BLOCK_LOAD_DIRECT, - .load_modifier = cub::CacheLoadModifier::LOAD_LDG, - .rle_compress = false, - .mem_preference = cub::SMEM, - .use_work_stealing = false, - .init_kernel_pdl_trigger_max_bins = 2048}; + .gmem = {.threads_per_block = 128, + .items_per_thread = 7, + .vec_size = 4, + .load_algorithm = cub::BLOCK_LOAD_DIRECT, + .load_modifier = cub::CacheLoadModifier::LOAD_LDG, + .rle_compress = false, + .work_stealing = false}, + .static_smem = {.threads_per_block = 96, + .items_per_thread = 3, + .vec_size = 4, + .load_algorithm = cub::BLOCK_LOAD_DIRECT, + .load_modifier = cub::CacheLoadModifier::LOAD_LDG, + .rle_compress = false, + .work_stealing = false}, + .dynamic_smem = {.threads_per_block = 128, + .items_per_thread = 7, + .vec_size = 4, + .load_algorithm = cub::BLOCK_LOAD_DIRECT, + .load_modifier = cub::CacheLoadModifier::LOAD_LDG, + .rle_compress = false, + .work_stealing = false}, + .init_threads_per_block = 256, + .max_privatized_static_smem_single_channel_bytes = 2052, + .max_privatized_dynamic_smem_single_channel_bytes = 12345, + .static_smem_min_blocks_per_sm = 2, + .max_privatized_dynamic_smem_multi_channel_range_bytes = 1024, + .max_privatized_dynamic_smem_2_channel_even_bytes = 4096, + .max_privatized_dynamic_smem_3_channel_even_bytes = 8192, + .max_privatized_dynamic_smem_4_channel_even_bytes = 16384, + .min_cached_search_gmem_range_bins = 32768, + .max_output_histogram_bytes_for_init_kernel_pdl = 2048}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 @@ -1777,9 +1942,201 @@ CUB_TEST("Test HistogramPolicy properties", "[histogram][device]", CUB_SMALL) os << p; return os.str(); }; - REQUIRE(to_string(p1) - == "HistogramPolicy { .threads_per_block = 128, .pixels_per_thread = 7, .vec_size = 4" - ", .load_algorithm = BLOCK_LOAD_DIRECT, .load_modifier = LOAD_LDG, .rle_compress = 0" - ", .mem_preference = SMEM, .use_work_stealing = 0, .init_kernel_pdl_trigger_max_bins = 2048 }"); + REQUIRE(to_string(p1) == to_string(p2)); +} + +CUB_TEST("Histogram architecture policies carry their dynamic shared-memory budgets", "[histogram][device]", CUB_SMALL) +{ + using selector_t = cub::detail::histogram::policy_selector_from_types; + + constexpr auto sm90_policy = selector_t{}(cuda::compute_capability{9, 0}); + constexpr auto sm100_policy = selector_t{}(cuda::compute_capability{10, 0}); + constexpr auto sm120_policy = selector_t{}(cuda::compute_capability{12, 0}); + constexpr auto sm90_range_u32_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{9, 0}); + constexpr auto sm100_range_u32_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_range_f64_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_range_wide_counter_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_wide_counter_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr int expected_single_channel_policy_bytes = 228352; + constexpr int expected_single_channel_limit_bytes = expected_single_channel_policy_bytes; + + STATIC_REQUIRE(cub::detail::histogram::max_privatized_smem_bins( + sm90_policy.max_privatized_static_smem_single_channel_bytes) + == 256); + STATIC_REQUIRE(cub::detail::histogram::max_privatized_smem_bins( + sm100_policy.max_privatized_static_smem_single_channel_bytes) + == 256); + STATIC_REQUIRE(sm90_policy.max_privatized_dynamic_smem_single_channel_bytes == 0); + STATIC_REQUIRE(sm100_policy.max_privatized_dynamic_smem_single_channel_bytes == expected_single_channel_policy_bytes); + STATIC_REQUIRE( + cub::detail::histogram::dynamic_smem_limit_bytes(sm100_policy) == expected_single_channel_limit_bytes); + STATIC_REQUIRE(sm100_wide_counter_policy.max_privatized_dynamic_smem_single_channel_bytes == 0); + STATIC_REQUIRE(cub::detail::histogram::max_privatized_smem_bins( + sm100_policy.max_privatized_dynamic_smem_single_channel_bytes) + == expected_single_channel_policy_bytes / int{sizeof(unsigned int)}); + STATIC_REQUIRE(sm100_policy.max_privatized_dynamic_smem_multi_channel_range_bytes == 8192); + STATIC_REQUIRE(sm100_policy.gmem.threads_per_block == 768); + STATIC_REQUIRE(sm100_policy.gmem.items_per_thread == 12); + STATIC_REQUIRE(sm100_policy.static_smem == sm100_policy.gmem); + + constexpr auto sm100_range_u64_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + STATIC_REQUIRE(sm100_range_u64_policy.gmem.threads_per_block == 768); + STATIC_REQUIRE(sm100_range_u64_policy.gmem.items_per_thread == 6); + STATIC_REQUIRE(sm100_range_u64_policy.static_smem.threads_per_block == 384); + STATIC_REQUIRE(sm100_range_u64_policy.static_smem.items_per_thread == 8); + STATIC_REQUIRE(sm100_range_u64_policy.static_smem_min_blocks_per_sm == 3); + STATIC_REQUIRE(sm90_range_u32_policy.min_cached_search_gmem_range_bins == 0); + STATIC_REQUIRE(sm100_range_u32_policy.min_cached_search_gmem_range_bins == 1); + STATIC_REQUIRE(sm100_range_f64_policy.min_cached_search_gmem_range_bins == 1); + STATIC_REQUIRE(sm100_range_wide_counter_policy.min_cached_search_gmem_range_bins == 0); + STATIC_REQUIRE(sm100_range_u64_policy.min_cached_search_gmem_range_bins == 1); + STATIC_REQUIRE(cub::detail::histogram::use_cached_search_for_gmem_range(sm100_range_u32_policy, 1)); + STATIC_REQUIRE(cub::detail::histogram::max_privatized_smem_bins( + sm100_range_u64_policy.max_privatized_static_smem_single_channel_bytes) + == 256); + + constexpr auto sm100_multi_range_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_multi_range_f64_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_even_2ch_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_even_3ch_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm100_even_4ch_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{10, 0}); + constexpr auto sm120_even_2ch_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{12, 0}); + constexpr auto sm120_even_3ch_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{12, 0}); + constexpr auto sm120_even_4ch_policy = + cub::detail::histogram::policy_selector_from_types{}( + cuda::compute_capability{12, 0}); + STATIC_REQUIRE(sm100_multi_range_policy.min_cached_search_gmem_range_bins == 16384); + STATIC_REQUIRE(sm100_multi_range_f64_policy.min_cached_search_gmem_range_bins == 16384); + STATIC_REQUIRE_FALSE(cub::detail::histogram::use_cached_search_for_gmem_range(sm100_multi_range_policy, 16383)); + STATIC_REQUIRE(cub::detail::histogram::use_cached_search_for_gmem_range(sm100_multi_range_policy, 16384)); + STATIC_REQUIRE(sm100_even_4ch_policy.min_cached_search_gmem_range_bins == 0); + constexpr int expected_even_2ch_policy_bytes = 65536; + constexpr int expected_even_3ch_policy_bytes = 98304; + constexpr int expected_even_4ch_policy_bytes = 131072; + constexpr int expected_even_2ch_limit_bytes = expected_even_2ch_policy_bytes; + constexpr int expected_even_3ch_limit_bytes = expected_even_3ch_policy_bytes; + constexpr int expected_even_4ch_limit_bytes = expected_even_4ch_policy_bytes; + STATIC_REQUIRE( + sm100_even_2ch_policy.max_privatized_dynamic_smem_2_channel_even_bytes == expected_even_2ch_policy_bytes); + STATIC_REQUIRE( + sm100_even_3ch_policy.max_privatized_dynamic_smem_3_channel_even_bytes == expected_even_3ch_policy_bytes); + STATIC_REQUIRE( + sm100_even_4ch_policy.max_privatized_dynamic_smem_4_channel_even_bytes == expected_even_4ch_policy_bytes); + STATIC_REQUIRE( + cub::detail::histogram::dynamic_smem_limit_bytes(sm100_even_2ch_policy) == expected_even_2ch_limit_bytes); + STATIC_REQUIRE( + cub::detail::histogram::dynamic_smem_limit_bytes(sm100_even_3ch_policy) == expected_even_3ch_limit_bytes); + STATIC_REQUIRE( + cub::detail::histogram::dynamic_smem_limit_bytes(sm100_even_4ch_policy) == expected_even_4ch_limit_bytes); + constexpr int expected_sm120_dynamic_smem_bytes = 99 * 1024; + STATIC_REQUIRE(sm120_policy.max_privatized_dynamic_smem_single_channel_bytes == expected_sm120_dynamic_smem_bytes); + STATIC_REQUIRE(sm120_even_2ch_policy.max_privatized_dynamic_smem_2_channel_even_bytes == 65536); + STATIC_REQUIRE(sm120_even_3ch_policy.max_privatized_dynamic_smem_3_channel_even_bytes == 98304); + STATIC_REQUIRE( + sm120_even_4ch_policy.max_privatized_dynamic_smem_4_channel_even_bytes == expected_sm120_dynamic_smem_bytes); + STATIC_REQUIRE(sm100_multi_range_policy.gmem.threads_per_block == 384); + STATIC_REQUIRE(sm100_multi_range_policy.gmem.items_per_thread == 5); + STATIC_REQUIRE(sm100_multi_range_policy.static_smem.threads_per_block == 384); + STATIC_REQUIRE(sm100_multi_range_policy.static_smem.items_per_thread == 5); + STATIC_REQUIRE(sm100_multi_range_policy.static_smem_min_blocks_per_sm == 3); + STATIC_REQUIRE(cub::detail::histogram::max_privatized_smem_bins( + sm100_multi_range_policy.max_privatized_static_smem_single_channel_bytes) + == 256); + + using cub::detail::histogram::privatization_mode; + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_policy, 256) + == privatization_mode::static_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_policy, 257) + == privatization_mode::dynamic_smem); + constexpr int sm100_single_channel_max_dynamic_bins = expected_single_channel_limit_bytes / int{sizeof(unsigned int)}; + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_policy, sm100_single_channel_max_dynamic_bins) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_policy, sm100_single_channel_max_dynamic_bins + 1) + == privatization_mode::gmem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_range_u64_policy, 256) + == privatization_mode::static_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_range_u64_policy, 257) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE( + cub::detail::histogram::select_privatization_mode(sm100_multi_range_policy, 256) + == privatization_mode::static_smem); + STATIC_REQUIRE( + cub::detail::histogram::select_privatization_mode(sm100_multi_range_policy, 257) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE( + cub::detail::histogram::select_privatization_mode(sm100_multi_range_policy, 2048) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE( + cub::detail::histogram::select_privatization_mode(sm100_multi_range_policy, 2049) + == privatization_mode::gmem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_even_4ch_policy, 256) + == privatization_mode::static_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_even_4ch_policy, 257) + == privatization_mode::dynamic_smem); + constexpr int sm100_even_4ch_max_dynamic_bins = expected_even_4ch_limit_bytes / sizeof(unsigned int) / 4; + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_even_4ch_policy, sm100_even_4ch_max_dynamic_bins) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_even_4ch_policy, sm100_even_4ch_max_dynamic_bins + 1) + == privatization_mode::gmem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_even_2ch_policy, 257) + == privatization_mode::dynamic_smem); + constexpr int sm100_even_2ch_max_dynamic_bins = expected_even_2ch_limit_bytes / sizeof(unsigned int) / 2; + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_even_2ch_policy, sm100_even_2ch_max_dynamic_bins) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_even_2ch_policy, sm100_even_2ch_max_dynamic_bins + 1) + == privatization_mode::gmem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode(sm100_even_3ch_policy, 257) + == privatization_mode::dynamic_smem); + constexpr int sm100_even_3ch_max_dynamic_bins = expected_even_3ch_limit_bytes / sizeof(unsigned int) / 3; + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_even_3ch_policy, sm100_even_3ch_max_dynamic_bins) + == privatization_mode::dynamic_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode( + sm100_even_3ch_policy, sm100_even_3ch_max_dynamic_bins + 1) + == privatization_mode::gmem); + STATIC_REQUIRE( + cub::detail::histogram::select_privatization_mode(sm100_wide_counter_policy, 128) + == privatization_mode::static_smem); + STATIC_REQUIRE( + cub::detail::histogram::select_privatization_mode(sm100_wide_counter_policy, 129) + == privatization_mode::gmem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode_for_counter_size( + sm100_policy, 128, sizeof(cuda::std::uint64_t)) + == privatization_mode::static_smem); + STATIC_REQUIRE(cub::detail::histogram::select_privatization_mode_for_counter_size( + sm100_policy, 28545, sizeof(cuda::std::uint64_t)) + == privatization_mode::gmem); } #endif // _CCCL_COMPILER(GCC, >=, 8) diff --git a/cub/test/catch2_test_device_histogram_env_api.cu b/cub/test/catch2_test_device_histogram_env_api.cu index c79fa6ce12a9..c0759791cfa5 100644 --- a/cub/test/catch2_test_device_histogram_env_api.cu +++ b/cub/test/catch2_test_device_histogram_env_api.cu @@ -418,15 +418,22 @@ struct HistogramPolicySelector { __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::HistogramPolicy { - return {.threads_per_block = 128, - .pixels_per_thread = cc > cuda::compute_capability{9, 0} ? 16 : 7, - .vec_size = 4, - .load_algorithm = cub::BLOCK_LOAD_DIRECT, - .load_modifier = cub::LOAD_LDG, - .rle_compress = false, - .mem_preference = cub::SMEM, - .use_work_stealing = false, - .init_kernel_pdl_trigger_max_bins = 2048}; + const auto sweep = cub::HistogramPrivatizationPolicy{ + 128, cc > cuda::compute_capability{9, 0} ? 16 : 7, 4, cub::BLOCK_LOAD_DIRECT, cub::LOAD_LDG, false, false}; + return { + .gmem = sweep, + .static_smem = sweep, + .dynamic_smem = sweep, + .init_threads_per_block = 256, + .max_privatized_static_smem_single_channel_bytes = 256 * sizeof(unsigned int), + .max_privatized_dynamic_smem_single_channel_bytes = 0, + .static_smem_min_blocks_per_sm = 0, + .max_privatized_dynamic_smem_multi_channel_range_bytes = 0, + .max_privatized_dynamic_smem_2_channel_even_bytes = 0, + .max_privatized_dynamic_smem_3_channel_even_bytes = 0, + .max_privatized_dynamic_smem_4_channel_even_bytes = 0, + .min_cached_search_gmem_range_bins = 0, + .max_output_histogram_bytes_for_init_kernel_pdl = 8192}; } }; // example-end histogram-even-policy-selector diff --git a/cub/test/catch2_test_enum_formatting.cu b/cub/test/catch2_test_enum_formatting.cu index 87c17e311729..e75d3d27eefc 100644 --- a/cub/test/catch2_test_enum_formatting.cu +++ b/cub/test/catch2_test_enum_formatting.cu @@ -45,15 +45,6 @@ struct FormatTester template void do_test(const Tester& tester) { - // BlockHistogramMemoryPreference - { - tester(cub::BlockHistogramMemoryPreference::GMEM, "GMEM"); - tester(cub::BlockHistogramMemoryPreference::SMEM, "SMEM"); - tester(cub::BlockHistogramMemoryPreference::BLEND, "BLEND"); - - tester(cub::BlockHistogramMemoryPreference(100), ""); - } - // RadixSortStoreAlgorithm { tester(cub::RadixSortStoreAlgorithm::RADIX_SORT_STORE_DIRECT, "RADIX_SORT_STORE_DIRECT");