From 0c57bab1d1195bc8b8a968c9896dc0f9a2dd11c9 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 28 Oct 2025 18:24:02 +0100 Subject: [PATCH 1/7] Add PDL to DeviceHistogram --- cub/cub/agent/agent_histogram.cuh | 4 ++++ cub/cub/device/dispatch/dispatch_histogram.cuh | 4 ++-- cub/cub/device/dispatch/kernels/histogram.cuh | 3 +++ 3 files changed, 9 insertions(+), 2 deletions(-) diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index b5116d1d8dd0..80e4d0b525f6 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -628,6 +628,8 @@ struct AgentHistogram ((size_t(d_native_samples) & pixel_mask) == 0) && // ptr is pixel-aligned ((row_bytes & pixel_mask) == 0); // number of row-samples is a multiple of the alignment of the pixel + _CCCL_PDL_GRID_DEPENDENCY_SYNC(); + // Whether rows are aligned and can be vectorized if ((d_native_samples != nullptr) && (vec_aligned_rows || pixel_aligned_rows)) { @@ -639,6 +641,8 @@ struct AgentHistogram ConsumeTiles( num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, bool_constant_v); } + + _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); } //! Initialize privatized bin counters. Specialized for privatized shared-memory counters diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index f51067023ee2..df82a6d8f4c0 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -227,7 +227,7 @@ struct dispatch_histogram #endif // CUB_DEBUG_LOG // Invoke histogram_init_kernel - launcher_factory(histogram_init_grid_dims, histogram_init_block_threads, 0, stream) + launcher_factory(histogram_init_grid_dims, histogram_init_block_threads, 0, stream, true) .doit(histogram_init_kernel, num_output_bins_wrapper, d_output_histograms, tile_queue); // Return if empty problem @@ -250,7 +250,7 @@ struct dispatch_histogram #endif // CUB_DEBUG_LOG // Invoke histogram_sweep_kernel - launcher_factory(sweep_grid_dims, block_threads, 0, stream) + launcher_factory(sweep_grid_dims, block_threads, 0, stream, true) .doit(histogram_sweep_kernel, d_samples, num_output_bins_wrapper, diff --git a/cub/cub/device/dispatch/kernels/histogram.cuh b/cub/cub/device/dispatch/kernels/histogram.cuh index da387ee6d679..1ccaf06a62eb 100644 --- a/cub/cub/device/dispatch/kernels/histogram.cuh +++ b/cub/cub/device/dispatch/kernels/histogram.cuh @@ -334,6 +334,9 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( ::cuda::std::array d_output_histograms_wrapper, GridQueue tile_queue) { + _CCCL_PDL_GRID_DEPENDENCY_SYNC(); + _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); + if ((threadIdx.x == 0) && (blockIdx.x == 0)) { tile_queue.ResetDrain(); From 6ca8be8001025ad3777a7871fb3f9d3cd15c91d5 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 28 Oct 2025 18:42:58 +0100 Subject: [PATCH 2/7] threshold --- cub/cub/agent/agent_histogram.cuh | 2 +- cub/cub/device/dispatch/kernels/histogram.cuh | 11 +++++++++-- 2 files changed, 10 insertions(+), 3 deletions(-) diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index 80e4d0b525f6..d49aecdc75d3 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -642,7 +642,7 @@ struct AgentHistogram num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, bool_constant_v); } - _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); + _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); // omitting makes no difference in cub.bench.histogram.even.base } //! Initialize privatized bin counters. Specialized for privatized shared-memory counters diff --git a/cub/cub/device/dispatch/kernels/histogram.cuh b/cub/cub/device/dispatch/kernels/histogram.cuh index 1ccaf06a62eb..3c8a9a38bd73 100644 --- a/cub/cub/device/dispatch/kernels/histogram.cuh +++ b/cub/cub/device/dispatch/kernels/histogram.cuh @@ -16,6 +16,8 @@ #include #include +#include + CUB_NAMESPACE_BEGIN namespace detail::histogram { @@ -334,8 +336,13 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( ::cuda::std::array d_output_histograms_wrapper, GridQueue tile_queue) { - _CCCL_PDL_GRID_DEPENDENCY_SYNC(); - _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); + _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 + + if (::cuda::std::reduce(num_output_bins_wrapper.begin(), num_output_bins_wrapper.end()) <= 2048) + { + _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); + } if ((threadIdx.x == 0) && (blockIdx.x == 0)) { From a18bc0782d969cf801da5cf15cbf04407b58ee6e Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 28 Oct 2025 18:55:52 +0100 Subject: [PATCH 3/7] Put trigger threshold into tuning policy --- cub/cub/device/dispatch/dispatch_histogram.cuh | 11 ++++++++--- cub/cub/device/dispatch/kernels/histogram.cuh | 12 +++++++----- cub/cub/device/dispatch/tuning/tuning_histogram.cuh | 4 ++++ 3 files changed, 19 insertions(+), 8 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index df82a6d8f4c0..d3fc11e29e64 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -54,10 +54,15 @@ namespace detail::histogram template struct DeviceHistogramKernelSource { - CUB_DEFINE_KERNEL_GETTER(HistogramInitKernel, DeviceHistogramInitKernel); - // We define this differently than the other kernel getters because there are // different dispatch paths which affect which policy is used and the decode operators. + + template + _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramInitKernel() + { + return &DeviceHistogramInitKernel; + } + template _CCCL_HIDE_FROM_ABI CUB_RUNTIME_FUNCTION static constexpr auto HistogramSweepKernel() { @@ -287,7 +292,7 @@ struct dispatch_histogram CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t Invoke(ActivePolicyT active_policy = {}) { return Invoke( - kernel_source.HistogramInitKernel(), + kernel_source.template HistogramInitKernel(), kernel_source .template HistogramSweepKernel(), active_policy); diff --git a/cub/cub/device/dispatch/kernels/histogram.cuh b/cub/cub/device/dispatch/kernels/histogram.cuh index 3c8a9a38bd73..10d627bf39b5 100644 --- a/cub/cub/device/dispatch/kernels/histogram.cuh +++ b/cub/cub/device/dispatch/kernels/histogram.cuh @@ -330,7 +330,7 @@ struct Transforms //! //! @param tile_queue //! Drain queue descriptor for dynamically mapping tile data onto thread blocks -template +template CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( ::cuda::std::array num_output_bins_wrapper, ::cuda::std::array d_output_histograms_wrapper, @@ -339,10 +339,12 @@ CUB_DETAIL_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 - if (::cuda::std::reduce(num_output_bins_wrapper.begin(), num_output_bins_wrapper.end()) <= 2048) - { - _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); - } + // 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()) + <= ChainedPolicyT::ActivePolicy::pld_trigger_next_launch_in_init_kernel_max_bin_count) { + _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); + })); if ((threadIdx.x == 0) && (blockIdx.x == 0)) { diff --git a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh index bd49dd9b58d1..56854c424918 100644 --- a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh @@ -219,6 +219,8 @@ struct policy_hub using AgentHistogramPolicyT = decltype(select_agent_policy< sm90_tuning()>>(0)); + + static constexpr int pld_trigger_next_launch_in_init_kernel_max_bin_count = 2048; }; struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900> @@ -242,6 +244,8 @@ struct policy_hub decltype(select_agent_policy< sm100_tuning()>>( 0)); + + static constexpr int pld_trigger_next_launch_in_init_kernel_max_bin_count = 2048; }; using MaxPolicy = Policy1000; From 2c68bf41f7dfde76770350648d8e1557583d56b9 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Sat, 1 Nov 2025 23:24:03 +0100 Subject: [PATCH 4/7] Fix --- c/parallel/src/histogram.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/c/parallel/src/histogram.cu b/c/parallel/src/histogram.cu index 35cb18c8ecca..656d44e72ef5 100644 --- a/c/parallel/src/histogram.cu +++ b/c/parallel/src/histogram.cu @@ -75,6 +75,7 @@ struct histogram_kernel_source { cccl_device_histogram_build_result_t& build; + template CUkernel HistogramInitKernel() const { return build.init_kernel; From 9199f0890825270e8565c45b6c1410fa7852398f Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 3 Nov 2025 13:15:59 +0100 Subject: [PATCH 5/7] Add missing include --- cub/cub/agent/agent_histogram.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index d49aecdc75d3..3855d5656efd 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -23,6 +23,7 @@ #include #include +#include #include #include #include From c04fcaa39816d776be527c56dea16fb7cd97bca4 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 3 Nov 2025 13:34:54 +0100 Subject: [PATCH 6/7] Duplicate tuning into cccl c --- c/parallel/src/histogram.cu | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/c/parallel/src/histogram.cu b/c/parallel/src/histogram.cu index 656d44e72ef5..4a52d1ced130 100644 --- a/c/parallel/src/histogram.cu +++ b/c/parallel/src/histogram.cu @@ -103,10 +103,15 @@ histogram_runtime_tuning_policy get_policy(int /*cc*/, cccl_type_info sample_t, return {384, pixels_per_thread}; } -std::string get_init_kernel_name(int num_active_channels, std::string_view counter_t, std::string_view offset_t) +std::string get_init_kernel_name( + std::string_view chained_policy_t, int num_active_channels, std::string_view counter_t, std::string_view offset_t) { return std::format( - "cub::detail::histogram::DeviceHistogramInitKernel<{0}, {1}, {2}>", num_active_channels, counter_t, offset_t); + "cub::detail::histogram::DeviceHistogramInitKernel<{0}, {1}, {2}, {3}>", + chained_policy_t, + num_active_channels, + counter_t, + offset_t); } std::string get_sweep_kernel_name( @@ -226,6 +231,7 @@ struct agent_policy_t {{ struct {5} {{ struct ActivePolicy {{ using AgentHistogramPolicyT = agent_policy_t; + static constexpr int pld_trigger_next_launch_in_init_kernel_max_bin_count = 2048; }}; }}; )XXX"; @@ -254,7 +260,8 @@ struct {5} {{ const bool is_byte_sample = d_samples.value_type.size == 1; - std::string init_kernel_name = histogram::get_init_kernel_name(num_active_channels, counter_cpp, offset_cpp); + std::string init_kernel_name = + histogram::get_init_kernel_name(chained_policy_t, num_active_channels, counter_cpp, offset_cpp); std::string sweep_kernel_name = histogram::get_sweep_kernel_name( chained_policy_t, privatized_smem_bins, From 874c44a0157b2c71194440408b32eca9b42d6674 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 3 Nov 2025 16:03:15 +0100 Subject: [PATCH 7/7] can't even spell --- c/parallel/src/histogram.cu | 2 +- cub/cub/device/dispatch/kernels/histogram.cuh | 2 +- cub/cub/device/dispatch/tuning/tuning_histogram.cuh | 4 ++-- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/c/parallel/src/histogram.cu b/c/parallel/src/histogram.cu index 4a52d1ced130..1fc20ac6d50d 100644 --- a/c/parallel/src/histogram.cu +++ b/c/parallel/src/histogram.cu @@ -231,7 +231,7 @@ struct agent_policy_t {{ struct {5} {{ struct ActivePolicy {{ using AgentHistogramPolicyT = agent_policy_t; - static constexpr int pld_trigger_next_launch_in_init_kernel_max_bin_count = 2048; + static constexpr int pdl_trigger_next_launch_in_init_kernel_max_bin_count = 2048; }}; }}; )XXX"; diff --git a/cub/cub/device/dispatch/kernels/histogram.cuh b/cub/cub/device/dispatch/kernels/histogram.cuh index 10d627bf39b5..e4520813221a 100644 --- a/cub/cub/device/dispatch/kernels/histogram.cuh +++ b/cub/cub/device/dispatch/kernels/histogram.cuh @@ -342,7 +342,7 @@ CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceHistogramInitKernel( // 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()) - <= ChainedPolicyT::ActivePolicy::pld_trigger_next_launch_in_init_kernel_max_bin_count) { + <= ChainedPolicyT::ActivePolicy::pdl_trigger_next_launch_in_init_kernel_max_bin_count) { _CCCL_PDL_TRIGGER_NEXT_LAUNCH(); })); diff --git a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh index 56854c424918..2949e73c5517 100644 --- a/cub/cub/device/dispatch/tuning/tuning_histogram.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_histogram.cuh @@ -220,7 +220,7 @@ struct policy_hub decltype(select_agent_policy< sm90_tuning()>>(0)); - static constexpr int pld_trigger_next_launch_in_init_kernel_max_bin_count = 2048; + static constexpr int pdl_trigger_next_launch_in_init_kernel_max_bin_count = 2048; }; struct Policy1000 : ChainedPolicy<1000, Policy1000, Policy900> @@ -245,7 +245,7 @@ struct policy_hub sm100_tuning()>>( 0)); - static constexpr int pld_trigger_next_launch_in_init_kernel_max_bin_count = 2048; + static constexpr int pdl_trigger_next_launch_in_init_kernel_max_bin_count = 2048; }; using MaxPolicy = Policy1000;