Tunings#

Device-scope algorithms in CUB have many knobs that significantly impact performance (without affecting correctness). For instance, the number of threads per block and items per thread can be tuned to maximize performance for a given device and data type. But also algorithmic choices such as the load or store algorithm, a load vectorization size, the used block-level algorithm, or enabling the use of the tensor memory accelerator can be tweaked. Most device-scope algorithms in CUB accept a set of such tuning parameters, by passing a policy selector, wrapped into cuda::execution::tune(...), as part of the environment of a CUB device-scope API. In the following, we describe this process in more detail.

Warning

CUB’s built-in tunings may change in every CCCL release, and so may the performance of algorithms using a custom tuning. In some cases, compilation may fail as well. See Stability of tunings for what this means for custom tunings.

Policy selectors#

Policy selectors are the mechanism by which CUB’s device algorithms and kernels select tuning parameters for a given workload and GPU compute capability. A policy selector is a stateless callable that maps a cuda::compute_capability to a policy struct containing the tuning values for that compute capability. Each set of CUB algorithms using a common underlying implementation defines a common policy struct, e.g. cub::ReducePolicy, which must be returned by a policy selector passed to those algorithms.

CUB employs internal default policy selectors providing tunings for known compute capabilities and workloads, which are not publicly accessible to users. Users can override CUB’s policy selector for a given algorithm by passing a custom policy selector through the algorithm’s environment parameter, or build upon CUB’s default tuning by passing a policy transformer.

For a description of how policy selectors are used internally in CUB’s dispatch layer and kernels see the corresponding developer documentation.

Defining a policy selector#

A policy selector is any type with a __host__, __device__, constexpr, and const call operator taking a cuda::compute_capability and returning the algorithm’s policy struct:

template <typename T>
struct my_reduce_tuning {
  __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::ReducePolicy {
    // tuning for Hopper and later
    if (cc >= cuda::compute_capability(9, 0)) {
      const auto pass = cub::ReducePassPolicy{
        .threads_per_block = 512,
        .items_per_thread = std::max<int>(64 / sizeof(T), 1), // 8 double, 16 float, 32 half_t, ...
        .vec_size = 2,
        .reduce_algorithm = cub::BLOCK_REDUCE_WARP_REDUCTIONS,
        .load_modifier = cub::LOAD_DEFAULT
      };
      return { .multi_tile = pass, .single_tile = pass };
    }
    // fallback for older GPUs
    const auto pass = cub::ReducePassPolicy{
      .threads_per_block = 256,
      .items_per_thread = 12,
      .vec_size = 1,
      .reduce_algorithm = cub::BLOCK_REDUCE_WARP_REDUCTIONS,
      .load_modifier = cub::LOAD_DEFAULT
    };
    return { .multi_tile = pass, .single_tile = pass };
  }
};

Warning

The policy selector must be stateless (std::is_empty_v<T> must be true) since only its type will be passed to a kernel — any captured state would be silently lost. Policy selectors are also freely constructed and copied where needed, so they must also be default constructible and copyable (std::semiregular<T> must be true).

It can be a class template, but then only a full specialization can be passed to CUB, e.g., my_reduce_tuning<float>. The implementation can use branches and helper functions arbitrarily, as long as they can be evaluated at compile-time. The returned tuning policy can, and probably should, contain different values for different compute capabilities or workloads.

Each CUB algorithm with an environment parameter searches the environment for a policy selector returning a matching policy struct. If one is found, the policy selector will be used to determine the tuning values for host-side dispatch and kernel compilation. Each CUB algorithm documents the policy struct to which it responds. If CUB does not find a matching policy selector in the environment, it falls back to its internal default policy selector.

Multiple policy selectors returning different policy structs can be passed as part of the same environment, and each algorithm will pick the one with the matching policy struct. This is useful if the same environment is reused across several algorithm calls.

The policy structs themselves are simple semiregular aggregates. They support C++20 designated initializers (i.e., the syntax { .threads_per_block = 512, ... }), comparison for (in-)equality, and serialization using operator<<. They may occasionally contain member functions that compute derived values from the contained tuning values. All policy structs are public types and will evolve in a non-breaking way, at least during minor releases, by only adding new data members at the end of the struct.

Policy transformers#

Often, only a few values of CUB’s default tuning need to change. For example, CUB’s tuning for your GPU and workload is good, except you want to choose a different block-level algorithm which suits your operator better. Instead of a policy selector, a policy transformer can be passed to cuda::execution::tune(...). A policy transformer is a stateless callable taking a cuda::compute_capability and the policy returned by CUB’s default policy selector for that compute capability and workload. It returns the (possibly modified) policy, which is then used to tune the algorithm:

struct AdjacentDifferencePolicyTransformer
{
  __host__ __device__ constexpr auto operator()(cuda::compute_capability, cub::AdjacentDifferencePolicy policy) const
    -> cub::AdjacentDifferencePolicy
  {
    // tweak block size, but keep all other values of CUB's default tuning
    policy.threads_per_block = 64;
    return policy;
  }
};

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(AdjacentDifferencePolicyTransformer{}));
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};

The same requirements as for policy selectors apply: a policy transformer must be stateless and semiregular. It must return the algorithm’s policy struct by value, and accept it as its second parameter. A callable passed to cuda::execution::tune(...) must be either a policy selector or a policy transformer.

Passing a policy selector to CUB device-scope algorithms#

Custom policy selectors are passed to CUB algorithms via the environment argument. They first need to be wrapped by passing them to cuda::execution::tune(...):

cub::DeviceReduce::Reduce(
  d_in, d_out, num_items, op, init,
  cuda::execution::tune(my_reduce_tuning<int>{}));

Multiple tunings for different algorithms can be combined in a single environment:

auto env = cuda::execution::tune(
  my_reduce_tuning<int>{},
  my_scan_tuning{}
);
cub::DeviceReduce::Reduce(d_in, d_out, num_items, op, init, env);
cub::DeviceScan::ExclusiveSum(d_in, d_out, num_items, env);

An environment can carry further properties like a stream or a memory resource. Policy selectors can simply be added to those:

auto env = cuda::std::execution::env{
  stream_ref,
  resource,
  cuda::execution::tune(
    my_reduce_tuning<int>{}, my_scan_tuning{})
};
cub::DeviceReduce::Reduce(d_in, d_out, num_items, op, init, env);

// alternatively, if we want to extend an env `other_env`
auto env = cuda::std::execution::env{
  other_env,
  cuda::execution::tune(my_scan_tuning{})
};
cub::DeviceReduce::Reduce(d_in, d_out, num_items, op, init, env);

CUB’s benchmarks also make heavy use of policy selectors for tuning. For more details on authoring benchmarks and their policy selectors for automatic tuning, see Automated Tuning Infrastructure.

Stability of tunings#

CCCL reserves the right to change CUB’s built-in tunings in every release, including minor and patch releases. Such changes are validated using CUB’s benchmarks, but may still regress workloads outside our benchmark coverage.

Algorithmic changes in CUB also sometimes require tunings to change to retain performance. For example, restructuring an algorithm to allow for higher load vectorization, may require different items processed per thread to be effective. A user-tuned value for the items per thread may now select a slower fallback path after updating CCCL, without any diagnostic.

In extreme cases CUB may change tunings in a way where a custom override could lead to a compilation failure. Suppose that changes to CUB’s tunings lead to a higher use of GPU resources like shared memory. A user-tuned value for the block size may now exceed the available shared memory and fail to compile.

Therefore, a specific tuning overriding CUB’s internal one may work for a given CCCL version, but may no longer be adequate or even fail to compile in a future CCCL version. We strongly advise users to regularly benchmark and reevaluate their custom tunings when upgrading.

The two ways of customizing a tuning are affected differently by such changes:

  • A policy selector specifies all tuning values, so changes to CUB’s tuning values do not affect it. But it locks you out of any improvements CUB makes to its tunings in newer CCCL versions, and it is still affected by changes to the algorithms themselves.

  • A policy transformer only modifies CUB’s current tuning and thus picks up such improvements. But the values it does not override can change underneath it, and CUB does not validate your modification in combination with future tunings. The result may perform worse or no longer compile.

Use a policy transformer for a small and targeted modification of CUB’s tuning. Use a policy selector to freeze a good tuning you found over time.