diff options
| -rw-r--r-- | clang/lib/Frontend/CompilerInvocation.cpp | 7 | ||||
| -rw-r--r-- | clang/lib/Headers/CMakeLists.txt | 2 | ||||
| -rw-r--r-- | clang/lib/Headers/__clang_cuda_builtin_vars.h (renamed from clang/lib/Headers/cuda_builtin_vars.h) | 0 | ||||
| -rw-r--r-- | clang/lib/Headers/__clang_cuda_runtime_wrapper.h | 16 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/cuda-builtin-vars.cu | 2 | ||||
| -rw-r--r-- | clang/test/SemaCUDA/cuda-builtin-vars.cu | 12 |
6 files changed, 20 insertions, 19 deletions
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index 0dd05b5d9a9..cf39f357012 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2012,9 +2012,10 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, // enabled for Microsoft Extensions or Borland Extensions, here. // // FIXME: __declspec is also currently enabled for CUDA, but isn't really a - // CUDA extension, however it is required for supporting cuda_builtin_vars.h, - // which uses __declspec(property). Once that has been rewritten in terms of - // something more generic, remove the Opts.CUDA term here. + // CUDA extension. However, it is required for supporting + // __clang_cuda_builtin_vars.h, which uses __declspec(property). Once that has + // been rewritten in terms of something more generic, remove the Opts.CUDA + // term here. Opts.DeclSpecKeyword = Args.hasFlag(OPT_fdeclspec, OPT_fno_declspec, (Opts.MicrosoftExt || Opts.Borland || Opts.CUDA)); diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index 958038e0cda..be18ea8dbf5 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -22,12 +22,12 @@ set(files avxintrin.h bmi2intrin.h bmiintrin.h + __clang_cuda_builtin_vars.h __clang_cuda_cmath.h __clang_cuda_intrinsics.h __clang_cuda_math_forward_declares.h __clang_cuda_runtime_wrapper.h cpuid.h - cuda_builtin_vars.h clflushoptintrin.h emmintrin.h f16cintrin.h diff --git a/clang/lib/Headers/cuda_builtin_vars.h b/clang/lib/Headers/__clang_cuda_builtin_vars.h index 6f5eb9c78d8..6f5eb9c78d8 100644 --- a/clang/lib/Headers/cuda_builtin_vars.h +++ b/clang/lib/Headers/__clang_cuda_builtin_vars.h diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 0cf8b17def5..6c6dff86adc 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -72,9 +72,9 @@ #define __CUDA_ARCH__ 350 #endif -#include "cuda_builtin_vars.h" +#include "__clang_cuda_builtin_vars.h" -// No need for device_launch_parameters.h as cuda_builtin_vars.h above +// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above // has taken care of builtin variables declared in the file. #define __DEVICE_LAUNCH_PARAMETERS_H__ @@ -283,8 +283,8 @@ __device__ static inline void *malloc(size_t __size) { } } // namespace std -// Out-of-line implementations from cuda_builtin_vars.h. These need to come -// after we've pulled in the definition of uint3 and dim3. +// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to +// come after we've pulled in the definition of uint3 and dim3. __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { uint3 ret; @@ -315,10 +315,10 @@ __device__ inline __cuda_builtin_gridDim_t::operator dim3() const { // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host // mode, giving them their "proper" types of dim3 and uint3. This is -// incompatible with the types we give in cuda_builtin_vars.h. As as hack, -// force-include the header (nvcc doesn't include it by default) but redefine -// dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are only -// used here for the redeclarations of blockDim and threadIdx.) +// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as +// hack, force-include the header (nvcc doesn't include it by default) but +// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are +// only used here for the redeclarations of blockDim and threadIdx.) #pragma push_macro("dim3") #pragma push_macro("uint3") #define dim3 __cuda_builtin_blockDim_t diff --git a/clang/test/CodeGenCUDA/cuda-builtin-vars.cu b/clang/test/CodeGenCUDA/cuda-builtin-vars.cu index c2159f5af14..c1edff936a0 100644 --- a/clang/test/CodeGenCUDA/cuda-builtin-vars.cu +++ b/clang/test/CodeGenCUDA/cuda-builtin-vars.cu @@ -1,6 +1,6 @@ // RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s -#include "cuda_builtin_vars.h" +#include "__clang_cuda_builtin_vars.h" // CHECK: define void @_Z6kernelPi(i32* %out) __attribute__((global)) diff --git a/clang/test/SemaCUDA/cuda-builtin-vars.cu b/clang/test/SemaCUDA/cuda-builtin-vars.cu index 108e75cae76..27a9c5abd70 100644 --- a/clang/test/SemaCUDA/cuda-builtin-vars.cu +++ b/clang/test/SemaCUDA/cuda-builtin-vars.cu @@ -1,6 +1,6 @@ // RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s -#include "cuda_builtin_vars.h" +#include "__clang_cuda_builtin_vars.h" __attribute__((global)) void kernel(int *out) { int i = 0; @@ -34,20 +34,20 @@ void kernel(int *out) { out[i++] = warpSize; warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}} - // expected-note@cuda_builtin_vars.h:* {{variable 'warpSize' declared const here}} + // expected-note@__clang_cuda_builtin_vars.h:* {{variable 'warpSize' declared const here}} // Make sure we can't construct or assign to the special variables. __cuda_builtin_threadIdx_t x; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}} - // expected-note@cuda_builtin_vars.h:* {{declared private here}} + // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}} __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}} - // expected-note@cuda_builtin_vars.h:* {{declared private here}} + // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}} threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_threadIdx_t'}} - // expected-note@cuda_builtin_vars.h:* {{declared private here}} + // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}} void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_threadIdx_t'}} - // expected-note@cuda_builtin_vars.h:* {{declared private here}} + // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}} // Following line should've caused an error as one is not allowed to // take address of a built-in variable in CUDA. Alas there's no way |
