NVIDIA / NVIDIA/cccl

[BUG]: False-positive memcheck failure when using `atomic_ref` on small data types (e.g., int8_t)

Open
#6,430 0 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.