Home | History | Annotate | Line # | Download | only in Headers 
      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 
     52 struct __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* 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 
     65 struct __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* 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 
     78 struct __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* 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 
     91 struct __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* 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))
    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 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 */
    122