122 lines · c
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-exception6 *7 *===-----------------------------------------------------------------------===8 */9 10#ifndef __CUDA_BUILTIN_VARS_H11#define __CUDA_BUILTIN_VARS_H12 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.aspx19// All read accesses of built-in variable fields get converted into calls to a20// getter function which in turn calls the appropriate builtin to fetch the21// value.22//23// Example:24// int x = threadIdx.x;25// IR output:26// %0 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() #327// 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 >= 201103L38#define __DELETE =delete39#else40#define __DELETE41#endif42 43// Make sure nobody can create instances of the special variable types. nvcc44// also disallows taking address of special variables, so we disable address-of45// 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 __DELETE51 52struct __cuda_builtin_threadIdx_t {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* a57 // 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:62 __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_threadIdx_t);63};64 65struct __cuda_builtin_blockIdx_t {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* a70 // 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:75 __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockIdx_t);76};77 78struct __cuda_builtin_blockDim_t {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* a83 // 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:88 __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_blockDim_t);89};90 91struct __cuda_builtin_gridDim_t {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* a96 // 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:101 __CUDA_DISALLOW_BUILTINVAR_ACCESS(__cuda_builtin_gridDim_t);102};103 104#define __CUDA_BUILTIN_VAR \105 extern const __attribute__((device)) __attribute__((weak))106__CUDA_BUILTIN_VAR __cuda_builtin_threadIdx_t threadIdx;107__CUDA_BUILTIN_VAR __cuda_builtin_blockIdx_t blockIdx;108__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim;109__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim;110 111// warpSize should translate to read of %WARP_SZ but there's currently no112// builtin to do so. According to PTX v4.2 docs 'to date, all target113// architectures have a WARP_SZ value of 32'.114__attribute__((device)) const int warpSize = 32;115 116#undef __CUDA_DEVICE_BUILTIN117#undef __CUDA_BUILTIN_VAR118#undef __CUDA_DISALLOW_BUILTINVAR_ACCESS119#undef __DELETE120 121#endif /* __CUDA_BUILTIN_VARS_H */122