Home | History | Annotate | Download | only in Headers 
      1 /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
      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 /*
     25  * WARNING: This header is intended to be directly -include'd by
     26  * the compiler and is not supposed to be included by users.
     27  *
     28  * CUDA headers are implemented in a way that currently makes it
     29  * impossible for user code to #include directly when compiling with
     30  * Clang. They present different view of CUDA-supplied functions
     31  * depending on where in NVCC's compilation pipeline the headers are
     32  * included. Neither of these modes provides function definitions with
     33  * correct attributes, so we use preprocessor to force the headers
     34  * into a form that Clang can use.
     35  *
     36  * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's
     37  * this file during every CUDA compilation.
     38  */
     39 
     40 #ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__
     41 #define __CLANG_CUDA_RUNTIME_WRAPPER_H__
     42 
     43 #if defined(__CUDA__) && defined(__clang__)
     44 
     45 // Include some forward declares that must come before cmath.
     46 #include <__clang_cuda_math_forward_declares.h>
     47 
     48 // Include some standard headers to avoid CUDA headers including them
     49 // while some required macros (like __THROW) are in a weird state.
     50 #include <cmath>
     51 #include <cstdlib>
     52 #include <stdlib.h>
     53 
     54 // Preserve common macros that will be changed below by us or by CUDA
     55 // headers.
     56 #pragma push_macro("__THROW")
     57 #pragma push_macro("__CUDA_ARCH__")
     58 
     59 // WARNING: Preprocessor hacks below are based on specific details of
     60 // CUDA-7.x headers and are not expected to work with any other
     61 // version of CUDA headers.
     62 #include "cuda.h"
     63 #if !defined(CUDA_VERSION)
     64 #error "cuda.h did not define CUDA_VERSION"
     65 #elif CUDA_VERSION < 7000 || CUDA_VERSION > 7050
     66 #error "Unsupported CUDA version!"
     67 #endif
     68 
     69 // Make largest subset of device functions available during host
     70 // compilation -- SM_35 for the time being.
     71 #ifndef __CUDA_ARCH__
     72 #define __CUDA_ARCH__ 350
     73 #endif
     74 
     75 #include "cuda_builtin_vars.h"
     76 
     77 // No need for device_launch_parameters.h as cuda_builtin_vars.h above
     78 // has taken care of builtin variables declared in the file.
     79 #define __DEVICE_LAUNCH_PARAMETERS_H__
     80 
     81 // {math,device}_functions.h only have declarations of the
     82 // functions. We don't need them as we're going to pull in their
     83 // definitions from .hpp files.
     84 #define __DEVICE_FUNCTIONS_H__
     85 #define __MATH_FUNCTIONS_H__
     86 #define __COMMON_FUNCTIONS_H__
     87 
     88 #undef __CUDACC__
     89 #define __CUDABE__
     90 // Disables definitions of device-side runtime support stubs in
     91 // cuda_device_runtime_api.h
     92 #include "driver_types.h"
     93 #include "host_config.h"
     94 #include "host_defines.h"
     95 
     96 #undef __CUDABE__
     97 #define __CUDACC__
     98 #include "cuda_runtime.h"
     99 
    100 #undef __CUDACC__
    101 #define __CUDABE__
    102 
    103 // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
    104 // not have at the moment. Emulate them with a builtin memcpy/memset.
    105 #define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n)
    106 #define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n)
    107 
    108 #include "crt/device_runtime.h"
    109 #include "crt/host_runtime.h"
    110 // device_runtime.h defines __cxa_* macros that will conflict with
    111 // cxxabi.h.
    112 // FIXME: redefine these as __device__ functions.
    113 #undef __cxa_vec_ctor
    114 #undef __cxa_vec_cctor
    115 #undef __cxa_vec_dtor
    116 #undef __cxa_vec_new2
    117 #undef __cxa_vec_new3
    118 #undef __cxa_vec_delete2
    119 #undef __cxa_vec_delete
    120 #undef __cxa_vec_delete3
    121 #undef __cxa_pure_virtual
    122 
    123 // We need decls for functions in CUDA's libdevice with __device__
    124 // attribute only. Alas they come either as __host__ __device__ or
    125 // with no attributes at all. To work around that, define __CUDA_RTC__
    126 // which produces HD variant and undef __host__ which gives us desided
    127 // decls with __device__ attribute.
    128 #pragma push_macro("__host__")
    129 #define __host__
    130 #define __CUDACC_RTC__
    131 #include "device_functions_decls.h"
    132 #undef __CUDACC_RTC__
    133 
    134 // Temporarily poison __host__ macro to ensure it's not used by any of
    135 // the headers we're about to include.
    136 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
    137 
    138 // device_functions.hpp and math_functions*.hpp use 'static
    139 // __forceinline__' (with no __device__) for definitions of device
    140 // functions. Temporarily redefine __forceinline__ to include
    141 // __device__.
    142 #pragma push_macro("__forceinline__")
    143 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
    144 #include "device_functions.hpp"
    145 
    146 // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
    147 // get the slow-but-accurate or fast-but-inaccurate versions of functions like
    148 // sin and exp.  This is controlled in clang by -fcuda-approx-transcendentals.
    149 //
    150 // device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
    151 // slow divides), so we need to scope our define carefully here.
    152 #pragma push_macro("__USE_FAST_MATH__")
    153 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
    154 #define __USE_FAST_MATH__
    155 #endif
    156 #include "math_functions.hpp"
    157 #pragma pop_macro("__USE_FAST_MATH__")
    158 
    159 #include "math_functions_dbl_ptx3.hpp"
    160 #pragma pop_macro("__forceinline__")
    161 
    162 // Pull in host-only functions that are only available when neither
    163 // __CUDACC__ nor __CUDABE__ are defined.
    164 #undef __MATH_FUNCTIONS_HPP__
    165 #undef __CUDABE__
    166 #include "math_functions.hpp"
    167 // Alas, additional overloads for these functions are hard to get to.
    168 // Considering that we only need these overloads for a few functions,
    169 // we can provide them here.
    170 static inline float rsqrt(float __a) { return rsqrtf(__a); }
    171 static inline float rcbrt(float __a) { return rcbrtf(__a); }
    172 static inline float sinpi(float __a) { return sinpif(__a); }
    173 static inline float cospi(float __a) { return cospif(__a); }
    174 static inline void sincospi(float __a, float *__b, float *__c) {
    175   return sincospif(__a, __b, __c);
    176 }
    177 static inline float erfcinv(float __a) { return erfcinvf(__a); }
    178 static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
    179 static inline float normcdf(float __a) { return normcdff(__a); }
    180 static inline float erfcx(float __a) { return erfcxf(__a); }
    181 
    182 // For some reason single-argument variant is not always declared by
    183 // CUDA headers. Alas, device_functions.hpp included below needs it.
    184 static inline __device__ void __brkpt(int __c) { __brkpt(); }
    185 
    186 // Now include *.hpp with definitions of various GPU functions.  Alas,
    187 // a lot of thins get declared/defined with __host__ attribute which
    188 // we don't want and we have to define it out. We also have to include
    189 // {device,math}_functions.hpp again in order to extract the other
    190 // branch of #if/else inside.
    191 
    192 #define __host__
    193 #undef __CUDABE__
    194 #define __CUDACC__
    195 #undef __DEVICE_FUNCTIONS_HPP__
    196 #include "device_atomic_functions.hpp"
    197 #include "device_functions.hpp"
    198 #include "sm_20_atomic_functions.hpp"
    199 #include "sm_20_intrinsics.hpp"
    200 #include "sm_32_atomic_functions.hpp"
    201 
    202 // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h.  These define the
    203 // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
    204 // define them using builtins so that the optimizer can reason about and across
    205 // these instructions.  In particular, using intrinsics for ldg gets us the
    206 // [addr+imm] addressing mode, which, although it doesn't actually exist in the
    207 // hardware, seems to generate faster machine code because ptxas can more easily
    208 // reason about our code.
    209 
    210 #undef __MATH_FUNCTIONS_HPP__
    211 
    212 // math_functions.hpp defines ::signbit as a __host__ __device__ function.  This
    213 // conflicts with libstdc++'s constexpr ::signbit, so we have to rename
    214 // math_function.hpp's ::signbit.  It's guarded by #undef signbit, but that's
    215 // conditional on __GNUC__.  :)
    216 #pragma push_macro("signbit")
    217 #pragma push_macro("__GNUC__")
    218 #undef __GNUC__
    219 #define signbit __ignored_cuda_signbit
    220 #include "math_functions.hpp"
    221 #pragma pop_macro("__GNUC__")
    222 #pragma pop_macro("signbit")
    223 
    224 #pragma pop_macro("__host__")
    225 
    226 #include "texture_indirect_functions.h"
    227 
    228 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
    229 #pragma pop_macro("__CUDA_ARCH__")
    230 #pragma pop_macro("__THROW")
    231 
    232 // Set up compiler macros expected to be seen during compilation.
    233 #undef __CUDABE__
    234 #define __CUDACC__
    235 
    236 extern "C" {
    237 // Device-side CUDA system calls.
    238 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
    239 // We need these declarations and wrappers for device-side
    240 // malloc/free/printf calls to work without relying on
    241 // -fcuda-disable-target-call-checks option.
    242 __device__ int vprintf(const char *, const char *);
    243 __device__ void free(void *) __attribute((nothrow));
    244 __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
    245 __device__ void __assertfail(const char *__message, const char *__file,
    246                              unsigned __line, const char *__function,
    247                              size_t __charSize) __attribute__((noreturn));
    248 
    249 // In order for standard assert() macro on linux to work we need to
    250 // provide device-side __assert_fail()
    251 __device__ static inline void __assert_fail(const char *__message,
    252                                             const char *__file, unsigned __line,
    253                                             const char *__function) {
    254   __assertfail(__message, __file, __line, __function, sizeof(char));
    255 }
    256 
    257 // Clang will convert printf into vprintf, but we still need
    258 // device-side declaration for it.
    259 __device__ int printf(const char *, ...);
    260 } // extern "C"
    261 
    262 // We also need device-side std::malloc and std::free.
    263 namespace std {
    264 __device__ static inline void free(void *__ptr) { ::free(__ptr); }
    265 __device__ static inline void *malloc(size_t __size) {
    266   return ::malloc(__size);
    267 }
    268 } // namespace std
    269 
    270 // Out-of-line implementations from cuda_builtin_vars.h.  These need to come
    271 // after we've pulled in the definition of uint3 and dim3.
    272 
    273 __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
    274   uint3 ret;
    275   ret.x = x;
    276   ret.y = y;
    277   ret.z = z;
    278   return ret;
    279 }
    280 
    281 __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
    282   uint3 ret;
    283   ret.x = x;
    284   ret.y = y;
    285   ret.z = z;
    286   return ret;
    287 }
    288 
    289 __device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
    290   return dim3(x, y, z);
    291 }
    292 
    293 __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
    294   return dim3(x, y, z);
    295 }
    296 
    297 #include <__clang_cuda_cmath.h>
    298 #include <__clang_cuda_intrinsics.h>
    299 
    300 // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
    301 // mode, giving them their "proper" types of dim3 and uint3.  This is
    302 // incompatible with the types we give in cuda_builtin_vars.h.  As as hack,
    303 // force-include the header (nvcc doesn't include it by default) but redefine
    304 // dim3 and uint3 to our builtin types.  (Thankfully dim3 and uint3 are only
    305 // used here for the redeclarations of blockDim and threadIdx.)
    306 #pragma push_macro("dim3")
    307 #pragma push_macro("uint3")
    308 #define dim3 __cuda_builtin_blockDim_t
    309 #define uint3 __cuda_builtin_threadIdx_t
    310 #include "curand_mtgp32_kernel.h"
    311 #pragma pop_macro("dim3")
    312 #pragma pop_macro("uint3")
    313 #pragma pop_macro("__USE_FAST_MATH__")
    314 
    315 #endif // __CUDA__
    316 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
    317