https://github.com/carodewig updated https://github.com/llvm/llvm-project/pull/235910
>From eb596f62e12f410e6541dc020da2b6f82803651d Mon Sep 17 00:00:00 2001 From: carodewig <[email protected]> Date: Mon, 26 Oct 2026 20:45:31 +0000 Subject: [PATCH 1/3] [clang][CUDA] Search include/cccl of CUDA 13 installations >From CUDA 13.0 on, the CUDA C++ Core Libraries (libcu++, CUB and Thrust) are in include/cccl of the CUDA installation rather than in include/. Some of the toolkit's own headers include them: cooperative_groups.h, for example, includes <cuda/std/type_traits>. Clang only searches include/, so with CUDA 13 a file which includes cooperative_groups.h fails to compile: cooperative_groups/details/info.h: fatal error: 'cuda/std/type_traits' file not found nvcc searches include/cccl as a system include directory right after include/. Clang now does the same, for both the device and the host compilation, when the directory exists. As with include/, -nogpuinc leaves it out, and include directories given by the user are searched first, so a separately installed CCCL still takes precedence over the toolkit's. This mirrors what the HIP toolchain does for include/libhipcxx of ROCm. Assisted-by: Devin --- clang/lib/Driver/ToolChains/Cuda.cpp | 22 ++++++++++----- .../Inputs/CUDA_130/usr/local/cuda/bin/.keep | 0 .../CUDA_130/usr/local/cuda/include/.keep | 0 .../usr/local/cuda/include/cccl/.keep | 0 .../CUDA_130/usr/local/cuda/include/cuda.h | 7 +++++ .../Inputs/CUDA_130/usr/local/cuda/lib/.keep | 0 .../CUDA_130/usr/local/cuda/lib64/.keep | 0 .../local/cuda/nvvm/libdevice/libdevice.10.bc | 0 clang/test/Driver/cuda-cccl-include.cu | 27 +++++++++++++++++++ 9 files changed, 50 insertions(+), 6 deletions(-) create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/bin/.keep create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/.keep create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cccl/.keep create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cuda.h create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/lib/.keep create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/lib64/.keep create mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/nvvm/libdevice/libdevice.10.bc create mode 100644 clang/test/Driver/cuda-cccl-include.cu diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp index acbbbb3b72c048..06aa6b13f7d964 100644 --- a/clang/lib/Driver/ToolChains/Cuda.cpp +++ b/clang/lib/Driver/ToolChains/Cuda.cpp @@ -1061,12 +1061,22 @@ void CudaToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs, HostTC.AddClangSystemIncludeArgs(DriverArgs, CC1Args); - if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc, - true) && - CudaInstallation.isValid()) - CC1Args.append( - {"-internal-isystem", - DriverArgs.MakeArgString(CudaInstallation.getIncludePath())}); + if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc, + true)) + return; + + if (!CudaInstallation.isValid()) + return; + + StringRef IncludePath = CudaInstallation.getIncludePath(); + CC1Args.append({"-internal-isystem", DriverArgs.MakeArgString(IncludePath)}); + + // From CUDA 13.0 on, libcu++, CUB and Thrust are in include/cccl, and some of + // the toolkit's own headers, such as cooperative_groups.h, include them. + SmallString<128> CCCLPath(IncludePath); + llvm::sys::path::append(CCCLPath, "cccl"); + if (getDriver().getVFS().exists(CCCLPath)) + CC1Args.append({"-internal-isystem", DriverArgs.MakeArgString(CCCLPath)}); } void CudaToolChain::AddClangCXXStdlibIncludeArgs(const ArgList &Args, diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/bin/.keep b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/bin/.keep new file mode 100644 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/.keep b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/.keep new file mode 100644 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cccl/.keep b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cccl/.keep new file mode 100644 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cuda.h b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cuda.h new file mode 100644 index 00000000000000..9ddfd9b011fdac --- /dev/null +++ b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/cuda.h @@ -0,0 +1,7 @@ +// +// Placeholder file for testing CUDA version detection +// + +#define CUDA_VERSION 13000 + +// diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/lib/.keep b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/lib/.keep new file mode 100644 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/lib64/.keep b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/lib64/.keep new file mode 100644 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/nvvm/libdevice/libdevice.10.bc b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/nvvm/libdevice/libdevice.10.bc new file mode 100644 index 00000000000000..e69de29bb2d1d6 diff --git a/clang/test/Driver/cuda-cccl-include.cu b/clang/test/Driver/cuda-cccl-include.cu new file mode 100644 index 00000000000000..6b2613dbf69a65 --- /dev/null +++ b/clang/test/Driver/cuda-cccl-include.cu @@ -0,0 +1,27 @@ +// From CUDA 13.0 on, the CUDA C++ Core Libraries are in include/cccl of the +// CUDA installation. Check that it is a system include directory, right after +// include/, for both the device and the host compilation. + +// RUN: %clang -### --target=x86_64-unknown-linux-gnu --cuda-gpu-arch=sm_75 \ +// RUN: --cuda-path=%S/Inputs/CUDA_130/usr/local/cuda %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=CCCL + +// CCCL: "-cc1" "-triple" "nvptx64-nvidia-cuda" +// CCCL-SAME: "-internal-isystem" "[[CUDA:[^"]*Inputs/CUDA_130/usr/local/cuda]]/include" "-internal-isystem" "[[CUDA]]/include{{[/\\]+}}cccl" +// CCCL: "-cc1" "-triple" "x86_64-unknown-linux-gnu" +// CCCL-SAME: "-internal-isystem" "[[CUDA]]/include" "-internal-isystem" "[[CUDA]]/include{{[/\\]+}}cccl" + +// Older CUDA installations have no include/cccl. +// RUN: %clang -### --target=x86_64-unknown-linux-gnu --cuda-gpu-arch=sm_75 \ +// RUN: --cuda-path=%S/Inputs/CUDA_111/usr/local/cuda %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NOCCCL + +// -nogpuinc leaves out all of the CUDA installation's include directories. +// RUN: %clang -### --target=x86_64-unknown-linux-gnu --cuda-gpu-arch=sm_75 \ +// RUN: -nogpuinc --cuda-path=%S/Inputs/CUDA_130/usr/local/cuda %s 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NOCCCL + +// NOCCCL: "-cc1" "-triple" "nvptx64-nvidia-cuda" +// NOCCCL-NOT: include{{[/\\]+}}cccl" +// NOCCCL: "-cc1" "-triple" "x86_64-unknown-linux-gnu" +// NOCCCL-NOT: include{{[/\\]+}}cccl" >From da4334419e6be66b3b16620d3ff17c13f9e6a373 Mon Sep 17 00:00:00 2001 From: carodewig <[email protected]> Date: Mon, 26 Oct 2026 21:29:51 +0000 Subject: [PATCH 2/3] Add a release note, and trim the diff Add a release note under CUDA Support. Keep the existing condition rather than splitting it into early returns, and drop include/.keep from the fake CUDA 13.0 installation, as cuda.h is already there. Assisted-by: Devin --- clang/docs/ReleaseNotes.md | 6 ++++ clang/lib/Driver/ToolChains/Cuda.cpp | 30 +++++++++---------- .../CUDA_130/usr/local/cuda/include/.keep | 0 3 files changed, 20 insertions(+), 16 deletions(-) delete mode 100644 clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/.keep diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index 200c41fb9e4a29..8c9b3d02a7fc51 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -1337,6 +1337,12 @@ documentation](CIR/index.md) for more information regarding ClangIR's status and functions with CUDA 13.4 and newer. Added `__nvvm_idp4a_*` and `__nvvm_idp2a_*` builtins. (#GH233434) +- CUDA compilations now add the `include/cccl` directory of the selected CUDA + installation to the header search path when it exists. From CUDA 13.0 on, + libcu++, CUB and Thrust are in that directory, and toolkit headers such as + `cooperative_groups.h` include them. The `-nogpuinc` option disables this + path together with the other CUDA include paths. + #### AIX Support #### NetBSD Support diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp index 06aa6b13f7d964..677da1f85fed2b 100644 --- a/clang/lib/Driver/ToolChains/Cuda.cpp +++ b/clang/lib/Driver/ToolChains/Cuda.cpp @@ -1061,22 +1061,20 @@ void CudaToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs, HostTC.AddClangSystemIncludeArgs(DriverArgs, CC1Args); - if (!DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc, - true)) - return; - - if (!CudaInstallation.isValid()) - return; - - StringRef IncludePath = CudaInstallation.getIncludePath(); - CC1Args.append({"-internal-isystem", DriverArgs.MakeArgString(IncludePath)}); - - // From CUDA 13.0 on, libcu++, CUB and Thrust are in include/cccl, and some of - // the toolkit's own headers, such as cooperative_groups.h, include them. - SmallString<128> CCCLPath(IncludePath); - llvm::sys::path::append(CCCLPath, "cccl"); - if (getDriver().getVFS().exists(CCCLPath)) - CC1Args.append({"-internal-isystem", DriverArgs.MakeArgString(CCCLPath)}); + if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc, + true) && + CudaInstallation.isValid()) { + StringRef IncludePath = CudaInstallation.getIncludePath(); + CC1Args.append( + {"-internal-isystem", DriverArgs.MakeArgString(IncludePath)}); + + // From CUDA 13.0 on, libcu++, CUB and Thrust are in include/cccl, and some + // of the toolkit's own headers, such as cooperative_groups.h, include them. + SmallString<128> CCCLPath(IncludePath); + llvm::sys::path::append(CCCLPath, "cccl"); + if (getDriver().getVFS().exists(CCCLPath)) + CC1Args.append({"-internal-isystem", DriverArgs.MakeArgString(CCCLPath)}); + } } void CudaToolChain::AddClangCXXStdlibIncludeArgs(const ArgList &Args, diff --git a/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/.keep b/clang/test/Driver/Inputs/CUDA_130/usr/local/cuda/include/.keep deleted file mode 100644 index e69de29bb2d1d6..00000000000000 >From 10cbc1193b32d6f4dba0b685465bef4f0aa11618 Mon Sep 17 00:00:00 2001 From: Caroline Rodewig <[email protected]> Date: Thu, 29 Oct 2026 15:52:42 -0400 Subject: [PATCH 3/3] test: allow backslash for include path Co-authored-by: Ambrose Leeb <[email protected]> --- clang/test/Driver/cuda-cccl-include.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/Driver/cuda-cccl-include.cu b/clang/test/Driver/cuda-cccl-include.cu index 6b2613dbf69a65..8916bbb7f7f652 100644 --- a/clang/test/Driver/cuda-cccl-include.cu +++ b/clang/test/Driver/cuda-cccl-include.cu @@ -7,9 +7,9 @@ // RUN: | FileCheck %s --check-prefix=CCCL // CCCL: "-cc1" "-triple" "nvptx64-nvidia-cuda" -// CCCL-SAME: "-internal-isystem" "[[CUDA:[^"]*Inputs/CUDA_130/usr/local/cuda]]/include" "-internal-isystem" "[[CUDA]]/include{{[/\\]+}}cccl" +// CCCL-SAME: "-internal-isystem" "[[CUDA:[^"]*Inputs/CUDA_130/usr/local/cuda]]{{[/\\]+}}include" "-internal-isystem" "[[CUDA]]/include{{[/\\]+}}cccl" // CCCL: "-cc1" "-triple" "x86_64-unknown-linux-gnu" -// CCCL-SAME: "-internal-isystem" "[[CUDA]]/include" "-internal-isystem" "[[CUDA]]/include{{[/\\]+}}cccl" +// CCCL-SAME: "-internal-isystem" "[[CUDA]]{{[/\\]+}}include" "-internal-isystem" "[[CUDA]]{{[/\\]+}}include{{[/\\]+}}cccl" // Older CUDA installations have no include/cccl. // RUN: %clang -### --target=x86_64-unknown-linux-gnu --cuda-gpu-arch=sm_75 \ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
