NVIDIA / NVIDIA/cccl

[BUG]: Calling cub::DeviceRadixSort::SortKeys fails with invalid device function

Open
#4,650 2 comments 0 reactions 1 assignee Claimed by @elstehle 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

Runtime Error

### Component

CUB

### Describe the bug

calling `cub::DeviceRadixSort::SortKeys` fails with `invalid device function`.

This happens since update to CUDA 12.9, reverting to CUDA 12.8 leads to other CUB related errors.

### How to Reproduce

```cpp
#include
#include
#include
#include

#define CUDA_CHECK(call) \
do \
{ \
cudaError_t err = call; \
if (err != cudaSuccess) \
{ \
std::fprintf(stderr, \
"CUDA error at %s:%d:\n" \
" call: %s\n" \
" code: %d\n" \
" reason: %s\n", \
__FILE__, __LINE__, #call, static_cast(err), \
cudaGetErrorString(err)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)

int main()
{
cudaStream_t stream = 0;
CUDA_CHECK(cudaStreamCreate(&stream));

void *d_temp_storage = nullptr;
size_t temp_storage_bytes = 0;
const float *d_keys_in = nullptr;
float *d_keys_out = nullptr;
size_t num_items = 512 * 320;
int begin_bit = 0;
int end_bit = 32;

CUDA_CHECK(cub::DeviceRadixSort::SortKeys(
d_temp_storage,
temp_storage_bytes,
d_keys_in,
d_keys_out,
num_items,
begin_bit,
end_bit,
stream));

return 0;
}
```

```cmake
cmake_minimum_required(VERSION 3.18)
project(bug LANGUAGES CUDA CXX)

add_executable(bug main.cu)

set_target_properties(bug PROPERTIES
CXX_STANDARD 20
CXX_STANDARD_REQUIRED YES
CXX_EXTENSIONS NO
CUDA_STANDARD 20
CUDA_STANDARD_REQUIRED YES
CUDA_SEPARABLE_COMPILATION ON
CUDA_RESOLVE_DEVICE_SYMBOLS ON
)

target_link_libraries(bug cudart)
```

```cmd
C:\Users\test\Documents\bug>build\Debug\bug.exe
CUDA error 98 [C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.9\include\cub/device/dispatch/dispatch_radix_sort.cuh, 1874]: invalid device function
CUDA error at C:\Users\test\Documents\bug\main.cu:44:
call: cub::DeviceRadixSort::SortKeys( d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, num_items, begin_bit, end_bit, stream)
code: 98
reason: invalid device function
```

The error seems to occure at `cub\util_device.cuh:732`:
```cpp
template
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t runtime_to_compiletime(int device_ptx_version, FunctorT& op)
{
// We instantiate invoke_static for each CudaArches, but only call the one matching device_ptx_version.
// If there's no exact match of the architectures in __CUDA_ARCH_LIST__/NV_TARGET_SM_INTEGER_LIST and the runtime
// queried ptx version (i.e., the closest ptx version to the current device's architecture that the EmptyKernel was
// compiled for), we return cudaErrorInvalidDeviceFunction. Such a scenario may arise if CUB_DISABLE_NAMESPACE_MAGIC
// is set and different TUs are compiled for different sets of architecture.
cudaError_t e = cudaErrorInvalidDeviceFunction;
const cudaError_t dummy[] = {
(device_ptx_version == (CudaArches * ArchMult)
? (e = invoke_static<(CudaArches * ArchMult)>(op, ::cuda::std::true_type{}))
: cudaSuccess)...};
(void) dummy;
return e;
}
```

### Expected behavior

Calling `cub::DeviceRadixSort::SortKeys` should not fail and return the `temp_storage_bytes` required.

### Reproduction link

_No response_

### Operating System

Windows 11 Pro

### nvidia-smi output

+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 576.02 Driver Version: 576.02 CUDA Version: 12.9 |
|-----------------------------------------+------------------------+----------------------+
| 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 3090 WDDM | 00000000:01:00.0 On | N/A |
| 0% 37C P0 110W / 350W | 1103MiB / 24576MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

### NVCC version

nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Apr__9_19:29:17_Pacific_Daylight_Time_2025
Cuda compilation tools, release 12.9, V12.9.41
Build cuda_12.9.r12.9/compiler.35813241_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.