clang  14.0.0git
__clang_cuda_builtin_vars.h
Go to the documentation of this file.
1 /*===---- cuda_builtin_vars.h - CUDA built-in variables ---------------------===
2  *
3  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4  * See https://llvm.org/LICENSE.txt for license information.
5  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6  *
7  *===-----------------------------------------------------------------------===
8  */
9 
10 #ifndef __CUDA_BUILTIN_VARS_H
11 #define __CUDA_BUILTIN_VARS_H
12 
13 // Forward declares from vector_types.h.
14 struct uint3;
15 struct dim3;
16 
17 // The file implements built-in CUDA variables using __declspec(property).
18 // https://msdn.microsoft.com/en-us/library/yhfk0thd.aspx
19 // All read accesses of built-in variable fields get converted into calls to a
20 // getter function which in turn calls the appropriate builtin to fetch the
21 // value.
22 //
23 // Example:
24 // int x = threadIdx.x;
25 // IR output:
26 // %0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #3
27 // PTX output:
28 // mov.u32 %r2, %tid.x;
29 
30 #define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC) \
31  __declspec(property(get = __fetch_builtin_##FIELD)) unsigned int FIELD; \
32  static inline __attribute__((always_inline)) \
33  __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) { \
34  return INTRINSIC; \
35  }
36 
37 #if __cplusplus >= 201103L
38 #define __DELETE =delete
39 #else
40 #define __DELETE
41 #endif
42 
43 // Make sure nobody can create instances of the special variable types. nvcc
44 // also disallows taking address of special variables, so we disable address-of
45 // operator as well.
46 #define __CUDA_DISALLOW_BUILTINVAR_ACCESS(TypeName) \
47  __attribute__((device)) TypeName() __DELETE; \
48  __attribute__((device)) TypeName(const TypeName &) __DELETE; \
49  __attribute__((device)) void operator=(const TypeName &) const __DELETE; \
50  __attribute__((device)) TypeName *operator&() const __DELETE
51 
53  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_tid_x());
54  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_tid_y());
55  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_tid_z());
56  // threadIdx should be convertible to uint3 (in fact in nvcc, it *is* a
57  // uint3). This function is defined after we pull in vector_types.h.
58  __attribute__((device)) operator dim3() const;
59  __attribute__((device)) operator uint3() const;
60 
61 private:
62  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_threadIdx_t);
63 };
64 
66  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ctaid_x());
67  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ctaid_y());
68  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ctaid_z());
69  // blockIdx should be convertible to uint3 (in fact in nvcc, it *is* a
70  // uint3). This function is defined after we pull in vector_types.h.
71  __attribute__((device)) operator dim3() const;
72  __attribute__((device)) operator uint3() const;
73 
74 private:
75  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockIdx_t);
76 };
77 
79  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_ntid_x());
80  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_ntid_y());
81  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_ntid_z());
82  // blockDim should be convertible to dim3 (in fact in nvcc, it *is* a
83  // dim3). This function is defined after we pull in vector_types.h.
84  __attribute__((device)) operator dim3() const;
85  __attribute__((device)) operator uint3() const;
86 
87 private:
88  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockDim_t);
89 };
90 
92  __CUDA_DEVICE_BUILTIN(x,__nvvm_read_ptx_sreg_nctaid_x());
93  __CUDA_DEVICE_BUILTIN(y,__nvvm_read_ptx_sreg_nctaid_y());
94  __CUDA_DEVICE_BUILTIN(z,__nvvm_read_ptx_sreg_nctaid_z());
95  // gridDim should be convertible to dim3 (in fact in nvcc, it *is* a
96  // dim3). This function is defined after we pull in vector_types.h.
97  __attribute__((device)) operator dim3() const;
98  __attribute__((device)) operator uint3() const;
99 
100 private:
101  __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_gridDim_t);
102 };
103 
104 #define __CUDA_BUILTIN_VAR \
105  extern const __attribute__((device)) __attribute__((weak))
110 
111 // warpSize should translate to read of %WARP_SZ but there's currently no
112 // builtin to do so. According to PTX v4.2 docs 'to date, all target
113 // architectures have a WARP_SZ value of 32'.
114 __attribute__((device)) const int warpSize = 32;
115 
116 #undef __CUDA_DEVICE_BUILTIN
117 #undef __CUDA_BUILTIN_VAR
118 #undef __CUDA_DISALLOW_BUILTINVAR_ACCESS
119 #undef __DELETE
120 
121 #endif /* __CUDA_BUILTIN_VARS_H */
__cuda_builtin_threadIdx_t
Definition: __clang_cuda_builtin_vars.h:52
__cuda_builtin_threadIdx_t::__attribute__
__attribute__((device)) operator dim3() const
blockDim
__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim
Definition: __clang_cuda_builtin_vars.h:108
__cuda_builtin_gridDim_t::__attribute__
__attribute__((device)) operator dim3() const
__attribute__
__attribute__((device)) const int warpSize
__cuda_builtin_gridDim_t::__CUDA_DEVICE_BUILTIN
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_nctaid_x())
__cuda_builtin_threadIdx_t::__CUDA_DEVICE_BUILTIN
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_tid_x())
__cuda_builtin_blockIdx_t
Definition: __clang_cuda_builtin_vars.h:65
x
IRgen optimization opportunities The common pattern of short x
Definition: README.txt:7
__cuda_builtin_gridDim_t
Definition: __clang_cuda_builtin_vars.h:91
__cuda_builtin_blockIdx_t::__CUDA_DEVICE_BUILTIN
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_ctaid_x())
gridDim
__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim
Definition: __clang_cuda_builtin_vars.h:109
threadIdx
__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx
Definition: __clang_cuda_builtin_vars.h:106
__cuda_builtin_blockDim_t
Definition: __clang_cuda_builtin_vars.h:78
blockIdx
__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx
Definition: __clang_cuda_builtin_vars.h:107
__cuda_builtin_blockIdx_t::__attribute__
__attribute__((device)) operator dim3() const
__CUDA_BUILTIN_VAR
#define __CUDA_BUILTIN_VAR
Definition: __clang_cuda_builtin_vars.h:104
__cuda_builtin_blockDim_t::__CUDA_DEVICE_BUILTIN
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_ntid_x())
__cuda_builtin_blockDim_t::__attribute__
__attribute__((device)) operator dim3() const