From: Justin Lebar Date: Sat, 8 Oct 2016 22:16:08 +0000 (+0000) Subject: [CUDA] Rename cuda_builtin_vars.h to __clang_cuda_builtin_vars.h. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=204eb555222d86837d865007c1526018e391bf67;p=clang [CUDA] Rename cuda_builtin_vars.h to __clang_cuda_builtin_vars.h. Summary: This matches the idiom we use for our other CUDA wrapper headers. Reviewers: tra Subscribers: beanz, mgorny, cfe-commits Differential Revision: https://reviews.llvm.org/D24978 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283679 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp index 0dd05b5d9a..cf39f35701 100644 --- a/lib/Frontend/CompilerInvocation.cpp +++ b/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/lib/Headers/CMakeLists.txt b/lib/Headers/CMakeLists.txt index 958038e0cd..be18ea8dbf 100644 --- a/lib/Headers/CMakeLists.txt +++ b/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/lib/Headers/cuda_builtin_vars.h b/lib/Headers/__clang_cuda_builtin_vars.h similarity index 100% rename from lib/Headers/cuda_builtin_vars.h rename to lib/Headers/__clang_cuda_builtin_vars.h diff --git a/lib/Headers/__clang_cuda_runtime_wrapper.h b/lib/Headers/__clang_cuda_runtime_wrapper.h index 0cf8b17def..6c6dff86ad 100644 --- a/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/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/test/CodeGenCUDA/cuda-builtin-vars.cu b/test/CodeGenCUDA/cuda-builtin-vars.cu index c2159f5af1..c1edff936a 100644 --- a/test/CodeGenCUDA/cuda-builtin-vars.cu +++ b/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/test/SemaCUDA/cuda-builtin-vars.cu b/test/SemaCUDA/cuda-builtin-vars.cu index 108e75cae7..27a9c5abd7 100644 --- a/test/SemaCUDA/cuda-builtin-vars.cu +++ b/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