Index: lib/Headers/CMakeLists.txt =================================================================== --- lib/Headers/CMakeLists.txt +++ lib/Headers/CMakeLists.txt @@ -13,6 +13,7 @@ bmi2intrin.h bmiintrin.h cpuid.h + cuda/cuda_builtin_vars.h emmintrin.h f16cintrin.h float.h Index: lib/Headers/cuda/cuda_builtin_vars.h =================================================================== --- /dev/null +++ lib/Headers/cuda/cuda_builtin_vars.h @@ -0,0 +1,94 @@ +/*===---- 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 + +// 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 would call appropriate builtin to fetch the +// value. +// +// Example: +// int x = threadIdx.x; +// IR output: +// %0 = call i32 @llvm.ptx.read.tid.x() #3 +// PTX output: +// mov.u32 %r2, %tid.x; + +#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC) \ + __declspec(property(get = __fetch_builtin_##FIELD)) int FIELD; \ + static inline __attribute__((always_inline)) \ + __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) { \ + return INTRINSIC; \ + } + +struct __cuda_builtin_threadIdx_t { + __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_tid_x()); + __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_tid_y()); + __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_tid_z()); +private: + __cuda_builtin_threadIdx_t() {} +}; + +struct __cuda_builtin_blockIdx_t { + __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_ctaid_x()); + __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_ctaid_y()); + __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_ctaid_z()); +private: + __cuda_builtin_blockIdx_t() {} +}; + +struct __cuda_builtin_blockDim_t { + __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_ntid_x()); + __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_ntid_y()); + __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_ntid_z()); +private: + __cuda_builtin_blockDim_t() {} +}; + +struct __cuda_builtin_gridDim_t { + __CUDA_DEVICE_BUILTIN(x,__builtin_ptx_read_nctaid_x()); + __CUDA_DEVICE_BUILTIN(y,__builtin_ptx_read_nctaid_y()); + __CUDA_DEVICE_BUILTIN(z,__builtin_ptx_read_nctaid_z()); +private: + __cuda_builtin_gridDim_t() {} +}; + +#define __CUDA_BUILTIN_VAR extern const __attribute__((device)) + +__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'. +__CUDA_BUILTIN_VAR int warpSize = 32; + +#undef __CUDA_DEVICE_BUILTIN +#undef __CUDA_BUILTIN_VAR + +#endif /* __CUDA_BUILTIN_VARS_H */ Index: test/CodeGenCUDA/cuda-builtin-vars.cu =================================================================== --- /dev/null +++ test/CodeGenCUDA/cuda-builtin-vars.cu @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s + +#include "cuda/cuda_builtin_vars.h" + +// CHECK: define void @_Z6kernelPi(i32* %out) +__attribute__((global)) +void kernel(int *out) { + int i = 0; + out[i++] = threadIdx.x; // CHECK: call i32 @llvm.ptx.read.tid.x() + out[i++] = threadIdx.y; // CHECK: call i32 @llvm.ptx.read.tid.y() + out[i++] = threadIdx.z; // CHECK: call i32 @llvm.ptx.read.tid.z() + + out[i++] = blockIdx.x; // CHECK: call i32 @llvm.ptx.read.ctaid.x() + out[i++] = blockIdx.y; // CHECK: call i32 @llvm.ptx.read.ctaid.y() + out[i++] = blockIdx.z; // CHECK: call i32 @llvm.ptx.read.ctaid.z() + + out[i++] = blockDim.x; // CHECK: call i32 @llvm.ptx.read.ntid.x() + out[i++] = blockDim.y; // CHECK: call i32 @llvm.ptx.read.ntid.y() + out[i++] = blockDim.z; // CHECK: call i32 @llvm.ptx.read.ntid.z() + + out[i++] = gridDim.x; // CHECK: call i32 @llvm.ptx.read.nctaid.x() + out[i++] = gridDim.y; // CHECK: call i32 @llvm.ptx.read.nctaid.y() + out[i++] = gridDim.z; // CHECK: call i32 @llvm.ptx.read.nctaid.z() + + out[i++] = warpSize; // CHECK: store i32 32, + + // CHECK: ret void +}