Author: Corentin Kerisit Date: 2026-04-15T12:16:38-07:00 New Revision: cdfd0b60d2f30d91087227ea90f16e43af3ff875
URL: https://github.com/llvm/llvm-project/commit/cdfd0b60d2f30d91087227ea90f16e43af3ff875 DIFF: https://github.com/llvm/llvm-project/commit/cdfd0b60d2f30d91087227ea90f16e43af3ff875.diff LOG: [CUDA] Change __CUDACC__ definition to 1 (#189457) I recently encountered an issue where `nccl` used `#if __CUDACC__` , assuming `__CUDACC__` is not only defined but having a #if-able value. https://github.com/NVIDIA/nccl/blob/v2.28.3-1/src/include/nccl_device/coop.h#L18 Looking at nvcc invocation, I see that: ``` echo "" | nvcc -x cu -E -Xcompiler -dM - | grep __CUDACC__ #define __CUDACC__ 1 ``` Changing __CUDACC__ to 1 to match what NVIDIA downstream code assumptions. Added: Modified: clang/lib/Headers/__clang_cuda_runtime_wrapper.h Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 295f4191f9927..da3c92f33539c 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -34,7 +34,7 @@ // Define __CUDACC__ early as libstdc++ standard headers with GNU extensions // enabled depend on it to avoid using __float128, which is unsupported in // CUDA. -#define __CUDACC__ +#define __CUDACC__ 1 // Include some standard headers to avoid CUDA headers including them // while some required macros (like __THROW) are in a weird state. @@ -103,7 +103,7 @@ #if CUDA_VERSION < 9000 #define __CUDABE__ #else -#define __CUDACC__ +#define __CUDACC__ 1 #define __CUDA_LIBDEVICE__ #endif // Disables definitions of device-side runtime support stubs in @@ -120,7 +120,7 @@ #define nv_weak weak #undef __CUDABE__ #undef __CUDA_LIBDEVICE__ -#define __CUDACC__ +#define __CUDACC__ 1 #include "cuda_runtime.h" #pragma pop_macro("nv_weak") @@ -266,7 +266,7 @@ static inline __device__ void __brkpt(int __c) { __brkpt(); } // branch of #if/else inside. #define __host__ #undef __CUDABE__ -#define __CUDACC__ +#define __CUDACC__ 1 #if CUDA_VERSION >= 9000 // Some atomic functions became compiler builtins in CUDA-9 , so we need their // declarations. @@ -412,7 +412,7 @@ __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr, // Set up compiler macros expected to be seen during compilation. #undef __CUDABE__ -#define __CUDACC__ +#define __CUDACC__ 1 extern "C" { // Device-side CUDA system calls. _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
