__clang_cuda_runtime_wrapper.h revision 341825
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 > 9020 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// device_functions_decls is replaced by __clang_cuda_device_functions.h 88// included below. 89#define __DEVICE_FUNCTIONS_DECLS_H__ 90 91#undef __CUDACC__ 92#if CUDA_VERSION < 9000 93#define __CUDABE__ 94#else 95#define __CUDA_LIBDEVICE__ 96#endif 97// Disables definitions of device-side runtime support stubs in 98// cuda_device_runtime_api.h 99#include "driver_types.h" 100#include "host_config.h" 101#include "host_defines.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__ 144inline __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 became compiler 154// builtins in CUDA-9 and have been removed from the CUDA headers. Clang now 155// provides its own implementation of the wrappers. 156#if CUDA_VERSION >= 9000 157#include <__clang_cuda_device_functions.h> 158#endif 159 160// __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's 161// counterpart does not do it, so we need to make it empty here to keep 162// following CUDA includes happy. 163#undef __THROW 164#define __THROW 165 166// CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values. 167// Previous versions used to check whether they are defined or not. 168// CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it 169// here to detect the switch. 170 171#if defined(CU_DEVICE_INVALID) 172#if !defined(__USE_FAST_MATH__) 173#define __USE_FAST_MATH__ 0 174#endif 175 176#if !defined(__CUDA_PREC_DIV) 177#define __CUDA_PREC_DIV 0 178#endif 179#endif 180 181// Temporarily poison __host__ macro to ensure it's not used by any of 182// the headers we're about to include. 183#pragma push_macro("__host__") 184#define __host__ UNEXPECTED_HOST_ATTRIBUTE 185 186// device_functions.hpp and math_functions*.hpp use 'static 187// __forceinline__' (with no __device__) for definitions of device 188// functions. Temporarily redefine __forceinline__ to include 189// __device__. 190#pragma push_macro("__forceinline__") 191#define __forceinline__ __device__ __inline__ __attribute__((always_inline)) 192#if CUDA_VERSION < 9000 193#include "device_functions.hpp" 194#endif 195 196// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we 197// get the slow-but-accurate or fast-but-inaccurate versions of functions like 198// sin and exp. This is controlled in clang by -fcuda-approx-transcendentals. 199// 200// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs. 201// slow divides), so we need to scope our define carefully here. 202#pragma push_macro("__USE_FAST_MATH__") 203#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__) 204#define __USE_FAST_MATH__ 1 205#endif 206 207#if CUDA_VERSION >= 9000 208// CUDA-9.2 needs host-side memcpy for some host functions in 209// device_functions.hpp 210#if CUDA_VERSION >= 9020 211#include <string.h> 212#endif 213#include "crt/math_functions.hpp" 214#else 215#include "math_functions.hpp" 216#endif 217 218#pragma pop_macro("__USE_FAST_MATH__") 219 220#if CUDA_VERSION < 9000 221#include "math_functions_dbl_ptx3.hpp" 222#endif 223#pragma pop_macro("__forceinline__") 224 225// Pull in host-only functions that are only available when neither 226// __CUDACC__ nor __CUDABE__ are defined. 227#undef __MATH_FUNCTIONS_HPP__ 228#undef __CUDABE__ 229#if CUDA_VERSION < 9000 230#include "math_functions.hpp" 231#endif 232// Alas, additional overloads for these functions are hard to get to. 233// Considering that we only need these overloads for a few functions, 234// we can provide them here. 235static inline float rsqrt(float __a) { return rsqrtf(__a); } 236static inline float rcbrt(float __a) { return rcbrtf(__a); } 237static inline float sinpi(float __a) { return sinpif(__a); } 238static inline float cospi(float __a) { return cospif(__a); } 239static inline void sincospi(float __a, float *__b, float *__c) { 240 return sincospif(__a, __b, __c); 241} 242static inline float erfcinv(float __a) { return erfcinvf(__a); } 243static inline float normcdfinv(float __a) { return normcdfinvf(__a); } 244static inline float normcdf(float __a) { return normcdff(__a); } 245static inline float erfcx(float __a) { return erfcxf(__a); } 246 247#if CUDA_VERSION < 9000 248// For some reason single-argument variant is not always declared by 249// CUDA headers. Alas, device_functions.hpp included below needs it. 250static inline __device__ void __brkpt(int __c) { __brkpt(); } 251#endif 252 253// Now include *.hpp with definitions of various GPU functions. Alas, 254// a lot of thins get declared/defined with __host__ attribute which 255// we don't want and we have to define it out. We also have to include 256// {device,math}_functions.hpp again in order to extract the other 257// branch of #if/else inside. 258#define __host__ 259#undef __CUDABE__ 260#define __CUDACC__ 261#if CUDA_VERSION >= 9000 262// Some atomic functions became compiler builtins in CUDA-9 , so we need their 263// declarations. 264#include "device_atomic_functions.h" 265#endif 266#undef __DEVICE_FUNCTIONS_HPP__ 267#include "device_atomic_functions.hpp" 268#if CUDA_VERSION >= 9000 269#include "crt/device_functions.hpp" 270#include "crt/device_double_functions.hpp" 271#else 272#include "device_functions.hpp" 273#define __CUDABE__ 274#include "device_double_functions.h" 275#undef __CUDABE__ 276#endif 277#include "sm_20_atomic_functions.hpp" 278#include "sm_20_intrinsics.hpp" 279#include "sm_32_atomic_functions.hpp" 280 281// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the 282// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to 283// define them using builtins so that the optimizer can reason about and across 284// these instructions. In particular, using intrinsics for ldg gets us the 285// [addr+imm] addressing mode, which, although it doesn't actually exist in the 286// hardware, seems to generate faster machine code because ptxas can more easily 287// reason about our code. 288 289#if CUDA_VERSION >= 8000 290#pragma push_macro("__CUDA_ARCH__") 291#undef __CUDA_ARCH__ 292#include "sm_60_atomic_functions.hpp" 293#include "sm_61_intrinsics.hpp" 294#pragma pop_macro("__CUDA_ARCH__") 295#endif 296 297#undef __MATH_FUNCTIONS_HPP__ 298 299// math_functions.hpp defines ::signbit as a __host__ __device__ function. This 300// conflicts with libstdc++'s constexpr ::signbit, so we have to rename 301// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's 302// conditional on __GNUC__. :) 303#pragma push_macro("signbit") 304#pragma push_macro("__GNUC__") 305#undef __GNUC__ 306#define signbit __ignored_cuda_signbit 307 308// CUDA-9 omits device-side definitions of some math functions if it sees 309// include guard from math.h wrapper from libstdc++. We have to undo the header 310// guard temporarily to get the definitions we need. 311#pragma push_macro("_GLIBCXX_MATH_H") 312#pragma push_macro("_LIBCPP_VERSION") 313#if CUDA_VERSION >= 9000 314#undef _GLIBCXX_MATH_H 315// We also need to undo another guard that checks for libc++ 3.8+ 316#ifdef _LIBCPP_VERSION 317#define _LIBCPP_VERSION 3700 318#endif 319#endif 320 321#if CUDA_VERSION >= 9000 322#include "crt/math_functions.hpp" 323#else 324#include "math_functions.hpp" 325#endif 326#pragma pop_macro("_GLIBCXX_MATH_H") 327#pragma pop_macro("_LIBCPP_VERSION") 328#pragma pop_macro("__GNUC__") 329#pragma pop_macro("signbit") 330 331#pragma pop_macro("__host__") 332 333#include "texture_indirect_functions.h" 334 335// Restore state of __CUDA_ARCH__ and __THROW we had on entry. 336#pragma pop_macro("__CUDA_ARCH__") 337#pragma pop_macro("__THROW") 338 339// Set up compiler macros expected to be seen during compilation. 340#undef __CUDABE__ 341#define __CUDACC__ 342 343extern "C" { 344// Device-side CUDA system calls. 345// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls 346// We need these declarations and wrappers for device-side 347// malloc/free/printf calls to work without relying on 348// -fcuda-disable-target-call-checks option. 349__device__ int vprintf(const char *, const char *); 350__device__ void free(void *) __attribute((nothrow)); 351__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); 352__device__ void __assertfail(const char *__message, const char *__file, 353 unsigned __line, const char *__function, 354 size_t __charSize) __attribute__((noreturn)); 355 356// In order for standard assert() macro on linux to work we need to 357// provide device-side __assert_fail() 358__device__ static inline void __assert_fail(const char *__message, 359 const char *__file, unsigned __line, 360 const char *__function) { 361 __assertfail(__message, __file, __line, __function, sizeof(char)); 362} 363 364// Clang will convert printf into vprintf, but we still need 365// device-side declaration for it. 366__device__ int printf(const char *, ...); 367} // extern "C" 368 369// We also need device-side std::malloc and std::free. 370namespace std { 371__device__ static inline void free(void *__ptr) { ::free(__ptr); } 372__device__ static inline void *malloc(size_t __size) { 373 return ::malloc(__size); 374} 375} // namespace std 376 377// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to 378// come after we've pulled in the definition of uint3 and dim3. 379 380__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { 381 uint3 ret; 382 ret.x = x; 383 ret.y = y; 384 ret.z = z; 385 return ret; 386} 387 388__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { 389 uint3 ret; 390 ret.x = x; 391 ret.y = y; 392 ret.z = z; 393 return ret; 394} 395 396__device__ inline __cuda_builtin_blockDim_t::operator dim3() const { 397 return dim3(x, y, z); 398} 399 400__device__ inline __cuda_builtin_gridDim_t::operator dim3() const { 401 return dim3(x, y, z); 402} 403 404#include <__clang_cuda_cmath.h> 405#include <__clang_cuda_intrinsics.h> 406#include <__clang_cuda_complex_builtins.h> 407 408// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host 409// mode, giving them their "proper" types of dim3 and uint3. This is 410// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as 411// hack, force-include the header (nvcc doesn't include it by default) but 412// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are 413// only used here for the redeclarations of blockDim and threadIdx.) 414#pragma push_macro("dim3") 415#pragma push_macro("uint3") 416#define dim3 __cuda_builtin_blockDim_t 417#define uint3 __cuda_builtin_threadIdx_t 418#include "curand_mtgp32_kernel.h" 419#pragma pop_macro("dim3") 420#pragma pop_macro("uint3") 421#pragma pop_macro("__USE_FAST_MATH__") 422 423#endif // __CUDA__ 424#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ 425