[clang] cdfd0b6 - [CUDA] Change __CUDACC__ definition to 1 (#189457)

[via cfe-commits]

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.