contrib/llvm/tools/clang/lib/Headers/__clang_cuda_runtime_wrapper.h

   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 > 8000
  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 "__clang_cuda_builtin_vars.h"
  76
  77 // No need for device_launch_parameters.h as __clang_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_new
 117 #undef __cxa_vec_new2
 118 #undef __cxa_vec_new3
 119 #undef __cxa_vec_delete2
 120 #undef __cxa_vec_delete
 121 #undef __cxa_vec_delete3
 122 #undef __cxa_pure_virtual
 123
 124 // math_functions.hpp expects this host function be defined on MacOS, but it
 125 // ends up not being there because of the games we play here.  Just define it
 126 // ourselves; it's simple enough.
 127 #ifdef __APPLE__
 128 inline __host__ double __signbitd(double x) {
 129   return std::signbit(x);
 130 }
 131 #endif
 132
 133 // We need decls for functions in CUDA's libdevice with __device__
 134 // attribute only. Alas they come either as __host__ __device__ or
 135 // with no attributes at all. To work around that, define __CUDA_RTC__
 136 // which produces HD variant and undef __host__ which gives us desided
 137 // decls with __device__ attribute.
 138 #pragma push_macro("__host__")
 139 #define __host__
 140 #define __CUDACC_RTC__
 141 #include "device_functions_decls.h"
 142 #undef __CUDACC_RTC__
 143
 144 // Temporarily poison __host__ macro to ensure it's not used by any of
 145 // the headers we're about to include.
 146 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
 147
 148 // CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values.
 149 // Previous versions used to check whether they are defined or not.
 150 // CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it
 151 // here to detect the switch.
 152
 153 #if defined(CU_DEVICE_INVALID)
 154 #if !defined(__USE_FAST_MATH__)
 155 #define __USE_FAST_MATH__ 0
 156 #endif
 157
 158 #if !defined(__CUDA_PREC_DIV)
 159 #define __CUDA_PREC_DIV 0
 160 #endif
 161 #endif
 162
 163 // device_functions.hpp and math_functions*.hpp use 'static
 164 // __forceinline__' (with no __device__) for definitions of device
 165 // functions. Temporarily redefine __forceinline__ to include
 166 // __device__.
 167 #pragma push_macro("__forceinline__")
 168 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
 169 #include "device_functions.hpp"
 170
 171 // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
 172 // get the slow-but-accurate or fast-but-inaccurate versions of functions like
 173 // sin and exp.  This is controlled in clang by -fcuda-approx-transcendentals.
 174 //
 175 // device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
 176 // slow divides), so we need to scope our define carefully here.
 177 #pragma push_macro("__USE_FAST_MATH__")
 178 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
 179 #define __USE_FAST_MATH__ 1
 180 #endif
 181 #include "math_functions.hpp"
 182 #pragma pop_macro("__USE_FAST_MATH__")
 183
 184 #include "math_functions_dbl_ptx3.hpp"
 185 #pragma pop_macro("__forceinline__")
 186
 187 // Pull in host-only functions that are only available when neither
 188 // __CUDACC__ nor __CUDABE__ are defined.
 189 #undef __MATH_FUNCTIONS_HPP__
 190 #undef __CUDABE__
 191 #include "math_functions.hpp"
 192 // Alas, additional overloads for these functions are hard to get to.
 193 // Considering that we only need these overloads for a few functions,
 194 // we can provide them here.
 195 static inline float rsqrt(float __a) { return rsqrtf(__a); }
 196 static inline float rcbrt(float __a) { return rcbrtf(__a); }
 197 static inline float sinpi(float __a) { return sinpif(__a); }
 198 static inline float cospi(float __a) { return cospif(__a); }
 199 static inline void sincospi(float __a, float *__b, float *__c) {
 200   return sincospif(__a, __b, __c);
 201 }
 202 static inline float erfcinv(float __a) { return erfcinvf(__a); }
 203 static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
 204 static inline float normcdf(float __a) { return normcdff(__a); }
 205 static inline float erfcx(float __a) { return erfcxf(__a); }
 206
 207 // For some reason single-argument variant is not always declared by
 208 // CUDA headers. Alas, device_functions.hpp included below needs it.
 209 static inline __device__ void __brkpt(int __c) { __brkpt(); }
 210
 211 // Now include *.hpp with definitions of various GPU functions.  Alas,
 212 // a lot of thins get declared/defined with __host__ attribute which
 213 // we don't want and we have to define it out. We also have to include
 214 // {device,math}_functions.hpp again in order to extract the other
 215 // branch of #if/else inside.
 216
 217 #define __host__
 218 #undef __CUDABE__
 219 #define __CUDACC__
 220 #undef __DEVICE_FUNCTIONS_HPP__
 221 #include "device_atomic_functions.hpp"
 222 #include "device_functions.hpp"
 223 #include "sm_20_atomic_functions.hpp"
 224 #include "sm_20_intrinsics.hpp"
 225 #include "sm_32_atomic_functions.hpp"
 226
 227 // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h.  These define the
 228 // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
 229 // define them using builtins so that the optimizer can reason about and across
 230 // these instructions.  In particular, using intrinsics for ldg gets us the
 231 // [addr+imm] addressing mode, which, although it doesn't actually exist in the
 232 // hardware, seems to generate faster machine code because ptxas can more easily
 233 // reason about our code.
 234
 235 #if CUDA_VERSION >= 8000
 236 #include "sm_60_atomic_functions.hpp"
 237 #include "sm_61_intrinsics.hpp"
 238 #endif
 239
 240 #undef __MATH_FUNCTIONS_HPP__
 241
 242 // math_functions.hpp defines ::signbit as a __host__ __device__ function.  This
 243 // conflicts with libstdc++'s constexpr ::signbit, so we have to rename
 244 // math_function.hpp's ::signbit.  It's guarded by #undef signbit, but that's
 245 // conditional on __GNUC__.  :)
 246 #pragma push_macro("signbit")
 247 #pragma push_macro("__GNUC__")
 248 #undef __GNUC__
 249 #define signbit __ignored_cuda_signbit
 250 #include "math_functions.hpp"
 251 #pragma pop_macro("__GNUC__")
 252 #pragma pop_macro("signbit")
 253
 254 #pragma pop_macro("__host__")
 255
 256 #include "texture_indirect_functions.h"
 257
 258 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
 259 #pragma pop_macro("__CUDA_ARCH__")
 260 #pragma pop_macro("__THROW")
 261
 262 // Set up compiler macros expected to be seen during compilation.
 263 #undef __CUDABE__
 264 #define __CUDACC__
 265
 266 extern "C" {
 267 // Device-side CUDA system calls.
 268 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
 269 // We need these declarations and wrappers for device-side
 270 // malloc/free/printf calls to work without relying on
 271 // -fcuda-disable-target-call-checks option.
 272 __device__ int vprintf(const char *, const char *);
 273 __device__ void free(void *) __attribute((nothrow));
 274 __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
 275 __device__ void __assertfail(const char *__message, const char *__file,
 276                              unsigned __line, const char *__function,
 277                              size_t __charSize) __attribute__((noreturn));
 278
 279 // In order for standard assert() macro on linux to work we need to
 280 // provide device-side __assert_fail()
 281 __device__ static inline void __assert_fail(const char *__message,
 282                                             const char *__file, unsigned __line,
 283                                             const char *__function) {
 284   __assertfail(__message, __file, __line, __function, sizeof(char));
 285 }
 286
 287 // Clang will convert printf into vprintf, but we still need
 288 // device-side declaration for it.
 289 __device__ int printf(const char *, ...);
 290 } // extern "C"
 291
 292 // We also need device-side std::malloc and std::free.
 293 namespace std {
 294 __device__ static inline void free(void *__ptr) { ::free(__ptr); }
 295 __device__ static inline void *malloc(size_t __size) {
 296   return ::malloc(__size);
 297 }
 298 } // namespace std
 299
 300 // Out-of-line implementations from __clang_cuda_builtin_vars.h.  These need to
 301 // come after we've pulled in the definition of uint3 and dim3.
 302
 303 __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
 304   uint3 ret;
 305   ret.x = x;
 306   ret.y = y;
 307   ret.z = z;
 308   return ret;
 309 }
 310
 311 __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
 312   uint3 ret;
 313   ret.x = x;
 314   ret.y = y;
 315   ret.z = z;
 316   return ret;
 317 }
 318
 319 __device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
 320   return dim3(x, y, z);
 321 }
 322
 323 __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
 324   return dim3(x, y, z);
 325 }
 326
 327 #include <__clang_cuda_cmath.h>
 328 #include <__clang_cuda_intrinsics.h>
 329 #include <__clang_cuda_complex_builtins.h>
 330
 331 // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
 332 // mode, giving them their "proper" types of dim3 and uint3.  This is
 333 // incompatible with the types we give in __clang_cuda_builtin_vars.h.  As as
 334 // hack, force-include the header (nvcc doesn't include it by default) but
 335 // redefine dim3 and uint3 to our builtin types.  (Thankfully dim3 and uint3 are
 336 // only used here for the redeclarations of blockDim and threadIdx.)
 337 #pragma push_macro("dim3")
 338 #pragma push_macro("uint3")
 339 #define dim3 __cuda_builtin_blockDim_t
 340 #define uint3 __cuda_builtin_threadIdx_t
 341 #include "curand_mtgp32_kernel.h"
 342 #pragma pop_macro("dim3")
 343 #pragma pop_macro("uint3")
 344 #pragma pop_macro("__USE_FAST_MATH__")
 345
 346 #endif // __CUDA__
 347 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__