clang 20.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.
14struct uint3;
15struct 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
61private:
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
74private:
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
87private:
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
100private:
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_VAR __cuda_builtin_blockDim_t blockDim
__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim
__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx
#define __CUDA_BUILTIN_VAR
#define __CUDA_DISALLOW_BUILTINVAR_ACCESS(TypeName)
__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
vector< uint, 3 > uint3
__attribute__((device)) operator dim3() const
__attribute__((device)) operator uint3() const
__CUDA_DEVICE_BUILTIN(z, __nvvm_read_ptx_sreg_ntid_z())
__CUDA_DEVICE_BUILTIN(y, __nvvm_read_ptx_sreg_ntid_y())
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_ntid_x())
__attribute__((device)) operator uint3() const
__attribute__((device)) operator dim3() const
__CUDA_DEVICE_BUILTIN(y, __nvvm_read_ptx_sreg_ctaid_y())
__CUDA_DEVICE_BUILTIN(z, __nvvm_read_ptx_sreg_ctaid_z())
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_ctaid_x())
__CUDA_DEVICE_BUILTIN(y, __nvvm_read_ptx_sreg_nctaid_y())
__attribute__((device)) operator dim3() const
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_nctaid_x())
__CUDA_DEVICE_BUILTIN(z, __nvvm_read_ptx_sreg_nctaid_z())
__attribute__((device)) operator uint3() const
__CUDA_DEVICE_BUILTIN(y, __nvvm_read_ptx_sreg_tid_y())
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_tid_x())
__attribute__((device)) operator uint3() const
__attribute__((device)) operator dim3() const
__CUDA_DEVICE_BUILTIN(z, __nvvm_read_ptx_sreg_tid_z())