| Issue |
141626
|
| Summary |
Regression: Cannot compile cuda code that calls CUB function with clang 20
|
| Labels |
clang
|
| Assignees |
|
| Reporter |
jabraham17
|
I am running some cuda code which makes use of CUB on an NVIDIA A100 GPU. I am compiling the code using clang++ 20 and running into issues that were not present in clang++ 19.
The following is my simplified test case, compiled as `clang++ bug.cu -I.... -L... -lcuda -lcudart`. I am using cuda 12.3.
```cuda
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cub/cub.cuh>
int main() {
int deviceIdx = 0;
cudaSetDevice(deviceIdx);
int N = 100;
float* arr;
cudaMalloc(&arr, N * sizeof(float));
float* sum;
cudaMalloc(&sum, sizeof(float));
void *temp_storage = nullptr;
size_t n_temp_storage = 0;
auto res = cub::DeviceReduce::Sum(temp_storage, n_temp_storage, arr, sum, N, (CUstream)0);
if ((int)res != CUDA_SUCCESS) {
printf("Error: %d: %s\n", res, cudaGetErrorString((cudaError_t)res));
return -1;
}
cudaMalloc(&temp_storage, n_temp_storage);
auto res2 = cub::DeviceReduce::Sum(temp_storage, n_temp_storage, arr, sum, N, (CUstream)0);
if ((int)res2 != CUDA_SUCCESS) {
printf("Error: %d %s\n", res2, cudaGetErrorString((cudaError_t)res2));
return -1;
}
float sum_host;
cudaMemcpy(&sum_host, sum, sizeof(float), cudaMemcpyDeviceToHost);
printf("Sum: %f\n", sum_host);
}
```
With clang 19.1.3, the code compiles and runs, reaching the `printf("Sum: ...` line. With clang 20.1.4, `cub::DeviceReduce::Sum` returns an error and exits
```
Error: 209: no kernel image is available for execution on the device
```
If I compile the exact same code with `nvcc`, it also works with cuda 12.3. I have tested with multiple versions of cuda and got the same results. I am reporting 12.3 here because that is the newest cuda version that both clang 19 and 20 compile without warning about cuda version.
To my mind, this all points to a regression from clang 19 in LLVM's cuda support.
![]()
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
