Allow public tuning of cub::DeviceAdjacentDifference

cub/benchmarks/bench/adjacent_difference/subtract_left.cu

@@ -12,7 +12,7 @@
 struct policy_selector_t
 {
   [[nodiscard]] _CCCL_HOST_DEVICE constexpr auto operator()(cuda::compute_capability) const
-    -> cub::detail::adjacent_difference::adjacent_difference_policy
+    -> cub::AdjacentDifferencePolicy
   {
     return {TUNE_THREADS_PER_BLOCK,
             TUNE_ITEMS_PER_THREAD,

cub/cub/device/device_adjacent_difference.cuh

@@ -87,6 +87,23 @@ CUB_NAMESPACE_BEGIN
 //!
 //!    // d_values <-- [1, 1, -1, 1, -1, 1, -1, 1]
 //!
+//! @par Tuning
+//! All algorithms in DeviceAdjacentDifference that accept an environment can be tuned by passing a custom
+//! :ref:`policy selector <cub-policy-selectors>` that returns an @ref AdjacentDifferencePolicy, as shown in the
+//! example below:
+//!
+//!  .. literalinclude:: ../../../cub/test/catch2_test_device_adjacent_difference_env_api.cu
+//!      :language: c++
+//!      :dedent:
+//!      :start-after: example-begin subtract-left-copy-policy-selector
+//!      :end-before: example-end subtract-left-copy-policy-selector
+//!
+//!  .. literalinclude:: ../../../cub/test/catch2_test_device_adjacent_difference_env_api.cu
+//!      :language: c++
+//!      :dedent:
+//!      :start-after: example-begin subtract-left-copy-tuning
+//!      :end-before: example-end subtract-left-copy-tuning
+//!
 //! @endrst
 struct DeviceAdjacentDifference
 {

cub/cub/device/dispatch/dispatch_adjacent_difference.cuh

@@ -63,7 +63,7 @@ _CCCL_KERNEL_ATTRIBUTES void DeviceAdjacentDifferenceDifferenceKernel(
   _CCCL_GRID_CONSTANT const OffsetT num_items)
 {
   static_assert(::cuda::std::is_empty_v<PolicySelector>);
-  static constexpr adjacent_difference_policy policy = current_policy<PolicySelector>();
+  static constexpr AdjacentDifferencePolicy policy = current_policy<PolicySelector>();
   using AdjacentDifferencePolicyT =
     AgentAdjacentDifferencePolicy<policy.threads_per_block,
                                   policy.items_per_thread,
@@ -100,10 +100,10 @@ template <typename PolicyHub>
 struct policy_selector_from_hub
 {
   // this is only called in device code, so we can ignore the cc parameter
-  _CCCL_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const -> adjacent_difference_policy
+  _CCCL_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const -> AdjacentDifferencePolicy
   {
     using p = typename PolicyHub::MaxPolicy::ActivePolicy::AdjacentDifferencePolicy;
-    return adjacent_difference_policy{
+    return AdjacentDifferencePolicy{
       p::BLOCK_THREADS, p::ITEMS_PER_THREAD, p::LOAD_ALGORITHM, p::LOAD_MODIFIER, p::STORE_ALGORITHM};
   }
 };
@@ -123,7 +123,7 @@ template <typename InputIteratorT,
           MayAlias AliasOpt,
           ReadOption ReadOpt,
           typename PolicyHub = detail::adjacent_difference::policy_hub<InputIteratorT, AliasOpt == MayAlias::Yes>>
-struct DispatchAdjacentDifference
+struct CCCL_DEPRECATED_BECAUSE("Please use DeviceAdjacentDifference") DispatchAdjacentDifference
 {
   using InputT = detail::it_value_t<InputIteratorT>;
 
@@ -344,7 +344,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE auto dispatch(
     return error;
   }
 
-  const adjacent_difference_policy active_policy = policy_selector(cc);
+  const AdjacentDifferencePolicy active_policy = policy_selector(cc);
 #if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG)
   NV_IF_TARGET(NV_IS_HOST, ({
                  ::std::stringstream ss;

cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh

@@ -22,43 +22,44 @@
 
 CUB_NAMESPACE_BEGIN
 
-namespace detail::adjacent_difference
-{
-struct adjacent_difference_policy
+//! The tuning policy for all algorithms in @ref DeviceAdjacentDifference.
+struct AdjacentDifferencePolicy
 {
-  int threads_per_block;
-  int items_per_thread;
-  BlockLoadAlgorithm load_algorithm;
-  CacheLoadModifier load_modifier;
-  BlockStoreAlgorithm store_algorithm;
-
-  _CCCL_HOST_DEVICE_API constexpr friend bool
-  operator==(const adjacent_difference_policy& lhs, const adjacent_difference_policy& rhs)
+  int threads_per_block; //!< Number of threads in a CUDA block
+  int items_per_thread; //!< Number of items processed per thread
+  BlockLoadAlgorithm load_algorithm; //!< The @ref BlockLoadAlgorithm used for loading items from global memory
+  CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
+  BlockStoreAlgorithm store_algorithm; //!< The @ref BlockStoreAlgorithm used for storing items to global memory
+
+  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  operator==(const AdjacentDifferencePolicy& lhs, const AdjacentDifferencePolicy& rhs)
   {
     return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
         && lhs.load_algorithm == rhs.load_algorithm && lhs.load_modifier == rhs.load_modifier
         && lhs.store_algorithm == rhs.store_algorithm;
   }
 
-  _CCCL_HOST_DEVICE_API constexpr friend bool
-  operator!=(const adjacent_difference_policy& lhs, const adjacent_difference_policy& rhs)
+  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
+  operator!=(const AdjacentDifferencePolicy& lhs, const AdjacentDifferencePolicy& rhs)
   {
     return !(lhs == rhs);
   }
 
 #if _CCCL_HOSTED()
-  friend ::std::ostream& operator<<(::std::ostream& os, const adjacent_difference_policy& p)
+  friend ::std::ostream& operator<<(::std::ostream& os, const AdjacentDifferencePolicy& p)
   {
-    return os << "adjacent_difference_policy { .threads_per_block = " << p.threads_per_block
+    return os << "AdjacentDifferencePolicy { .threads_per_block = " << p.threads_per_block
               << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm
               << ", .load_modifier = " << p.load_modifier << ", .store_algorithm = " << p.store_algorithm << " }";
   }
 #endif // _CCCL_HOSTED()
 };
 
+namespace detail::adjacent_difference
+{
 #if _CCCL_HAS_CONCEPTS()
 template <typename T>
-concept adjacent_difference_policy_selector = policy_selector<T, adjacent_difference_policy>;
+concept adjacent_difference_policy_selector = policy_selector<T, AdjacentDifferencePolicy>;
 #endif // _CCCL_HAS_CONCEPTS()
 
 struct policy_selector
@@ -67,9 +68,9 @@ struct policy_selector
   bool may_alias;
 
   [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const
-    -> adjacent_difference_policy
+    -> AdjacentDifferencePolicy
   {
-    return adjacent_difference_policy{
+    return AdjacentDifferencePolicy{
       128,
       nominal_8B_items_to_items(7, value_type_size),
       BLOCK_LOAD_WARP_TRANSPOSE,
@@ -87,7 +88,7 @@ template <typename InputIteratorT, bool MayAlias>
 struct policy_selector_from_types
 {
   [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const
-    -> adjacent_difference_policy
+    -> AdjacentDifferencePolicy
   {
     constexpr auto policies = policy_selector{static_cast<int>(sizeof(it_value_t<InputIteratorT>)), MayAlias};
     return policies(cc);

cub/test/catch2_test_device_adjacent_difference_custom_policy_hub.cu

@@ -1,6 +1,11 @@
 // SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
+// TODO(bgruber): drop this test with CCCL 4.0 when we drop the adjacent difference dispatcher
+
+// disable deprecation warnings for DispatchAdjacentDifference
+#define CCCL_IGNORE_DEPRECATED_API
+
 #include "insert_nested_NVTX_range_guard.h"
 
 #include <cub/device/device_adjacent_difference.cuh>
@@ -14,9 +19,6 @@
 
 using namespace cub;
 
-// TODO(bgruber): drop this test with CCCL 4.0 when we drop the adjacent difference dispatcher after publishing the
-// tuning API
-
 template <typename InputIteratorT>
 struct my_policy_hub
 {

cub/test/catch2_test_device_adjacent_difference_env.cu

@@ -168,8 +168,7 @@ struct block_size_extracting_minus_t
 template <int ThreadsPerBlock>
 struct adj_diff_tuning
 {
-  _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const
-    -> cub::detail::adjacent_difference::adjacent_difference_policy
+  _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::AdjacentDifferencePolicy
   {
     return {ThreadsPerBlock, 1, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT, cub::BLOCK_STORE_DIRECT};
   }
@@ -241,3 +240,35 @@ C2H_TEST("DeviceAdjacentDifference::SubtractRight can be tuned", "[adjacent_diff
 }
 
 #endif // TEST_LAUNCH != 1
+
+#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 cannot preserve constexpr-ness from p1 to p2
+C2H_TEST("AdjacentDifferencePolicy", "[adjacent_difference][device]")
+{
+  STATIC_REQUIRE(::cuda::std::semiregular<cub::AdjacentDifferencePolicy>);
+  STATIC_REQUIRE(::cuda::std::is_aggregate_v<cub::AdjacentDifferencePolicy>);
+
+  // aggregate init
+  constexpr auto p1 = cub::AdjacentDifferencePolicy{
+    128,
+    7,
+    cub::BlockLoadAlgorithm::BLOCK_LOAD_WARP_TRANSPOSE,
+    cub::CacheLoadModifier::LOAD_LDG,
+    cub::BlockStoreAlgorithm::BLOCK_STORE_WARP_TRANSPOSE};
+
+#  if _CCCL_STD_VER >= 2020
+  // designated init
+  constexpr auto p2 = cub::AdjacentDifferencePolicy{
+    .threads_per_block = 128,
+    .items_per_thread  = 7,
+    .load_algorithm    = cub::BlockLoadAlgorithm::BLOCK_LOAD_WARP_TRANSPOSE,
+    .load_modifier     = cub::CacheLoadModifier::LOAD_LDG,
+    .store_algorithm   = cub::BlockStoreAlgorithm::BLOCK_STORE_WARP_TRANSPOSE};
+#  else // _CCCL_STD_VER >= 2020
+  constexpr auto p2 = p1;
+#  endif // _CCCL_STD_VER >= 2020
+
+  // comparison
+  STATIC_REQUIRE(p1 == p2);
+  STATIC_REQUIRE_FALSE(p1 != p2);
+}
+#endif // _CCCL_COMPILER(GCC, >=, 8)

cub/test/catch2_test_device_adjacent_difference_env_api.cu

@@ -7,6 +7,7 @@
 
 #include <thrust/device_vector.h>
 
+#include <cuda/__execution/tune.h>
 #include <cuda/devices>
 #include <cuda/stream>
 
@@ -104,3 +105,45 @@ C2H_TEST("cub::DeviceAdjacentDifference::SubtractRight accepts stream", "[adjace
   REQUIRE(error == cudaSuccess);
   REQUIRE(data == expected);
 }
+
+#if _CCCL_STD_VER >= 2020
+
+// example-begin subtract-left-copy-policy-selector
+struct AdjacentDifferencePolicySelector
+{
+  __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::AdjacentDifferencePolicy
+  {
+    return {.threads_per_block = 128,
+            .items_per_thread  = cc > cuda::compute_capability{9, 0} ? 11 : 7,
+            .load_algorithm    = cub::BLOCK_LOAD_WARP_TRANSPOSE,
+            .load_modifier     = cub::LOAD_LDG,
+            .store_algorithm   = cub::BLOCK_STORE_WARP_TRANSPOSE};
+  }
+};
+// example-end subtract-left-copy-policy-selector
+
+C2H_TEST("cub::DeviceAdjacentDifference::SubtractLeftCopy env-based API with tuning", "[adjacent_difference][env]")
+{
+  // example-begin subtract-left-copy-tuning
+  auto input  = thrust::device_vector<int>{1, 2, 1, 2, 1, 2, 1, 2};
+  auto output = thrust::device_vector<int>(8, thrust::no_init);
+
+  const auto error = cub::DeviceAdjacentDifference::SubtractLeftCopy(
+    input.begin(),
+    output.begin(),
+    input.size(),
+    cuda::std::minus{},
+    cuda::execution::tune(AdjacentDifferencePolicySelector{}));
+  if (error != cudaSuccess)
+  {
+    std::cerr << "cub::DeviceAdjacentDifference::SubtractLeftCopy failed with status: " << error << '\n';
+  }
+
+  thrust::device_vector<int> expected{1, 1, -1, 1, -1, 1, -1, 1};
+  // example-end subtract-left-copy-tuning
+
+  REQUIRE(error == cudaSuccess);
+  REQUIRE(output == expected);
+}
+
+#endif // _CCCL_STD_VER >= 2020