__clang_cuda_builtin_vars.h revision 311143 
1/*===---- cuda_builtin_vars.h - CUDA built-in variables ---------------------===
2 *
3 * Permission is hereby granted, free of charge, to any person obtaining a copy
4 * of this software and associated documentation files (the "Software"), to deal
5 * in the Software without restriction, including without limitation the rights
6 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 * copies of the Software, and to permit persons to whom the Software is
8 * furnished to do so, subject to the following conditions:
9 *
10 * The above copyright notice and this permission notice shall be included in
11 * all copies or substantial portions of the Software.
12 *
13 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 * THE SOFTWARE.
20 *
21 *===-----------------------------------------------------------------------===
22 */
23
24#ifndef __CUDA_BUILTIN_VARS_H
25#define __CUDA_BUILTIN_VARS_H
26
27// Forward declares from vector_types.h.
28struct uint3;
29struct dim3;
30
31// The file implements built-in CUDA variables using __declspec(property).
32// https://msdn.microsoft.com/en-us/library/yhfk0thd.aspx
33// All read accesses of built-in variable fields get converted into calls to a
34// getter function which in turn calls the appropriate builtin to fetch the
35// value.
36//
37// Example:
38//    int x = threadIdx.x;
39// IR output:
40//  %0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #3
41// PTX output:
42//  mov.u32     %r2, %tid.x;
43
44#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC)                                \
45  __declspec(property(get = __fetch_builtin_##FIELD)) unsigned int FIELD;      \
46  static inline __attribute__((always_inline))                                 \
47      __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) {     \
48    return INTRINSIC;                                                          \
49  }
50
51#if __cplusplus >= 201103L
52#define __DELETE =delete
53#else
54#define __DELETE
55#endif
56
57// Make sure nobody can create instances of the special varible types.  nvcc
58// also disallows taking address of special variables, so we disable address-of
59// operator as well.
60#define __CUDA_DISALLOW_BUILTINVAR_ACCESS(TypeName)                            \
61  __attribute__((device)) TypeName() __DELETE;                                 \
62  __attribute__((device)) TypeName(const TypeName &) __DELETE;                 \
63  __attribute__((device)) void operator=(const TypeName &) const __DELETE;     \
64  __attribute__((device)) TypeName *operator&() const __DELETE
65
66struct __cuda_builtin_threadIdx_t {
67  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_tid_x());
68  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_tid_y());
69  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_tid_z());
70  // threadIdx should be convertible to uint3 (in fact in nvcc, it *is* a
71  // uint3).  This function is defined after we pull in vector_types.h.
72  __attribute__((device)) operator uint3() const;
73private:
74  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_threadIdx_t);
75};
76
77struct __cuda_builtin_blockIdx_t {
78  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ctaid_x());
79  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ctaid_y());
80  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ctaid_z());
81  // blockIdx should be convertible to uint3 (in fact in nvcc, it *is* a
82  // uint3).  This function is defined after we pull in vector_types.h.
83  __attribute__((device)) operator uint3() const;
84private:
85  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockIdx_t);
86};
87
88struct __cuda_builtin_blockDim_t {
89  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ntid_x());
90  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ntid_y());
91  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ntid_z());
92  // blockDim should be convertible to dim3 (in fact in nvcc, it *is* a
93  // dim3).  This function is defined after we pull in vector_types.h.
94  __attribute__((device)) operator dim3() const;
95private:
96  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockDim_t);
97};
98
99struct __cuda_builtin_gridDim_t {
100  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_nctaid_x());
101  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_nctaid_y());
102  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_nctaid_z());
103  // gridDim should be convertible to dim3 (in fact in nvcc, it *is* a
104  // dim3).  This function is defined after we pull in vector_types.h.
105  __attribute__((device)) operator dim3() const;
106private:
107  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_gridDim_t);
108};
109
110#define __CUDA_BUILTIN_VAR                                                     \
111  extern const __attribute__((device)) __attribute__((weak))
112__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx;
113__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx;
114__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim;
115__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim;
116
117// warpSize should translate to read of %WARP_SZ but there's currently no
118// builtin to do so. According to PTX v4.2 docs 'to date, all target
119// architectures have a WARP_SZ value of 32'.
120__attribute__((device)) const int warpSize = 32;
121
122#undef __CUDA_DEVICE_BUILTIN
123#undef __CUDA_BUILTIN_VAR
124#undef __CUDA_DISALLOW_BUILTINVAR_ACCESS
125
126#endif /* __CUDA_BUILTIN_VARS_H */
127