Use the new tuning API internally for detail::find::dispatch

[gonidelis]

fixes #9237

• check that there are no/minimal sass diffs

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 21ddb9de-0e35-4ad0-b3aa-713bab13c80d

📥 Commits

Reviewing files that changed from the base of the PR and between ea94bad and 7e663df.

📒 Files selected for processing (1)

• cub/test/catch2_test_device_find_env.cu

🚧 Files skipped from review as they are similar to previous changes (1)

• cub/test/catch2_test_device_find_env.cu

---

Note: CodeRabbit is enabled on this repository as a convenience for maintainers
and contributors. Use your best judgment when considering its review comments and
suggestions — a suggested change may be inadequate, unnecessary, or safe to ignore.
Contributors are not expected to address every comment. Human reviews are what
ultimately matter for merging.

Overview

This PR migrates internal dispatch for DeviceFind::FindIf to use the new tuning API and environment-based dispatching so that tuned policies can be selected and applied from execution environments. No external/public API signatures were changed.

What I verified (files inspected)

• cub/benchmarks/bench/find_if/base.cu
• cub/cub/device/device_find.cuh
• cub/test/catch2_test_device_find_env.cu

Changes

• DeviceFind::FindIf (cub/cub/device/device_find.cuh)

  • The public FindIf entry now routes to detail::find::dispatch as before but the env-based launch path selects a default policy selector (based on iterator value type) and uses dispatch_with_env_and_tuning so that a policy_selector is computed and forwarded into detail::find::dispatch. No public API signature changes.

• Benchmark (cub/benchmarks/bench/find_if/base.cu)

  • Adds optional compile-time tuning ranges and a bench_policy_selector that constructs a cub::detail::find::find_policy from TUNE_* parameters (threads-per-block pow2, items-per-thread, and load modifier).
  • Reworks nvbench execution to create a cub_bench_env (with cuda::execution::tune when !TUNE_BASE) and calls DeviceFind::FindIf with that env via _CCCL_TRY_CUDA_API, consolidating global memory accounting and replacing the manual two-phase sizing/allocation flow.

• Tests (cub/test/catch2_test_device_find_env.cu)

  • Adds a find_tuning policy selector and a block-size-extracting predicate used by a new parameterized test ("Device FindIf can be tuned") that:
    • Constructs a tuned cuda::execution::tune(...) env, launches device_find_if with the tuned env, and verifies (1) no-match behavior (d_out == num_items) and (2) that the block size recorded on-device equals the requested ThreadsPerBlock.
  • Existing environment-usage tests for FindIf/LowerBound/UpperBound are left in place and still exercise env-based launches.

Impact / Notes

• Behavior: Enables runtime/launch-time tuning for FindIf via execution environments. This increases flexibility for policy selection without changing public APIs.
• Tests: A new test verifies the tuning path actually selects the requested block size.
• Benchmark: nvbench benchmark supports compile-time tuning knobs and uses the new env-based invocation.
• Review request: bernhardmgruber asked the author to run the "find if" benchmark SASS check to ensure there are no unintended SASS diffs—this PR adds tuning paths and a benchmark variant, so validating SASS output (or confirming only intended kernels are emitted) is recommended before merging.

suggestion:

Walkthrough

Updates DeviceFind::FindIf to compute and pass a tuning policy selector at dispatch; adds a test that exercises cuda::execution::tune(...) to enforce block size and verify it was applied; updates benchmark to optionally construct a compile-time policy selector and run under the tuning environment.

Changes

Tuning API Integration for DeviceFind::FindIf

Layer / File(s)	Summary
Policy dispatch with tuning cub/cub/device/device_find.cuh	DeviceFind::FindIf env path computes default_policy_selector from input value type and routes through dispatch_with_env_and_tuning, passing policy selector to detail::find::dispatch.
Tuning validation and test cub/test/catch2_test_device_find_env.cu	Adds #include <cuda/functional>, defines find_tuning<ThreadsPerBlock> and a block-size-extracting never-matching predicate, and a gated test that runs device_find_if under cuda::execution::tune(...) to validate output and observed block size.
Benchmark tuning integration cub/benchmarks/bench/find_if/base.cu	Adds bench_policy_selector<T> when !TUNE_BASE and refactors execution to use cub_bench_env, _CCCL_TRY_CUDA_API, and optional cuda::execution::tune(...), with updated nvbench tags and memory accounting.

Assessment against linked issues

Objective	Addressed	Explanation
Use new tuning API internally for detail::find::dispatch [#9237]	✅

Suggested reviewers

• shwina
• pauleonix
• davebayer
• bernhardmgruber

---

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[bernhardmgruber]

@gonidelis Please still do a SASS check of the find if benchmark!

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 1h 44m: Pass: 100%/284 | Total: 1d 22h | Max: 53m 27s | Hits: 98%/196763

See results here.

[gonidelis]

No SASS diffs observed for sm120