1 /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------=== 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 /* 11 * WARNING: This header is intended to be directly -include'd by 12 * the compiler and is not supposed to be included by users. 13 * 14 * CUDA headers are implemented in a way that currently makes it 15 * impossible for user code to #include directly when compiling with 16 * Clang. They present different view of CUDA-supplied functions 17 * depending on where in NVCC's compilation pipeline the headers are 18 * included. Neither of these modes provides function definitions with 19 * correct attributes, so we use preprocessor to force the headers 20 * into a form that Clang can use. 21 * 22 * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's 23 * this file during every CUDA compilation. 24 */ 25 26 #ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__ 27 #define __CLANG_CUDA_RUNTIME_WRAPPER_H__ 28 29 #if defined(__CUDA__) && defined(__clang__) 30 31 // Include some forward declares that must come before cmath. 32 #include <__clang_cuda_math_forward_declares.h> 33 34 // Define __CUDACC__ early as libstdc++ standard headers with GNU extensions 35 // enabled depend on it to avoid using __float128, which is unsupported in 36 // CUDA. 37 #define __CUDACC__ 38 39 // Include some standard headers to avoid CUDA headers including them 40 // while some required macros (like __THROW) are in a weird state. 41 #include <cmath> 42 #include <cstdlib> 43 #include <stdlib.h> 44 #include <string.h> 45 #undef __CUDACC__ 46 47 // Preserve common macros that will be changed below by us or by CUDA 48 // headers. 49 #pragma push_macro("__THROW") 50 #pragma push_macro("__CUDA_ARCH__") 51 52 // WARNING: Preprocessor hacks below are based on specific details of 53 // CUDA-7.x headers and are not expected to work with any other 54 // version of CUDA headers. 55 #include "cuda.h" 56 #if !defined(CUDA_VERSION) 57 #error "cuda.h did not define CUDA_VERSION" 58 #elif CUDA_VERSION < 7000 59 #error "Unsupported CUDA version!" 60 #endif 61 62 #pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") 63 #if CUDA_VERSION >= 10000 64 #define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__ 65 #endif 66 67 // Make largest subset of device functions available during host 68 // compilation. 69 #ifndef __CUDA_ARCH__ 70 #define __CUDA_ARCH__ 9999 71 #endif 72 73 #include "__clang_cuda_builtin_vars.h" 74 75 // No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above 76 // has taken care of builtin variables declared in the file. 77 #define __DEVICE_LAUNCH_PARAMETERS_H__ 78 79 // {math,device}_functions.h only have declarations of the 80 // functions. We don't need them as we're going to pull in their 81 // definitions from .hpp files. 82 #define __DEVICE_FUNCTIONS_H__ 83 #define __MATH_FUNCTIONS_H__ 84 #define __COMMON_FUNCTIONS_H__ 85 // device_functions_decls is replaced by __clang_cuda_device_functions.h 86 // included below. 87 #define __DEVICE_FUNCTIONS_DECLS_H__ 88 89 #undef __CUDACC__ 90 #if CUDA_VERSION < 9000 91 #define __CUDABE__ 92 #else 93 #define __CUDACC__ 94 #define __CUDA_LIBDEVICE__ 95 #endif 96 // Disables definitions of device-side runtime support stubs in 97 // cuda_device_runtime_api.h 98 #include "host_defines.h" 99 #undef __CUDACC__ 100 #include "driver_types.h" 101 #include "host_config.h" 102 103 // Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in 104 // cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the 105 // functional equivalent of what we need. 106 #pragma push_macro("nv_weak") 107 #define nv_weak weak 108 #undef __CUDABE__ 109 #undef __CUDA_LIBDEVICE__ 110 #define __CUDACC__ 111 #include "cuda_runtime.h" 112 113 #pragma pop_macro("nv_weak") 114 #undef __CUDACC__ 115 #define __CUDABE__ 116 117 // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does 118 // not have at the moment. Emulate them with a builtin memcpy/memset. 119 #define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n) 120 #define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n) 121 122 #if CUDA_VERSION < 9000 123 #include "crt/device_runtime.h" 124 #endif 125 #include "crt/host_runtime.h" 126 // device_runtime.h defines __cxa_* macros that will conflict with 127 // cxxabi.h. 128 // FIXME: redefine these as __device__ functions. 129 #undef __cxa_vec_ctor 130 #undef __cxa_vec_cctor 131 #undef __cxa_vec_dtor 132 #undef __cxa_vec_new 133 #undef __cxa_vec_new2 134 #undef __cxa_vec_new3 135 #undef __cxa_vec_delete2 136 #undef __cxa_vec_delete 137 #undef __cxa_vec_delete3 138 #undef __cxa_pure_virtual 139 140 // math_functions.hpp expects this host function be defined on MacOS, but it 141 // ends up not being there because of the games we play here. Just define it 142 // ourselves; it's simple enough. 143 #ifdef __APPLE__ 144 inline __host__ double __signbitd(double x) { 145 return std::signbit(x); 146 } 147 #endif 148 149 // CUDA 9.1 no longer provides declarations for libdevice functions, so we need 150 // to provide our own. 151 #include <__clang_cuda_libdevice_declares.h> 152 153 // Wrappers for many device-side standard library functions, incl. math 154 // functions, became compiler builtins in CUDA-9 and have been removed from the 155 // CUDA headers. Clang now provides its own implementation of the wrappers. 156 #if CUDA_VERSION >= 9000 157 #include <__clang_cuda_device_functions.h> 158 #include <__clang_cuda_math.h> 159 #endif 160 161 // __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's 162 // counterpart does not do it, so we need to make it empty here to keep 163 // following CUDA includes happy. 164 #undef __THROW 165 #define __THROW 166 167 // CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values. 168 // Previous versions used to check whether they are defined or not. 169 // CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it 170 // here to detect the switch. 171 172 #if defined(CU_DEVICE_INVALID) 173 #if !defined(__USE_FAST_MATH__) 174 #define __USE_FAST_MATH__ 0 175 #endif 176 177 #if !defined(__CUDA_PREC_DIV) 178 #define __CUDA_PREC_DIV 0 179 #endif 180 #endif 181 182 // Temporarily poison __host__ macro to ensure it's not used by any of 183 // the headers we're about to include. 184 #pragma push_macro("__host__") 185 #define __host__ UNEXPECTED_HOST_ATTRIBUTE 186 187 // device_functions.hpp and math_functions*.hpp use 'static 188 // __forceinline__' (with no __device__) for definitions of device 189 // functions. Temporarily redefine __forceinline__ to include 190 // __device__. 191 #pragma push_macro("__forceinline__") 192 #define __forceinline__ __device__ __inline__ __attribute__((always_inline)) 193 #if CUDA_VERSION < 9000 194 #include "device_functions.hpp" 195 #endif 196 197 // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we 198 // get the slow-but-accurate or fast-but-inaccurate versions of functions like 199 // sin and exp. This is controlled in clang by -fgpu-approx-transcendentals. 200 // 201 // device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs. 202 // slow divides), so we need to scope our define carefully here. 203 #pragma push_macro("__USE_FAST_MATH__") 204 #if defined(__CLANG_GPU_APPROX_TRANSCENDENTALS__) 205 #define __USE_FAST_MATH__ 1 206 #endif 207 208 #if CUDA_VERSION >= 9000 209 #include "crt/math_functions.hpp" 210 #else 211 #include "math_functions.hpp" 212 #endif 213 214 #pragma pop_macro("__USE_FAST_MATH__") 215 216 #if CUDA_VERSION < 9000 217 #include "math_functions_dbl_ptx3.hpp" 218 #endif 219 #pragma pop_macro("__forceinline__") 220 221 // Pull in host-only functions that are only available when neither 222 // __CUDACC__ nor __CUDABE__ are defined. 223 #undef __MATH_FUNCTIONS_HPP__ 224 #undef __CUDABE__ 225 #if CUDA_VERSION < 9000 226 #include "math_functions.hpp" 227 #endif 228 // Alas, additional overloads for these functions are hard to get to. 229 // Considering that we only need these overloads for a few functions, 230 // we can provide them here. 231 static inline float rsqrt(float __a) { return rsqrtf(__a); } 232 static inline float rcbrt(float __a) { return rcbrtf(__a); } 233 static inline float sinpi(float __a) { return sinpif(__a); } 234 static inline float cospi(float __a) { return cospif(__a); } 235 static inline void sincospi(float __a, float *__b, float *__c) { 236 return sincospif(__a, __b, __c); 237 } 238 static inline float erfcinv(float __a) { return erfcinvf(__a); } 239 static inline float normcdfinv(float __a) { return normcdfinvf(__a); } 240 static inline float normcdf(float __a) { return normcdff(__a); } 241 static inline float erfcx(float __a) { return erfcxf(__a); } 242 243 #if CUDA_VERSION < 9000 244 // For some reason single-argument variant is not always declared by 245 // CUDA headers. Alas, device_functions.hpp included below needs it. 246 static inline __device__ void __brkpt(int __c) { __brkpt(); } 247 #endif 248 249 // Now include *.hpp with definitions of various GPU functions. Alas, 250 // a lot of thins get declared/defined with __host__ attribute which 251 // we don't want and we have to define it out. We also have to include 252 // {device,math}_functions.hpp again in order to extract the other 253 // branch of #if/else inside. 254 #define __host__ 255 #undef __CUDABE__ 256 #define __CUDACC__ 257 #if CUDA_VERSION >= 9000 258 // Some atomic functions became compiler builtins in CUDA-9 , so we need their 259 // declarations. 260 #include "device_atomic_functions.h" 261 #endif 262 #undef __DEVICE_FUNCTIONS_HPP__ 263 #include "device_atomic_functions.hpp" 264 #if CUDA_VERSION >= 9000 265 #include "crt/device_functions.hpp" 266 #include "crt/device_double_functions.hpp" 267 #else 268 #include "device_functions.hpp" 269 #define __CUDABE__ 270 #include "device_double_functions.h" 271 #undef __CUDABE__ 272 #endif 273 #include "sm_20_atomic_functions.hpp" 274 // Predicate functions used in `__builtin_assume` need to have no side effect. 275 // However, sm_20_intrinsics.hpp doesn't define them with neither pure nor 276 // const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define 277 // them as pure ones. 278 #pragma push_macro("__isGlobal") 279 #pragma push_macro("__isShared") 280 #pragma push_macro("__isConstant") 281 #pragma push_macro("__isLocal") 282 #define __isGlobal __ignored_cuda___isGlobal 283 #define __isShared __ignored_cuda___isShared 284 #define __isConstant __ignored_cuda___isConstant 285 #define __isLocal __ignored_cuda___isLocal 286 #include "sm_20_intrinsics.hpp" 287 #pragma pop_macro("__isGlobal") 288 #pragma pop_macro("__isShared") 289 #pragma pop_macro("__isConstant") 290 #pragma pop_macro("__isLocal") 291 #pragma push_macro("__DEVICE__") 292 #define __DEVICE__ static __device__ __forceinline__ __attribute__((const)) 293 __DEVICE__ unsigned int __isGlobal(const void *p) { 294 return __nvvm_isspacep_global(p); 295 } 296 __DEVICE__ unsigned int __isShared(const void *p) { 297 return __nvvm_isspacep_shared(p); 298 } 299 __DEVICE__ unsigned int __isConstant(const void *p) { 300 return __nvvm_isspacep_const(p); 301 } 302 __DEVICE__ unsigned int __isLocal(const void *p) { 303 return __nvvm_isspacep_local(p); 304 } 305 #pragma pop_macro("__DEVICE__") 306 #include "sm_32_atomic_functions.hpp" 307 308 // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the 309 // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to 310 // define them using builtins so that the optimizer can reason about and across 311 // these instructions. In particular, using intrinsics for ldg gets us the 312 // [addr+imm] addressing mode, which, although it doesn't actually exist in the 313 // hardware, seems to generate faster machine code because ptxas can more easily 314 // reason about our code. 315 316 #if CUDA_VERSION >= 8000 317 #pragma push_macro("__CUDA_ARCH__") 318 #undef __CUDA_ARCH__ 319 #include "sm_60_atomic_functions.hpp" 320 #include "sm_61_intrinsics.hpp" 321 #pragma pop_macro("__CUDA_ARCH__") 322 #endif 323 324 #undef __MATH_FUNCTIONS_HPP__ 325 326 // math_functions.hpp defines ::signbit as a __host__ __device__ function. This 327 // conflicts with libstdc++'s constexpr ::signbit, so we have to rename 328 // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's 329 // conditional on __GNUC__. :) 330 #pragma push_macro("signbit") 331 #pragma push_macro("__GNUC__") 332 #undef __GNUC__ 333 #define signbit __ignored_cuda_signbit 334 335 // CUDA-9 omits device-side definitions of some math functions if it sees 336 // include guard from math.h wrapper from libstdc++. We have to undo the header 337 // guard temporarily to get the definitions we need. 338 #pragma push_macro("_GLIBCXX_MATH_H") 339 #pragma push_macro("_LIBCPP_VERSION") 340 #if CUDA_VERSION >= 9000 341 #undef _GLIBCXX_MATH_H 342 // We also need to undo another guard that checks for libc++ 3.8+ 343 #ifdef _LIBCPP_VERSION 344 #define _LIBCPP_VERSION 3700 345 #endif 346 #endif 347 348 #if CUDA_VERSION >= 9000 349 #include "crt/math_functions.hpp" 350 #else 351 #include "math_functions.hpp" 352 #endif 353 #pragma pop_macro("_GLIBCXX_MATH_H") 354 #pragma pop_macro("_LIBCPP_VERSION") 355 #pragma pop_macro("__GNUC__") 356 #pragma pop_macro("signbit") 357 358 #pragma pop_macro("__host__") 359 360 // __clang_cuda_texture_intrinsics.h must be included first in order to provide 361 // implementation for __nv_tex_surf_handler that CUDA's headers depend on. 362 // The implementation requires c++11 and only works with CUDA-9 or newer. 363 #if __cplusplus >= 201103L && CUDA_VERSION >= 9000 364 // clang-format off 365 #include <__clang_cuda_texture_intrinsics.h> 366 // clang-format on 367 #else 368 #if CUDA_VERSION >= 9000 369 // Provide a hint that texture support needs C++11. 370 template <typename T> struct __nv_tex_needs_cxx11 { 371 const static bool value = false; 372 }; 373 template <class T> 374 __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr, 375 cudaTextureObject_t obj, 376 float x) { 377 _Static_assert(__nv_tex_needs_cxx11<T>::value, 378 "Texture support requires C++11"); 379 } 380 #else 381 // Textures in CUDA-8 and older are not supported by clang.There's no 382 // convenient way to intercept texture use in these versions, so we can't 383 // produce a meaningful error. The source code that attempts to use textures 384 // will continue to fail as it does now. 385 #endif // CUDA_VERSION 386 #endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000 387 #include "texture_fetch_functions.h" 388 #include "texture_indirect_functions.h" 389 390 // Restore state of __CUDA_ARCH__ and __THROW we had on entry. 391 #pragma pop_macro("__CUDA_ARCH__") 392 #pragma pop_macro("__THROW") 393 394 // Set up compiler macros expected to be seen during compilation. 395 #undef __CUDABE__ 396 #define __CUDACC__ 397 398 extern "C" { 399 // Device-side CUDA system calls. 400 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls 401 // We need these declarations and wrappers for device-side 402 // malloc/free/printf calls to work without relying on 403 // -fcuda-disable-target-call-checks option. 404 __device__ int vprintf(const char *, const char *); 405 __device__ void free(void *) __attribute((nothrow)); 406 __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); 407 408 // __assertfail() used to have a `noreturn` attribute. Unfortunately that 409 // contributed to triggering the longstanding bug in ptxas when assert was used 410 // in sufficiently convoluted code. See 411 // https://bugs.llvm.org/show_bug.cgi?id=27738 for the details. 412 __device__ void __assertfail(const char *__message, const char *__file, 413 unsigned __line, const char *__function, 414 size_t __charSize); 415 416 // In order for standard assert() macro on linux to work we need to 417 // provide device-side __assert_fail() 418 __device__ static inline void __assert_fail(const char *__message, 419 const char *__file, unsigned __line, 420 const char *__function) { 421 __assertfail(__message, __file, __line, __function, sizeof(char)); 422 } 423 424 // Clang will convert printf into vprintf, but we still need 425 // device-side declaration for it. 426 __device__ int printf(const char *, ...); 427 } // extern "C" 428 429 // We also need device-side std::malloc and std::free. 430 namespace std { 431 __device__ static inline void free(void *__ptr) { ::free(__ptr); } 432 __device__ static inline void *malloc(size_t __size) { 433 return ::malloc(__size); 434 } 435 } // namespace std 436 437 // Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to 438 // come after we've pulled in the definition of uint3 and dim3. 439 440 __device__ inline __cuda_builtin_threadIdx_t::operator dim3() const { 441 return dim3(x, y, z); 442 } 443 444 __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { 445 return {x, y, z}; 446 } 447 448 __device__ inline __cuda_builtin_blockIdx_t::operator dim3() const { 449 return dim3(x, y, z); 450 } 451 452 __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { 453 return {x, y, z}; 454 } 455 456 __device__ inline __cuda_builtin_blockDim_t::operator dim3() const { 457 return dim3(x, y, z); 458 } 459 460 __device__ inline __cuda_builtin_blockDim_t::operator uint3() const { 461 return {x, y, z}; 462 } 463 464 __device__ inline __cuda_builtin_gridDim_t::operator dim3() const { 465 return dim3(x, y, z); 466 } 467 468 __device__ inline __cuda_builtin_gridDim_t::operator uint3() const { 469 return {x, y, z}; 470 } 471 472 #include <__clang_cuda_cmath.h> 473 #include <__clang_cuda_intrinsics.h> 474 #include <__clang_cuda_complex_builtins.h> 475 476 // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host 477 // mode, giving them their "proper" types of dim3 and uint3. This is 478 // incompatible with the types we give in __clang_cuda_builtin_vars.h. As as 479 // hack, force-include the header (nvcc doesn't include it by default) but 480 // redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are 481 // only used here for the redeclarations of blockDim and threadIdx.) 482 #pragma push_macro("dim3") 483 #pragma push_macro("uint3") 484 #define dim3 __cuda_builtin_blockDim_t 485 #define uint3 __cuda_builtin_threadIdx_t 486 #include "curand_mtgp32_kernel.h" 487 #pragma pop_macro("dim3") 488 #pragma pop_macro("uint3") 489 #pragma pop_macro("__USE_FAST_MATH__") 490 #pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") 491 492 // CUDA runtime uses this undocumented function to access kernel launch 493 // configuration. The declaration is in crt/device_functions.h but that file 494 // includes a lot of other stuff we don't want. Instead, we'll provide our own 495 // declaration for it here. 496 #if CUDA_VERSION >= 9020 497 extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, 498 size_t sharedMem = 0, 499 void *stream = 0); 500 #endif 501 502 #endif // __CUDA__ 503 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ 504