Fix racecheck in segmented_offset_bitmask_binop kernel

[davidwendt]

Description

Fixes a racecheck in the segmented_offset_bitmask_binop libcudf kernel. This was found in the weekly compute-sanitizer run: https://github.com/NVIDIA/cudf/actions/runs/33247027974/job/99086397670#step:5:3039

[ RUN      ] MergeBitmaskTest.TestSegmentedBitmaskAndSingleSegment
========= Error: Race reported between Write access at decltype((param#3(param#2, param#2))) cooperative_groups::__v1::details::tile_reduce_dispatch<(unsigned int)5>::reduce<(unsigned int)256, cooperative_groups::__v1::thread_block, int &, cooperative_groups::__v1::plus<int>>(const cooperative_groups::__v1::thread_block_tile<T1, T2> &, T3 &&, T4 &&)::[lambda(const cooperative_groups::__v1::details::internal_thread_block_tile<(unsigned int)32, cooperative_groups::__v1::__static_size_multi_warp_tile_base<(unsigned int)256>> &, int *) (instance 1)]::operator ()(const cooperative_groups::__v1::details::internal_thread_block_tile<(unsigned int)32, cooperative_groups::__v1::__static_size_multi_warp_tile_base<(unsigned int)256>> &, int *) const+0xeb0 in reduce.h:283
=========     and Read access at decltype((param#3(param#2, param#2))) cooperative_groups::__v1::details::coalesced_reduce<int &, cooperative_groups::__v1::plus<int>, (unsigned int)8, cooperative_groups::__v1::details::internal_thread_block_tile<(unsigned int)32, cooperative_groups::__v1::__static_size_multi_warp_tile_base<(unsigned int)256>>>(const cooperative_groups::__v1::__single_warp_thread_block_tile<T3, T4> &, T1 &&, T2 &&)+0x10d0 in coalesced_reduce.h:66 [28 hazards]
=========

The race compute-sanitizer detects is a real shared-memory hazard: within cg::reduce's hierarchical reduction tree (e.g., 8 warps → 4 → 2 → 1), one level writes partial sums to shared memory and the next level reads them. The CG library uses __syncthreads() internally to order these, but compute-sanitizer racecheck operates at a warp granularity — it sees one warp writing shared memory and another warp reading that same location, and flags it as a race because the __syncthreads barrier is between those two events but racecheck's memory access tracking can fire before the barrier is resolved in its model.

The solution employed here replaces the CG calls with equivalent CUB BlockReduce calls instead which seems to make compute-sanitizer happy. The benchmarks showed no performance regressions with the modified code.

Checklist

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

[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: 3cfe8c31-2fbc-407a-8a4a-1ccd616fcfb7

📥 Commits

Reviewing files that changed from the base of the PR and between 75ec3b9 and 7de3437.

📒 Files selected for processing (1)

• cpp/include/cudf/detail/null_mask.cuh

Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.

---

📝 Walkthrough

Summary by CodeRabbit

• Performance
  • Improved segmented bitmask operations by simplifying GPU thread and block assignment.
  • Updated null-count aggregation for more efficient device-side reduction.

Walkthrough

segmented_offset_bitmask_binop now uses direct CUDA block and thread indices. It replaces cooperative-groups reduction with cub::BlockReduce and uses thread 0 for relaxed device-scope atomic null-count accumulation.

Changes

Null mask bitmask operation

Layer / File(s)	Summary
Block indexing and null-count reduction cpp/include/cudf/detail/null_mask.cuh	The operation assigns segments and destination words with blockIdx.x and threadIdx.x. It reduces counts with cub::BlockReduce, then atomically accumulates the positive result from thread 0.

Estimated code review effort: 2 (Simple) | ~10 minutes

Merge Risk: ⚪ Minimal · up to 7de34

This localized kernel change replaces the reduction implementation to address the reported racecheck issue without changing the intended behavior; no actionable merge-blocking risk remains after normal checks and review.

Suggested reviewers: mythrocks, nirandaperera, pointkernel

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Description check	✅ Passed	The description clearly explains that the change fixes a compute-sanitizer racecheck report by replacing cooperative-groups reductions with CUB BlockReduce in the targeted kernel.
Title check	✅ Passed	The title clearly identifies the primary change: fixing a racecheck issue in the segmented_offset_bitmask_binop kernel.
Docstring Coverage	✅ Passed	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	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

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 PR with unit tests

---

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