[BUG]: #include <thrust/device_vector.h> causes compile error

[Olli1080]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

Thrust

Describe the bug

When including thrust/device_vector in anywhere visible to a cpp compiled file it fails due to including device only code.
This issue is present in thrust shipped with CUDA 12.5 but was non existent in CUDA 12.3.

How to Reproduce

1. Include #include <thrust/device_vector.h> in any file related to c++ compiler.
2. Compile
3. C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\include\cub\util_ptx.cuh(90,29): error C2059: Syntaxerror: ":" (20 more lines omitted)

Expected behavior

1. Include #include <thrust/device_vector.h> in any file related to c++ compiler.
2. Compile
3. Compilation is succesfull

Reproduction link

No response

Operating System

Windows 11 Pro 10.0.22621

nvidia-smi output

+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 555.85 Driver Version: 555.85 CUDA Version: 12.5 |
|-----------------------------------------+------------------------+----------------------+
| 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 GTX 970 WDDM | 00000000:01:00.0 On | N/A |
| 36% 41C P0 51W / 179W | 1123MiB / 4096MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Wed_Apr_17_19:36:51_Pacific_Daylight_Time_2024
Cuda compilation tools, release 12.5, V12.5.40
Build cuda_12.5.r12.5/compiler.34177558_0

[jrhemstad]

Thanks for raising this @Olli1080.

Unfortunately, this is expected. See #1374

TL;DR: It was never intended that thrust/device_vector.h (or any Thrust header for that matter) would work with a host-only compiler when the device system is configured to be CUDA.

Out of curiosity, in your use case, what did you expect to be able to do with the thrust/device_vector.h header in a host-only translation unit? Simply include the header but not use anything from it?

[Olli1080]

Thank you for the clarification.
I've been trying very hard to find such guarantees.

My use case is to have thrust/device_vector as a private member of a class which can be consumed by C++ translation units e.g. outside of a library but to compile the definition of the class using CUDA.
That way i can cache certain data from e.g. point clouds in device_vectors but read out results to either other CUDA units or through methods as host_vectors.
This also makes it way easier to handle storage and to debug.

TL;DR: Exposing consumption of classes in host-only code, but also having the functionality inside the CUDA translation unit.

[Olli1080]

For example here is a one such class i'm compiling with cuda but consuming in host-only code.
https://github.com/Olli1080/gpu-voxels/blob/ar_integration/packages/gpu_voxels/src/gpu_voxels/helpers/MetaPointCloud.h
https://github.com/Olli1080/gpu-voxels/blob/ar_integration/packages/gpu_voxels/src/gpu_voxels/helpers/MetaPointCloud.cu

[bernhardmgruber]

@Olli1080 you can apply the PIMPL idiom to workaround this. Something along these lines:

Header:

class MetaPointCloud {
  struct Impl;
  std::unique_ptr<Impl> pimpl;
};

Source:

struct MetaPointCloud::Impl {
  thrust::device_vector<...> ...;
};

Make sure all constructors and destructors of MetaPointCloud are also defined in the source file. In C++26 you can switch std::unique_ptr for std::indirect. The corresponding paper has a section on PIMPL.

[Olli1080]

I had some digging around in the code (using /showincludes) and found out that the errors are caused by the changes in <cub/util_device.cuh>, namely adding <cub/util_ptx.cuh> and <cuda/discard_memory>. Looking at the main branch both of these includes are now gone. I'll test if including host_vector or device_vector works again asap.

For anyone interested the relevant include tree is:

CUDA\thrust/host_vector.h
CUDA\thrust/detail/vector_base.h
CUDA\thrust/detail/contiguous_storage.h
CUDA\thrust/detail/contiguous_storage.inl
CUDA\thrust/detail/allocator/copy_construct_range.h
CUDA\thrust/detail/allocator/copy_construct_range.inl
CUDA\thrust/detail/copy.h
CUDA\thrust/detail/copy.inl
CUDA\thrust/system/detail/adl/copy.h
CUDA\thrust/system/cuda/detail/copy.h
CUDA\thrust/system/cuda/detail/internal/copy_cross_system.h
CUDA\thrust/system/cuda/detail/util.h
CUDA\cub/util_device.cuh
CUDA\cub/util_ptx.cuh

[Olli1080]

I've tested CUDA 12.5 with replaced cccl headers and the issue does not exist anymore.
Commit as of this comment is bf1a71a318a3dbc4ee522a42e2b82b6d2a5410e4