[BUG]: cub::DeviceCopy mdspan path requires cuda::cuda_error declaration in HostJIT freestanding mode
- 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
Compile-time Error
### Component
CUB
### Describe the bug
When compiling a HostJIT-generated translation unit that `#include ` and calls `cub::DeviceCopy::Copy` on `cuda::std::mdspan` inputs, compilation fails in HostJIT freestanding mode because CUB references `::cuda::cuda_error` in an `_CCCL_CATCH` branch. In this mode `_CCCL_HOSTJIT() && !_CCCL_HOSTED()` is true, and libcudacxx does not define `cuda::cuda_error`.
**Observed compilation error**
```
error: no type named 'cuda_error' in namespace 'cuda'
```
This seems like a CCCL issue because `_CCCL_CATCH(...)` expands to a disabled/unreachable branch when exceptions are unavailable, but the catch declaration still has to be syntactically valid. Since `cuda::cuda_error` is only defined for `_CCCL_HOSTED()`, CUB’s freestanding HostJIT path references a type that is unavailable.
**Current workaround**
User code can inject a dummy declaration in generated HostJIT source:
```cpp
#if _CCCL_HOSTJIT() && !_CCCL_HOSTED()
namespace cuda {
class cuda_error {
public:
_CCCL_HOST_API ::cudaError_t status() const noexcept { return ::cudaErrorUnknown; }
};
}
#endif
```
but this should not be necessary.
### How to Reproduce
```cpp
#include
#include
#include
#include
namespace
{
constexpr int target_sm = 89;
constexpr bool inject_cuda_error_workaround = false;
constexpr bool keep_hostjit_artifacts = true;
constexpr const char* generated_function_name = "repro_device_copy";
std::string generated_source()
{
std::string source = R"cpp(#include
#include
#include
#include
#include
#include
)cpp";
if constexpr (inject_cuda_error_workaround)
{
source += R"cpp(#if _CCCL_HOSTJIT() && !_CCCL_HOSTED()
namespace cuda
{
class cuda_error
{
public:
_CCCL_HOST_API ::cudaError_t status() const noexcept { return ::cudaErrorUnknown; }
};
} // namespace cuda
#endif
)cpp";
}
source += R"cpp(#include
extern "C" _CCCL_VISIBILITY_EXPORT int repro_device_copy(
void* source_base,
void* destination_base,
const long long* shape,
void* stream)
{
using element_t = cuda::std::uint32_t;
using extents_t = cuda::std::extents;
using mdspan_t = cuda::std::mdspan;
extents_t extents{static_cast(shape[0])};
mdspan_t source{static_cast(source_base), extents};
mdspan_t destination{static_cast(destination_base), extents};
const auto status = cub::DeviceCopy::Copy(
source,
destination,
cuda::stream_ref{static_cast(stream)});
return static_cast(status);
}
)cpp";
return source;
}
} // namespace
int main()
{
auto config = hostjit::detectDefaultConfig();
config.entry_point_name = generated_function_name;
config.sm_version = target_sm;
config.optimization_level = 3;
config.keep_artifacts = keep_hostjit_artifacts;
std::string validation_error;
if (!hostjit::validateConfig(config, &validation_error))
{
std::cerr << "invalid HostJIT config: " << validation_error << "\n";
return 2;
}
const auto source = generated_source();
hostjit::JITCompiler compiler{config};
if (!compiler.compile(source))
{
std::cerr << "HostJIT compilation failed:\n" << compiler.getLastError() << "\n";
return 1;
}
using function_t = int (*)(void*, void*, const long long*, void*);
const auto function = compiler.getFunction(generated_function_name);
if (function == nullptr)
{
std::cerr << "compiled library did not export " << generated_function_name << ":\n"
<< compiler.getLastError() << "\n";
return 1;
}
std::cout << "HostJIT compilation succeeded.\n";
std::cout << "artifacts_path: " << compiler.getArtifactsPath() << "\n";
return 0;
}
```
### Expected behavior
CUB/libcudacxx should compile this path in HostJIT freestanding mode without user-provided declarations. Possible fixes might include guarding the `cuda::cuda_error` catch, making the exception type available as a lightweight freestanding declaration, or adjusting _CCCL_CATCH usage so unavailable exception types are not parsed in no-exception/freestanding mode.
### Reproduction link
_No response_
### Operating System
Ubuntu Linux 24.04
### nvidia-smi output
```
(ccclrt-dev) opavlyk@NV-22T4X34:~/repos/array/examples/array_copy_kernels/repro/cub_device_copy_hostjit_cuda_error$ nvidia-smi
Thu Aug 20 08:36:51 2026
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 580.112 Driver Version: 581.95 CUDA Version: 13.0 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA RTX 3500 Ada Gene... On | 00000000:01:00.0 Off | Off |
| N/A 43C P3 18W / 74W | 0MiB / 12282MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
+-----------------------------------------------------------------------------------------+
```
### NVCC version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2026 NVIDIA Corporation
Built on Tue_Jun_09_02:43:40_PM_PDT_2026
Cuda compilation tools, release 13.3, V13.3.73
Build cuda_13.3.r13.3/compiler.38244171_0
Contributor guide
Research direction
Start with cub/device/device_copy.cuh and trace the mdspan DeviceCopy path to the _CCCL_CATCH branch that references cuda::cuda_error. Compare that path with the freestanding HostJIT conditions and the declaration in cuda/std/__exception/cuda_error.h. Done means the provided HostJIT reproduction compiles without injecting a user-defined cuda_error declaration.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 55/100