[BUG]: Can't use cudaStreamTailLaunch in device code for CUB kernel stream
- 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
Assessment
This issue has not been assessed yet.