Skip to content

Bound HashCSR groupby build kernel registers on Ampere and Hopper #24747

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
rapids-bot merged 4 commits into from
Oct 26, 2026
Merged

Bound HashCSR groupby build kernel registers on Ampere and Hopper #24747

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
8f51788
Bound HashCSR groupby build kernel registers on Ampere and Hopper
PointKernel Oct 21, 2026
5665f68
Apply HashCSR build launch bounds without a macro
PointKernel Oct 21, 2026
6f7e142
Merge branch 'main' into fix/hashcsr-build-pre-blackwell-occupancy
PointKernel Oct 23, 2026
1d27eed
Merge branch 'main' into fix/hashcsr-build-pre-blackwell-occupancy
PointKernel Oct 26, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
25 changes: 16 additions & 9 deletions cpp/src/groupby/hash/hash_csr_kernels.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -102,15 +102,22 @@ struct hash_set_ref {
* the rows still to come are skipped: the caller then rebuilds with a larger set.
*/
template <typename Equal, typename Hasher>
CUDF_KERNEL void hash_csr_build_kernel(size_type num_rows,
bitmask_type const* valid_rows,
build_position_type* positions,
size_type* slot_counts,
bool count_by_representative,
hash_set_ref set,
Equal equal,
Hasher hasher,
cuda::std::int32_t* overflow)
CUDF_KERNEL void
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 && __CUDA_ARCH__ < 1000
Comment thread
mhaseeb123 marked this conversation as resolved.
// Ampere and Hopper compile this kernel to over 120 registers, two blocks per SM; bounding it
// trades a few spills for 2.5x the occupancy. Any bound raises Blackwell's register count, and
// Turing cannot fit five 256-thread blocks per SM.
__launch_bounds__(hash_csr_block_size, 5)
#endif
hash_csr_build_kernel(size_type num_rows,
bitmask_type const* valid_rows,
build_position_type* positions,
size_type* slot_counts,
bool count_by_representative,
hash_set_ref set,
Equal equal,
Hasher hasher,
cuda::std::int32_t* overflow)
{
auto const block = cooperative_groups::this_thread_block();
auto const warp = cooperative_groups::tiled_partition<cudf::detail::warp_size>(block);
Expand Down
Loading