[PointKernel]

Description

Followup of #24050

This PR adds __launch_bounds__(hash_csr_block_size, 5) to hash_csr_build_kernel on Ampere and Hopper, which fixes most of the HashCSR groupby slowdown on Hopper.

Before Blackwell, NVVM compiles this kernel to 121 registers (252 with nested keys), so only two blocks fit per SM. The same source compiles to 48 to 56 registers on Blackwell. The kernel is latency bound, so at 100K groups it takes 3.9 ms on GH200 and 1.4 ms on B200. This is codegen rather than tuning: every CTK from 12.6 to 13.3 shows the same split, and sm_90 PTX assembled for sm_100 still needs 121 registers. With the bound, sm_80 to sm_90 compile to 48 registers and fit five blocks per SM, at the cost of some spills.

The bound is applied per device pass through __CUDA_ARCH__ because any bound changes Blackwell codegen. (256, 4) raises sm_100 from 48 to 64 registers, and even (256, 1) raises it to 76. Blackwell SASS is unchanged. Turing is excluded because it cannot fit five 256-thread blocks per SM.

GROUPBY_NVBENCH on GH200, GPU time vs main:

Benchmark	Configs	Change
groupby_max_cardinality	14	-20% to -35%
groupby_max	2	-35%
sum	4	-26% to -39%
groupby_m2_var_std	6	-30% to -31%
complex_int_keys	4	-33% to -42%
groupby_struct_keys	18	4 faster (-11% to -13%), the rest within 40 us

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

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

[coderabbitai]

Important

Review skipped

Review was skipped as selected files did not have any reviewable changes.

Run configuration

• Source mode: Auto
• Repository: NVIDIA/cudf
• Branch: main
• File: .coderabbit.yaml
• Commit: 57d7fe5
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Enterprise
• Run ID: b3baba57-ecf0-4734-aeb1-4b6b144ae1b8

Commits

Reviewing files that changed from the base of the PR and between 6f7e142 and 1d27eed.

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

• 🔍 Trigger review

No actionable comments were generated in the recent review.

ℹ️ Recent review info

⚙️ Run configuration

• Configuration used: Repository: NVIDIA/cudf/.coderabbit.yaml
• Review profile: CHILL
• Plan: Enterprise
• Run ID: 1d6ecacf-9f97-425e-9544-e7f87af42176

📥 Commits

Reviewing files that changed from the base of the PR and between 8fc209e and 5665f68.

📒 Files selected for processing: 1

• cpp/src/groupby/hash/hash_csr_kernels.cuh

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.

📝 Summary

Summary by CodeRabbit

• Performance
  • Optimized kernel launch configuration for CUDA architectures 800–999.
  • Runtime behavior and kernel parameters remain unchanged.

The hash_csr_build_kernel declaration applies __launch_bounds__(hash_csr_block_size, 5) on CUDA architectures 800–999. Its parameters remain unchanged.

Changes: Hash CSR kernel launch bounds

Layer / File(s)	Summary
Architecture-conditional launch bounds cpp/src/groupby/hash/hash_csr_kernels.cuh	The kernel declaration applies a five-block-per-SM launch bound on CUDA architectures 800–999. Its parameters remain unchanged.

Priority: Normal

Estimated code review effort: 1 (Trivial) | ~3 minutes

Change: Bug fix

Suggested reviewers: asuhan

Merge Risk: Minimal · up to 5665f

No actionable merge-blocking risk was identified in this focused change.

Pre-merge checks: 5

Check name	Status	Explanation
Title check		The title clearly identifies the main change: applying register bounds to the HashCSR groupby build kernel on Ampere and Hopper.
Docstring Coverage		No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check. Docstring coverage is scoped to functions touched by this diff. Analyzed 0 functions across 0…
Linked Issues check		Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check		Check skipped because no linked issues were found for this pull request.
Description check		The pull request description directly explains the HashCSR kernel launch-bound changes, targeted GPU architectures, performance rationale, benchmark results, and testing scope.

Full details: Docstring Coverage

Explanation

No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check. Docstring coverage is scoped to functions touched by this diff. Analyzed 0 functions across 0 files. (1 skipped: 1 unsupported.)

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create a new PR

• Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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

[mhaseeb123]

Speedups look good enough and the change within the directive seems small enough to justify.

[PointKernel]

/merge