[BUG]: Race reported in CUB's merge sort implementation

[alliepiper]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Something else

Component

CUB

Describe the bug

https://github.com/NVIDIA/cccl/actions/runs/9606424817/job/26497170740?pr=1879

  	compute-sanitizer --tool racecheck --check-device-heap yes --leak-check full --padding 512 --track-stream-ordered-races all --check-warpgroup-mma yes --require-cuda-init no --check-exit-code yes --error-exitcode 1 --nvtx true /home/coder/cccl/build/cuda12.4-gcc13/cub-cpp17/bin/cub.cpp17.test.device_merge_sort.lid_2
  ========= COMPUTE-SANITIZER
  ========= Error: Race reported between Write access at void cub::CUB_200500_600_700_800_NS::AgentMerge<cub::CUB_200500_600_700_800_NS::AgentMergeSortPolicy<(int)256, (int)17, (cub::CUB_200500_600_700_800_NS::BlockLoadAlgorithm)4, (cub::CUB_200500_600_700_800_NS::CacheLoadModifier)0, (cub::CUB_200500_600_700_800_NS::BlockStoreAlgorithm)4>, short *, float *, int, custom_less_op_t, short, float>::reg_to_shared<short, short *>(T2, T1 (&)[17])+0xa120 in /home/coder/cccl/cub/cub/agent/agent_merge_sort.cuh:487
  =========     and Write access at void cub::CUB_200500_600_700_800_NS::detail::uninitialized_copy<float, float &, (int)0>(T1 *, T2 &&)+0xd400 in /home/coder/cccl/cub/cub/detail/uninitialized_copy.cuh:61 [12 hazards]
  =========     and Write access at void cub::CUB_200500_600_700_800_NS::detail::uninitialized_copy<float, float &, (int)0>(T1 *, T2 &&)+0xd500 in /home/coder/cccl/cub/cub/detail/uninitialized_copy.cuh:61 [12 hazards]
  =========     and Read access at void cub::CUB_200500_600_700_800_NS::BlockExchange<float, (int)256, (int)17, (bool)0, (int)1, (int)1, (int)0>::BlockedToWarpStriped<float>(float (&)[17], T1 (&)[17], cub::CUB_200500_600_700_800_NS::Int2Type<(int)0>)+0xd530 in /home/coder/cccl/cub/cub/block/block_exchange.cuh:346 [124 hazards]
  =========     and Read access at void cub::CUB_200500_600_700_800_NS::BlockExchange<float, (int)256, (int)17, (bool)0, (int)1, (int)1, (int)0>::BlockedToWarpStriped<float>(float (&)[17], T1 (&)[17], cub::CUB_200500_600_700_800_NS::Int2Type<(int)0>)+0xd630 in /home/coder/cccl/cub/cub/block/block_exchange.cuh:346 [120 hazards]
  =========     and Read access at void cub::CUB_200500_600_700_800_NS::BlockExchange<float, (int)256, (int)17, (bool)0, (int)1, (int)1, (int)0>::BlockedToWarpStriped<float>(float (&)[17], T1 (&)[17], cub::CUB_200500_600_700_800_NS::Int2Type<(int)0>)+0xd640 in /home/coder/cccl/cub/cub/block/block_exchange.cuh:346 [118 hazards]
  =========     and Read access at void cub::CUB_200500_600_700_800_NS::BlockExchange<float, (int)256, (int)17, (bool)0, (int)1, (int)1, (int)0>::BlockedToWarpStriped<float>(float (&)[17], T1 (&)[17], cub::CUB_200500_600_700_800_NS::Int2Type<(int)0>)+0xd650 in /home/coder/cccl/cub/cub/block/block_exchange.cuh:346 [124 hazards]

How to Reproduce

compute-sanitizer --tool racecheck --check-device-heap yes --leak-check full --padding 512 --track-stream-ordered-races all --check-warpgroup-mma yes --require-cuda-init no --check-exit-code yes --error-exitcode 1 --nvtx true /home/coder/cccl/build/cuda12.4-gcc13/cub-cpp17/bin/cub.cpp17.test.device_merge_sort.lid_2

Expected behavior

No diagnostic emitted.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response