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