compute-sanitizer reports invalid shared memory read in warpspeed DeviceScanKernel on Blackwell (sm_120a)
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Type of Bug
Something else
### Component
CUB
### Describe the bug
Running `compute-sanitizer --tool memcheck` against a program that calls `cub::DeviceScan::ExclusiveSum` (or `thrust::exclusive_scan`) on Blackwell GPU (sm_120a) produces the following error:
```
========= Invalid __shared__ read of size 16 bytes
========= at cp_async_bulk_cp_mask in cp_async_bulk.h:236
========= by thread (128,0,0) in block (0,0,0)
========= Address 0x400 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= ...
========= in function decode_page_headers ...
=========
========= ERROR SUMMARY: 1 error
```
A more full stack trace is available at this PR: https://github.com/rapidsai/cudf/pull/21999
The access address is 0x400 (1024 bytes). The `warpspeed` kernel's dynamic shared memory allocation is 8064 bytes (`smem_for_stages` with policy `{use_warpspeed=true, num_total_threads=352, items_per_thread=63, tile_size=8064})`, so the address is in-bounds. The program produces correct results.
The error originates inside `cub::detail::warpspeed::squadStoreBulkSync` → `cub::detail::scan::kernelBody` → `device_scan_lookahead_body` → `squadDispatch` → `DeviceScanKernel`.
### How to Reproduce
The error occurs in a libcudf example run with `compute-sanitizer` and can be seen here: https://github.com/rapidsai/cudf/actions/runs/23918157779/job/69767002847?pr=21999#step:12:835
The PR https://github.com/rapidsai/cudf/pull/21999 also contains additional information about reproducing the issue.
Unfortunately, we have not been able to create a standalone reproducer.
### Expected behavior
`compute-sanitizer` to pass with no errors
### Reproduction link
_No response_
### Operating System
_No response_
### nvidia-smi output
```
Fri Apr 3 17:11:46 2026
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 580.95.05 Driver Version: 580.95.05 CUDA Version: 13.0 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA RTX PRO 6000 Blac... Off | 00000000:3D:00.0 Off | 0 |
| N/A 32C P8 31W / 600W | 0MiB / 97887MiB | 0% Default |
| | | Disabled |
```
### NVCC version
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Tue_Dec_16_07:23:41_PM_PST_2025
Cuda compilation tools, release 13.1, V13.1.115
Build cuda_13.1.r13.1/compiler.37061995_0
$ compute-sanitizer --version
NVIDIA (R) Compute Sanitizer
Copyright (c) 2020-2025 NVIDIA Corporation
Version 2025.4.1.0 (build 37093031) (public-release)
Also fails on
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2026 NVIDIA Corporation
Built on Mon_Mar_02_09:52:23_PM_PST_2026
Cuda compilation tools, release 13.2, V13.2.51
Build cuda_13.2.r13.2/compiler.37434383_0
$ compute-sanitizer --version
NVIDIA (R) Compute Sanitizer
Copyright (c) 2020-2026 NVIDIA Corporation
Version 2026.1.0.0 (build 37182542) (public-release)
Contributor guide
Assessment
This issue has not been assessed yet.