[BUG]: False-positive memcheck failure when using `atomic_ref` on small data types (e.g., int8_t)
- 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
libcu++
### Describe the bug
When using `atomic_ref::min` (or other atomic operations) on small data types such as `int8_t`, `compute-sanitizer` (memcheck) reports an **invalid global read of size 4 bytes**.
This occurs when the underlying memory allocation is tightly packed (i.e., no 4-byte padding). The issue stems from CCCL’s use of a 4-byte CAS loop to emulate atomic operations on types smaller than 4 bytes.
While the sanitizer reports an out-of-bounds load, the operation is known to be **safe on all supported hardware**—making this a **false positive**.
libcudf currently works around this by rounding small-type column buffer sizes to a multiple of 4 bytes to silence the failures in our nightly CI.
### How to Reproduce
Compile and run the following kernel under `compute-sanitizer` (memcheck) with input data size of 5:
```cpp
__global__ void test(int8_t* data) {
cuda::atomic_ref ref(*data);
ref.fetch_min(1);
}
```
**Observed output:**
```
========= Invalid __global__ read of size 4 bytes
========= at void cuda::std::__cuda_atomic_load(...)
```
This only happens when `data` points to a buffer without 4-byte padding.
### Expected behavior
* No invalid global memory access should be reported by `compute-sanitizer`.
* The sanitizer should recognize that the 4B CAS operation does not modify bytes outside the target value.
### Additonal notes
* The operation is safe and behaves correctly on all supported GPU architectures.
* CCCL maintainers confirmed this is not a bug in `atomic_ref`, but a sanitizer false positive.
* Potential solution: NVVM intrinsic or PTX pragma to annotate these safe atomic operations so memcheck can suppress the warning.
Contributor guide
Assessment
This issue has not been assessed yet.