NVIDIA / NVIDIA/cccl

[BUG]: Can't use cudaStreamTailLaunch in device code for CUB kernel stream

Open
#6,537 1 comment 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

Runtime Error

### Component

CUB

### Describe the bug

CUB reports an error when launching a Sum kernel from device code on the `cudaStreamTailLaunch` stream.
Error: CUDA error 1 [block (0,0,0) thread (0,0,0)

### How to Reproduce

Godbolt doesn't seem to work with dynamic parallelism so I can't post an easy to repro code but the following code should clearly demonstrate the issue (it needs to be compiled with enabled dynamic parallelism):
```
#include
#include
#include
#include

__device__ int dVar = 0;

static constexpr int kNumElements = 10;

__global__ void C0(int *v, uint8_t *temp_storage, size_t temp_storage_bytes)
{
// ... launch kernel that would fill `v` ...
cub::DeviceReduce::Sum(temp_storage, temp_storage_bytes, v, &dVar, kNumElements, cudaStreamTailLaunch);
// ... launch another cudaStreamTailLaunch kernel that would read `dVar` ...
}

int main()
{
thrust::device_vector vec(10, 0);
size_t temp_storage_bytes = 0;
cub::DeviceReduce::Sum(nullptr, temp_storage_bytes, thrust::raw_pointer_cast(vec.data()), thrust::raw_pointer_cast(vec.data()), vec.size());
thrust::device_vector temp_storage(temp_storage_bytes);
C0<<<1, 1>>>(thrust::raw_pointer_cast(vec.data()), thrust::raw_pointer_cast(temp_storage.data()), temp_storage_bytes);
cudaDeviceSynchronize();
return 1;
}
```

All works fine when the `cudaStreamTailLaunch` is removed.

### Expected behavior

CUB kernel should launch after the previous kernels to be finished (due to the cudaStreamTailLaunch).

### Reproduction link

_No response_

### Operating System

Windows 11

### nvidia-smi output
```
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 581.57 Driver Version: 581.57 CUDA Version: 13.0 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Driver-Model | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA GeForce RTX 5070 Ti WDDM | 00000000:01:00.0 On | N/A |
| 0% 31C P8 17W / 300W | 1227MiB / 16303MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
```

### NVCC version

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.