Implement the new tuning API for DeviceTransform

[bernhardmgruber]

Fixes: #6919
Fixes: #5057
Fixes: #3017

• CUB transform tests pass
• CCCL.C transform tests pass
• Rebase on top of Implement cuda::__all_arch_ids and cuda::__is_specific_arch #6916
• Update benchmarks
• Rename anything arch_policy related to policy_selector
• Merge before: Rename new tuning API policies and fix MSVC warning #7103
• No SASS difference for cub.bench.transform.babelstream.base on SM 75;80;90;100;120
• No relevant SASS difference for cub.test.device.transform.lid_0 on SM 75;80;90;100;120
• Compile-time comparison between before and after this PR (including clang in CUDA, because we don't have __CUDA_ARCH_LIST__.)
• Resolve decision on changing value_type of CCCL.C output iterators from void to the type they write to. See also: [BUG] CCCL.C output iterator cccl_iterator_t::value_type::size does not match JIT-compiled sizeof(iterator::value_type) #7169

Compile time of cub.test.device.transform.lid_0 using nvcc 13.0 and clang 20 for sm86, sm120

branch:
1m49.900s
1m50.615s
1m50.255s

main:
1m56.917s
1m57.378s
1m59.371s

Compile time of cub.test.device.transform.lid_0 for sm86, sm120 using clang 20 in CUDA mode:

branch:
real 1m40.627s
real 1m40.675s
real 1m40.912s

main:
real 1m39.273s
real 1m39.669s
real 1m39.835s

[copy-pr-bot]

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

[miscco]

Nitpick: I believe we should either use the concept emulation or plain SFINAE in C++17 too

[bernhardmgruber]

Hmm. We could also static_assert, but ArchPolicies is already used in the kernel attributes before we reach the body. And using a static_assert would only be evaluated in the device path.

How would I write that using concept emulation and have the concept check before the __launch_bounds__?

[miscco]

We could write:

Suggested change

	#if _CCCL_HAS_CONCEPTS()
	requires transform_policy_hub<ArchPolicies>
	#endif // _CCCL_HAS_CONCEPTS()
	_CCCL_TEMPLATE(typename PolicySelector,
	typename Offset,
	typename Predicate,
	typename F,
	typename RandomAccessIteratorOut,
	typename... RandomAccessIteratorsIn)
	_CCCL_REQUIRES(transform_policy_selector<PolicySelector>)

[bernhardmgruber]

Yeah, but as discussed on Slack before, we would need to get transform_policy_selector and then policy_selector working, which we couldn't of the is_constant_expression check. Let's leave it.

[miscco]

Nitpick: While it is technically more efficient, I believe it would improve readability if we did

    const bool all_inputs_contiguous = ::cuda::std::all_of(input.begin(), input.end(), [](const auto& input) { return input.is_contiguous; })

[bernhardmgruber]

Can I do this later? Maybe we have std::ranges::all_of by then.

[bernhardmgruber]

I see tiny changes in the generated SASS for cub.bench.transform.babelstream.base, notable in the filling kernels (no inputs) for complex<float>. The compiler now generates STG.E.ENL2.256, which it didn't do before.

The fill lernel for int128 seems to have degraded from generating STG.E.128 to a lot more STG.E.

All kernels with a functor marked as __callable_permitting_copied_arguments show no changes. That's good.

It feels a bit like the items per thread changed for the fill kernels.

[bernhardmgruber]

It feels a bit like the items per thread changed for the fill kernels.

They did. Before we had a tuning policy for sm_120, that was not taken into account :D This PR now uses it.

[bernhardmgruber]

I disabled the sm120 fill policy and now the only SASS diff for filling is on:

void cub::_V_300300_SM_1200::detail::transform::transform_kernel<cub::_V_300300_SM_1200::detail::transform::policy_hub<false, true, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, 0, 0>>, unsigned long*>::policy1000, long, cub::_V_300300_SM_1200::detail::transform::always_true_predicate, cuda::__4::__callable_permitting_copied_arguments<(anonymous namespace)::lognormal_adjust_t<unsigned long>>, unsigned long*, cuda::__4::counting_iterator<long, 0, 0>>(long, int, bool, cub::_V_300300_SM_1200::detail::transform::always_true_predicate, cuda::__4::__callable_permitting_copied_arguments<(anonymous namespace)::lognormal_adjust_t<unsigned long>>, unsigned long*, cub::_V_300300_SM_1200::detail::transform::kernel_arg<cuda::__4::counting_iterator<long, 0, 0>>)

which is a thrust::tabulate of a counting_iterator<long> and an unsigned long*.

[bernhardmgruber]

Found the final issue with the fill kernels. Disabled the vectorized tunings when we have input streams (they were tuned for output only use cases). SASS of cub.bench.transform.fill.base now matches baseline on sm120.

[gevtushenko]

question: what was the intention for __needs_a_constexpr_value here? Do you want to check if hub's operator() can return compile-time information? If so, I don't think this works. Let's add a test for non-constexpr operator to cover this. For instance, the following type satisfies this concept:

struct policy_selector_all {
  auto operator()(arch_id id) const -> a_policy {
    int r = rand() % 5;
    return a_policy{static_cast<arch_id>(r)};
  }
};

maybe something along the following lines would fix this:

template <auto> inline constexpr bool __needs_a_constexpr_value = true;
// ...
requires __needs_a_constexpr_value<T{}(arch_id{})>;

[bernhardmgruber]

While fixing this I got this error and I am massively impressed:

/home/bgruber/dev/cccl/lib/cmake/cub/../../../cub/cub/device/dispatch/tuning/tuning_reduce.cuh(474): error: expression must have a constant value
  static_assert(__needs_a_constexpr_value2<policy_selector{}(::cuda::arch_id::sm_60)>);
                                           ^
/home/bgruber/dev/cccl/lib/cmake/cub/../../../cub/cub/util_arch.cuh(150): note #61-D: integer operation result is out of range
      ::cuda::std::clamp(nominal_4B_items_per_thread * 4 / target_type_size, 1, nominal_4B_items_per_thread * 2);
                                                         ^
/home/bgruber/dev/cccl/lib/cmake/cub/../../../cub/cub/device/dispatch/tuning/tuning_reduce.cuh(440): note #2693-D: called from:
        auto [scaled_items, scaled_threads] = scale_mem_bound(threads_per_block, items_per_thread, accum_size);
                                                             ^

The concept check did actually validate whether the computation of the tuning policy was sound. The reason it fails here is because a default constructed policy selector like reduce::policy_selector{} has a bunch of zero data members now, like accum_size, leading to the division by zero later. So we cannot test from just the type of a policy selector whether it returns a policy at compile-time. I will drop the constexpr test from the concept.

I was btw doing this investigation while I had Cursor (with claude-4.5-opus-high) trying to fix it as well. Cursor was fast in figuring out how to run the tests, but got really lost when the first two attempts of changing the concept definition didn't fix the problem.

[gevtushenko]

Excited to see the new tuning machinery at work! Code is much more readable now and we no longer have to parse PTX 🎉

[gevtushenko]

important: to my understanding, we no longer have to parse policy from the PTX. Let's drop this.

[bernhardmgruber]

This is for caching the items per thread we computed based on the occupancy in the transform dispatch, which involves a bunch of CUDA API calls and that's why we cache it. In ordinary CUB we can do this in a static variable, since there is one such variable for each template instantiation of the CUB algo, and each template instantiation with a distinct set of types may have a different config. In CCCL.C, we type-erase the iterators so the caching mechanism has to work differently, which is handled by the above cache. In my understanding, this is still needed.

[github-actions]

😬 CI Workflow Results

🟥 Finished in 6h 01m: Pass: 98%/133 | Total: 7d 08h | Max: 6h 01m | Hits: 52%/174121

See results here.