Index: cfe/trunk/lib/Frontend/CompilerInvocation.cpp =================================================================== --- cfe/trunk/lib/Frontend/CompilerInvocation.cpp +++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp @@ -2012,9 +2012,10 @@ // 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)); Index: cfe/trunk/lib/Headers/CMakeLists.txt =================================================================== --- cfe/trunk/lib/Headers/CMakeLists.txt +++ cfe/trunk/lib/Headers/CMakeLists.txt @@ -22,12 +22,12 @@ 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 Index: cfe/trunk/lib/Headers/__clang_cuda_builtin_vars.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_builtin_vars.h +++ cfe/trunk/lib/Headers/__clang_cuda_builtin_vars.h @@ -0,0 +1,126 @@ +/*===---- cuda_builtin_vars.h - CUDA built-in variables ---------------------=== + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + *===-----------------------------------------------------------------------=== + */ + +#ifndef __CUDA_BUILTIN_VARS_H +#define __CUDA_BUILTIN_VARS_H + +// Forward declares from vector_types.h. +struct uint3; +struct dim3; + +// The file implements built-in CUDA variables using __declspec(property). +// https://msdn.microsoft.com/en-us/library/yhfk0thd.aspx +// All read accesses of built-in variable fields get converted into calls to a +// getter function which in turn calls the appropriate builtin to fetch the +// value. +// +// Example: +// int x = threadIdx.x; +// IR output: +// %0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #3 +// PTX output: +// mov.u32 %r2, %tid.x; + +#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC) \ + __declspec(property(get = __fetch_builtin_##FIELD)) unsigned int FIELD; \ + static inline __attribute__((always_inline)) \ + __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) { \ + return INTRINSIC; \ + } + +#if __cplusplus >= 201103L +#define __DELETE =delete +#else +#define __DELETE +#endif + +// Make sure nobody can create instances of the special varible types. nvcc +// also disallows taking address of special variables, so we disable address-of +// operator as well. +#define __CUDA_DISALLOW_BUILTINVAR_ACCESS(TypeName) \ + __attribute__((device)) TypeName() __DELETE; \ + __attribute__((device)) TypeName(const TypeName &) __DELETE; \ + __attribute__((device)) void operator=(const TypeName &) const __DELETE; \ + __attribute__((device)) TypeName *operator&() const __DELETE + +struct __cuda_builtin_threadIdx_t { + __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_tid_x()); + __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_tid_y()); + __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_tid_z()); + // threadIdx should be convertible to uint3 (in fact in nvcc, it *is* a + // uint3). This function is defined after we pull in vector_types.h. + __attribute__((device)) operator uint3() const; +private: + __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_threadIdx_t); +}; + +struct __cuda_builtin_blockIdx_t { + __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ctaid_x()); + __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ctaid_y()); + __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ctaid_z()); + // blockIdx should be convertible to uint3 (in fact in nvcc, it *is* a + // uint3). This function is defined after we pull in vector_types.h. + __attribute__((device)) operator uint3() const; +private: + __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockIdx_t); +}; + +struct __cuda_builtin_blockDim_t { + __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ntid_x()); + __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ntid_y()); + __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ntid_z()); + // blockDim should be convertible to dim3 (in fact in nvcc, it *is* a + // dim3). This function is defined after we pull in vector_types.h. + __attribute__((device)) operator dim3() const; +private: + __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockDim_t); +}; + +struct __cuda_builtin_gridDim_t { + __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_nctaid_x()); + __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_nctaid_y()); + __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_nctaid_z()); + // gridDim should be convertible to dim3 (in fact in nvcc, it *is* a + // dim3). This function is defined after we pull in vector_types.h. + __attribute__((device)) operator dim3() const; +private: + __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_gridDim_t); +}; + +#define __CUDA_BUILTIN_VAR \ + extern const __attribute__((device)) __attribute__((weak)) +__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx; +__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx; +__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim; +__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim; + +// warpSize should translate to read of %WARP_SZ but there's currently no +// builtin to do so. According to PTX v4.2 docs 'to date, all target +// architectures have a WARP_SZ value of 32'. +__attribute__((device)) const int warpSize = 32; + +#undef __CUDA_DEVICE_BUILTIN +#undef __CUDA_BUILTIN_VAR +#undef __CUDA_DISALLOW_BUILTINVAR_ACCESS + +#endif /* __CUDA_BUILTIN_VARS_H */ Index: cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h =================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_runtime_wrapper.h +++ cfe/trunk/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 @@ } } // 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 @@ // 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 Index: cfe/trunk/lib/Headers/cuda_builtin_vars.h =================================================================== --- cfe/trunk/lib/Headers/cuda_builtin_vars.h +++ cfe/trunk/lib/Headers/cuda_builtin_vars.h @@ -1,126 +0,0 @@ -/*===---- cuda_builtin_vars.h - CUDA built-in variables ---------------------=== - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - * - *===-----------------------------------------------------------------------=== - */ - -#ifndef __CUDA_BUILTIN_VARS_H -#define __CUDA_BUILTIN_VARS_H - -// Forward declares from vector_types.h. -struct uint3; -struct dim3; - -// The file implements built-in CUDA variables using __declspec(property). -// https://msdn.microsoft.com/en-us/library/yhfk0thd.aspx -// All read accesses of built-in variable fields get converted into calls to a -// getter function which in turn calls the appropriate builtin to fetch the -// value. -// -// Example: -// int x = threadIdx.x; -// IR output: -// %0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #3 -// PTX output: -// mov.u32 %r2, %tid.x; - -#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC) \ - __declspec(property(get = __fetch_builtin_##FIELD)) unsigned int FIELD; \ - static inline __attribute__((always_inline)) \ - __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) { \ - return INTRINSIC; \ - } - -#if __cplusplus >= 201103L -#define __DELETE =delete -#else -#define __DELETE -#endif - -// Make sure nobody can create instances of the special varible types. nvcc -// also disallows taking address of special variables, so we disable address-of -// operator as well. -#define __CUDA_DISALLOW_BUILTINVAR_ACCESS(TypeName) \ - __attribute__((device)) TypeName() __DELETE; \ - __attribute__((device)) TypeName(const TypeName &) __DELETE; \ - __attribute__((device)) void operator=(const TypeName &) const __DELETE; \ - __attribute__((device)) TypeName *operator&() const __DELETE - -struct __cuda_builtin_threadIdx_t { - __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_tid_x()); - __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_tid_y()); - __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_tid_z()); - // threadIdx should be convertible to uint3 (in fact in nvcc, it *is* a - // uint3). This function is defined after we pull in vector_types.h. - __attribute__((device)) operator uint3() const; -private: - __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_threadIdx_t); -}; - -struct __cuda_builtin_blockIdx_t { - __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ctaid_x()); - __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ctaid_y()); - __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ctaid_z()); - // blockIdx should be convertible to uint3 (in fact in nvcc, it *is* a - // uint3). This function is defined after we pull in vector_types.h. - __attribute__((device)) operator uint3() const; -private: - __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockIdx_t); -}; - -struct __cuda_builtin_blockDim_t { - __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ntid_x()); - __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ntid_y()); - __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ntid_z()); - // blockDim should be convertible to dim3 (in fact in nvcc, it *is* a - // dim3). This function is defined after we pull in vector_types.h. - __attribute__((device)) operator dim3() const; -private: - __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockDim_t); -}; - -struct __cuda_builtin_gridDim_t { - __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_nctaid_x()); - __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_nctaid_y()); - __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_nctaid_z()); - // gridDim should be convertible to dim3 (in fact in nvcc, it *is* a - // dim3). This function is defined after we pull in vector_types.h. - __attribute__((device)) operator dim3() const; -private: - __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_gridDim_t); -}; - -#define __CUDA_BUILTIN_VAR \ - extern const __attribute__((device)) __attribute__((weak)) -__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx; -__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx; -__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim; -__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim; - -// warpSize should translate to read of %WARP_SZ but there's currently no -// builtin to do so. According to PTX v4.2 docs 'to date, all target -// architectures have a WARP_SZ value of 32'. -__attribute__((device)) const int warpSize = 32; - -#undef __CUDA_DEVICE_BUILTIN -#undef __CUDA_BUILTIN_VAR -#undef __CUDA_DISALLOW_BUILTINVAR_ACCESS - -#endif /* __CUDA_BUILTIN_VARS_H */ Index: cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu =================================================================== --- cfe/trunk/test/CodeGenCUDA/cuda-builtin-vars.cu +++ cfe/trunk/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)) Index: cfe/trunk/test/SemaCUDA/cuda-builtin-vars.cu =================================================================== --- cfe/trunk/test/SemaCUDA/cuda-builtin-vars.cu +++ cfe/trunk/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 @@ 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