[BUG]: Race reported in cub::BlockStore

[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/26497170643?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.block_run_length_decode
  ========= COMPUTE-SANITIZER
  ========= Error: Race reported between Write access at void cub::CUB_200500_600_700_800_NS::BlockStore<float, (int)64, (int)1, (cub::CUB_200500_600_700_800_NS::BlockStoreAlgorithm)4, (int)1, (int)1, (int)0>::StoreInternal<(cub::CUB_200500_600_700_800_NS::BlockStoreAlgorithm)4, (int)0>::Store<float *>(T1, float (&)[1], int)+0xa70 in /home/coder/cccl/cub/cub/block/block_store.cuh:980
  =========     and Read access at void cub::CUB_200500_600_700_800_NS::BlockStore<float, (int)64, (int)1, (cub::CUB_200500_600_700_800_NS::BlockStoreAlgorithm)4, (int)1, (int)1, (int)0>::StoreInternal<(cub::CUB_200500_600_700_800_NS::BlockStoreAlgorithm)4, (int)0>::Store<float *>(T1, float (&)[1], int)+0xaa0 in /home/coder/cccl/cub/cub/block/block_store.cuh:983 [9208 hazards]

Took a quick look at this one since it seems unbelievable. Looks like a false positive, but we should figure out why this is getting flagged. Likely has something to do with temp_storage.valid_items being marked volatile.

Lines 980 and 983 are the valid_items accesses:

    template <typename OutputIteratorT>
    _CCCL_DEVICE _CCCL_FORCEINLINE void Store(OutputIteratorT block_itr, T (&items)[ITEMS_PER_THREAD], int valid_items)
    {
      BlockExchange(temp_storage).BlockedToWarpStriped(items);
      if (linear_tid == 0)
      {
        // Move through volatile smem as a workaround to prevent RF spilling on
        // subsequent loads
        temp_storage.valid_items = valid_items;
      }
      CTA_SYNC(); // note, turns into __syncthreads();
      StoreDirectWarpStriped(linear_tid, block_itr, items, temp_storage.valid_items);
    }

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

Expected behavior

No diagnostic emitted.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response