NVIDIA / NVIDIA/cccl

[BUG]: `may be used uninitialized` with cub/thrust reduce.

Open
#7,246 6 comments 0 reactions 1 assignee Claimed by @jrhemstad View on GitHub
needs triage
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

I'm running into compiler warnings about uninitialized values when testing XGBoost with the nightly RMM. I don't know the exact commit of cub/thurst yet, here are the version macros:
```
#define CUB_VERSION 300200
#define THRUST_VERSION 300200
```

There are multiple places triggering the warning, one of them is using tabulate iterator, as shown in the below example. I attached the full build log warning for all cases (1.3MB, so be careful before you open it with your editor), and the error log for the reproducible example.

[rmm-build-log.zip](https://github.com/user-attachments/files/24643679/rmm-build-log.zip)

[build-log-example.zip](https://github.com/user-attachments/files/24643683/build-log-example.zip)

### How to Reproduce

I wrote a small function in XGBoost test file to reproduce the issue. Please let me know if you want me to extract a clean example. Compiling the following function should reproduce the warning.
``` c++
void TestSegReduce() {
std::size_t n_segs = 4;
thrust::device_vector indptr{1, 3, 6, 9, 12};
auto it = thrust::make_transform_iterator(thrust::make_counting_iterator(0),
[] __host__ __device__(std::int32_t i) { return i; });
thrust::device_vector out(n_segs);
auto p_out = out.data();
auto out_it = cuda::make_tabulate_output_iterator(
[=] __host__ __device__(std::int32_t idx, std::int32_t v) { p_out[idx] = v; });

std::size_t n_bytes = 0;
cub::DeviceSegmentedReduce::Sum(nullptr, n_bytes, it, out_it, n_segs, indptr.data(),
indptr.data() + 1, cudaStreamPerThread);
thrust::device_vector alloc(n_bytes);
cub::DeviceSegmentedReduce::Sum(alloc.data().get(), n_bytes, it, out_it, n_segs, indptr.data(),
indptr.data() + 1, cudaStreamPerThread);
}
```

### Expected behavior

No warning.

### Reproduction link

_No response_

### Operating System

Ubuntu 24.04

### nvidia-smi output

_No response_

### NVCC version

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Aug_20_01:58:59_PM_PDT_2025
Cuda compilation tools, release 13.0, V13.0.88
Build cuda_13.0.r13.0/compiler.36424714_0

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.