From e27b31071b54069b8d90084f2daadcc72221ec00 Mon Sep 17 00:00:00 2001 From: Federico Busato <50413820+fbusato@users.noreply.github.com> Date: Tue, 4 Feb 2025 17:32:32 -0800 Subject: [PATCH 1/3] Internalize cub::KernelConfig (#3683) (cherry picked from commit 32d05c747fa72a234dfe5816bbca55fbbd0f956a) --- .../device/dispatch/dispatch_radix_sort.cuh | 8 +++---- cub/cub/device/dispatch/dispatch_reduce.cuh | 4 ++-- cub/cub/util_device.cuh | 21 ++++++++----------- 3 files changed, 15 insertions(+), 18 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 1ed25d21796..32284137c46 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -1184,11 +1184,11 @@ struct DispatchRadixSort struct PassConfig { UpsweepKernelT upsweep_kernel; - KernelConfig upsweep_config; + detail::KernelConfig upsweep_config; ScanKernelT scan_kernel; - KernelConfig scan_config; + detail::KernelConfig scan_config; DownsweepKernelT downsweep_kernel; - KernelConfig downsweep_config; + detail::KernelConfig downsweep_config; int radix_bits; int radix_digits; int max_downsweep_grid_size; @@ -2135,7 +2135,7 @@ struct DispatchSegmentedRadixSort struct PassConfig { SegmentedKernelT segmented_kernel; - KernelConfig segmented_config; + detail::KernelConfig segmented_config; int radix_bits; int radix_digits; diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index 170237d5646..9b0835198b5 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -474,7 +474,7 @@ struct DispatchReduce } // Init regular kernel configuration - KernelConfig reduce_config; + detail::KernelConfig reduce_config; error = CubDebug(reduce_config.Init(reduce_kernel, active_policy.Reduce(), launcher_factory)); if (cudaSuccess != error) { @@ -949,7 +949,7 @@ struct DispatchSegmentedReduce } // Init kernel configuration - KernelConfig segmented_reduce_config; + [[maybe_unused]] detail::KernelConfig segmented_reduce_config; error = CubDebug(segmented_reduce_config.Init(segmented_reduce_kernel)); if (cudaSuccess != error) diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index ca18aa24fb0..6150b1115fd 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -637,25 +637,18 @@ CUB_RUNTIME_FUNCTION PolicyWrapper MakePolicyWrapper(PolicyT policy) namespace detail { + struct TripleChevronFactory; -} /** * Kernel dispatch configuration */ struct KernelConfig { - int block_threads; - int items_per_thread; - int tile_size; - int sm_occupancy; - - CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE KernelConfig() - : block_threads(0) - , items_per_thread(0) - , tile_size(0) - , sm_occupancy(0) - {} + int block_threads{0}; + int items_per_thread{0}; + int tile_size{0}; + int sm_occupancy{0}; template CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t @@ -668,6 +661,10 @@ struct KernelConfig } }; +} // namespace detail + +using KernelConfig CCCL_DEPRECATED_BECAUSE("This class is considered an implementation detail") = detail::KernelConfig; + /// Helper for dispatching into a policy chain template struct ChainedPolicy From a826d5fadfdd3af9cd40f7071512e12530411542 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 5 Feb 2025 09:16:48 +0100 Subject: [PATCH 2/3] Use plain old void cast to suppress unused warnings --- cub/cub/device/dispatch/dispatch_reduce.cuh | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index 9b0835198b5..909ad4f5999 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -475,6 +475,7 @@ struct DispatchReduce // Init regular kernel configuration detail::KernelConfig reduce_config; + (void) reduce_config; error = CubDebug(reduce_config.Init(reduce_kernel, active_policy.Reduce(), launcher_factory)); if (cudaSuccess != error) { @@ -949,7 +950,8 @@ struct DispatchSegmentedReduce } // Init kernel configuration - [[maybe_unused]] detail::KernelConfig segmented_reduce_config; + detail::KernelConfig segmented_reduce_config; + (void) segmented_reduce_config; error = CubDebug(segmented_reduce_config.Init(segmented_reduce_kernel)); if (cudaSuccess != error) From 14a670a635e3b59092479d163b3863fa9d6e0d2b Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 5 Feb 2025 12:53:58 +0100 Subject: [PATCH 3/3] Suppress unused variable warning for nvcc 11.1 --- cub/cub/device/dispatch/dispatch_reduce.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index 909ad4f5999..aa20fd238e7 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -951,7 +951,7 @@ struct DispatchSegmentedReduce // Init kernel configuration detail::KernelConfig segmented_reduce_config; - (void) segmented_reduce_config; + (void) &segmented_reduce_config; error = CubDebug(segmented_reduce_config.Init(segmented_reduce_kernel)); if (cudaSuccess != error)