Parallelize segmented bitmask binop over words in addition to segments

[vuule]

Description

closes #23487

segmented_offset_bitmask_binop assigned one warp per segment, capping parallelism at num_segments * 32 threads regardless of mask size. With few segments and large masks -- the shape superimpose_nulls produces -- the kernel left almost the whole GPU idle.

The kernel now uses a 2D grid: x over destination words, y over segments. Each thread computes one destination word, so parallelism scales with the mask size. Since a segment can span many blocks, null counts are reduced per block with cub::BlockReduce and accumulated with atomicAdd, and inplace_segmented_bitmask_binop switched to make_zeroed_device_uvector_async to zero the accumulator. gridDim.y is capped at 65535, so blocks stride over segments when there are more than that.

Every destination word still depends only on the same word index of its sources, so a destination mask may also be a source of another segment.

This also fixes an out-of-bounds read: the kernel indexed segment_offsets and destinations before checking segment_id >= num_segments, and num_blocks rounds up. The bounds check is now the loop condition, ahead of every read.

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]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[vuule]

Performance

A100 80GB PCIe, BITMASK_NVBENCH -b segmented_bitmask_and -d 0. GPU time, and global memory bandwidth utilization in parentheses.

mask_size_bits = 1048576:

segments	masks/segment	before	after	speedup
10	4	2.306 ms (0.10%)	54.6 us (4.22%)	42x
100	4	3.122 ms (0.76%)	191.4 us (12.35%)	16x
1,000	4	4.283 ms (5.48%)	1.542 ms (15.22%)	2.8x
10	8	4.006 ms (0.13%)	56.6 us (8.85%)	71x
100	8	5.066 ms (1.00%)	228.3 us (22.21%)	22x
1,000	8	6.443 ms (7.85%)	1.895 ms (26.68%)	3.4x
10	16	7.814 ms (0.13%)	70.3 us (14.83%)	111x
100	16	8.975 ms (1.17%)	303.3 us (34.58%)	30x
1,000	16	10.785 ms (9.71%)	2.602 ms (40.26%)	4.1x

The gain shrinks as segments grow because the old mapping only starts to fill the device once num_segments reaches the thousands.

The shipped axes (mask_size_bits 32 to 128) show no regression -- every config is within noise or slightly faster. At those sizes a segment is one to four words, so there is no parallelism to gain and the per-segment output allocation dominates the runtime. A sample:

segments	masks/segment	mask_size_bits	before	after
100	4	32	127.9 us	127.2 us
1,000	4	32	999.7 us	970.7 us
10,000	4	32	9.608 ms	9.411 ms
100	16	128	138.2 us	137.2 us
1,000	16	128	1.039 ms	1.010 ms
10,000	16	128	9.886 ms	9.640 ms

Validation

BITMASK_TEST, STRUCTS_TEST, COLUMN_TEST, ORC_TEST, and PARQUET_TEST all pass, and compute-sanitizer --tool=memcheck on BITMASK_TEST reports 0 errors.

Note that memcheck did not flag the out-of-bounds read on main either -- the stray reads land inside the pool allocation, so it cannot see them.