Implement the new tuning API for deterministic (rfa) reduce dispatch

[bernhardmgruber]

• Merge before: Implement the new tuning API for DeviceSegmentedReduce #7334
• Merge before: Fix DeviceReduce env test for rfa #7481
• CUB tests pass
• No SASS diff for cub.bench.reduce.deterministic.base on SM75/SM86/SM90

[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.

[bernhardmgruber]

/ok to test 1962a21

[oleksandr-pavlyk]

Perhaps is preprocessor to define auto rfa_dispatch differently based on TUNE_BASE and use that throughout to reduce branching?

Suggested change

	#if !TUNE_BASE
	cub::detail::rfa::dispatch<input_it_t, output_it_t, int, init_t, transform_t, accum_t, policy_selector_t>(
	nullptr, temp_storage_bytes, d_in, d_out, elements, {}, 0);
	#else
	cub::detail::rfa::dispatch<input_it_t, output_it_t, int, init_t, transform_t, accum_t>(
	nullptr, temp_storage_bytes, d_in, d_out, elements, {}, 0);
	#endif
	#if !TUNE_BASE
	auto rfa_dispatch = cub::detail::rfa::dispatch<input_it_t, output_it_t, int, init_t, transform_t, accum_t, policy_selector_t>;
	#else
	auto rfa_dispatch = cub::detail::rfa::dispatch<input_it_t, output_it_t, int, init_t, transform_t, accum_t>;
	#endif
	rfa_dispatch(nullptr, temp_storage_bytes, d_in, d_out, elements, {}, 0);

Then similarly for the actual invocation few lines below.

[bernhardmgruber]

That's an interesting suggestion, but if I take the address of the template instantiation of dispatch when I need to later provide all function arguments (defaults no longer work). I didn't know that! I refactored it differently now.

[github-actions]

😬 CI Workflow Results

🟥 Finished in 2h 44m: Pass: 98%/99 | Total: 22h 42m | Max: 2h 44m | Hits: 99%/108245

See results here.