__clang_cuda_runtime_wrapper.h revision 314564
1292920Sdim/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------=== 2292920Sdim * 3292920Sdim * Permission is hereby granted, free of charge, to any person obtaining a copy 4292920Sdim * of this software and associated documentation files (the "Software"), to deal 5292920Sdim * in the Software without restriction, including without limitation the rights 6292920Sdim * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell 7292920Sdim * copies of the Software, and to permit persons to whom the Software is 8292920Sdim * furnished to do so, subject to the following conditions: 9292920Sdim * 10292920Sdim * The above copyright notice and this permission notice shall be included in 11292920Sdim * all copies or substantial portions of the Software. 12292920Sdim * 13292920Sdim * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 14292920Sdim * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 15292920Sdim * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 16292920Sdim * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 17292920Sdim * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 18292920Sdim * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN 19292920Sdim * THE SOFTWARE. 20292920Sdim * 21292920Sdim *===-----------------------------------------------------------------------=== 22292920Sdim */ 23292920Sdim 24292920Sdim/* 25292920Sdim * WARNING: This header is intended to be directly -include'd by 26292920Sdim * the compiler and is not supposed to be included by users. 27292920Sdim * 28292920Sdim * CUDA headers are implemented in a way that currently makes it 29292920Sdim * impossible for user code to #include directly when compiling with 30292920Sdim * Clang. They present different view of CUDA-supplied functions 31292920Sdim * depending on where in NVCC's compilation pipeline the headers are 32292920Sdim * included. Neither of these modes provides function definitions with 33292920Sdim * correct attributes, so we use preprocessor to force the headers 34292920Sdim * into a form that Clang can use. 35292920Sdim * 36292920Sdim * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's 37292920Sdim * this file during every CUDA compilation. 38292920Sdim */ 39292920Sdim 40292920Sdim#ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__ 41292920Sdim#define __CLANG_CUDA_RUNTIME_WRAPPER_H__ 42292920Sdim 43292920Sdim#if defined(__CUDA__) && defined(__clang__) 44292920Sdim 45309124Sdim// Include some forward declares that must come before cmath. 46309124Sdim#include <__clang_cuda_math_forward_declares.h> 47309124Sdim 48292920Sdim// Include some standard headers to avoid CUDA headers including them 49292920Sdim// while some required macros (like __THROW) are in a weird state. 50309124Sdim#include <cmath> 51309124Sdim#include <cstdlib> 52292920Sdim#include <stdlib.h> 53292920Sdim 54292920Sdim// Preserve common macros that will be changed below by us or by CUDA 55292920Sdim// headers. 56292920Sdim#pragma push_macro("__THROW") 57292920Sdim#pragma push_macro("__CUDA_ARCH__") 58292920Sdim 59292920Sdim// WARNING: Preprocessor hacks below are based on specific details of 60292920Sdim// CUDA-7.x headers and are not expected to work with any other 61292920Sdim// version of CUDA headers. 62292920Sdim#include "cuda.h" 63292920Sdim#if !defined(CUDA_VERSION) 64292920Sdim#error "cuda.h did not define CUDA_VERSION" 65314564Sdim#elif CUDA_VERSION < 7000 || CUDA_VERSION > 8000 66292920Sdim#error "Unsupported CUDA version!" 67292920Sdim#endif 68292920Sdim 69292920Sdim// Make largest subset of device functions available during host 70292920Sdim// compilation -- SM_35 for the time being. 71292920Sdim#ifndef __CUDA_ARCH__ 72292920Sdim#define __CUDA_ARCH__ 350 73292920Sdim#endif 74292920Sdim 75314564Sdim#include "__clang_cuda_builtin_vars.h" 76292920Sdim 77314564Sdim// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above 78292920Sdim// has taken care of builtin variables declared in the file. 79292920Sdim#define __DEVICE_LAUNCH_PARAMETERS_H__ 80292920Sdim 81292920Sdim// {math,device}_functions.h only have declarations of the 82292920Sdim// functions. We don't need them as we're going to pull in their 83292920Sdim// definitions from .hpp files. 84292920Sdim#define __DEVICE_FUNCTIONS_H__ 85292920Sdim#define __MATH_FUNCTIONS_H__ 86309124Sdim#define __COMMON_FUNCTIONS_H__ 87292920Sdim 88292920Sdim#undef __CUDACC__ 89292920Sdim#define __CUDABE__ 90292920Sdim// Disables definitions of device-side runtime support stubs in 91292920Sdim// cuda_device_runtime_api.h 92309124Sdim#include "driver_types.h" 93292920Sdim#include "host_config.h" 94292920Sdim#include "host_defines.h" 95292920Sdim 96292920Sdim#undef __CUDABE__ 97292920Sdim#define __CUDACC__ 98292920Sdim#include "cuda_runtime.h" 99292920Sdim 100292920Sdim#undef __CUDACC__ 101292920Sdim#define __CUDABE__ 102292920Sdim 103292920Sdim// CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does 104292920Sdim// not have at the moment. Emulate them with a builtin memcpy/memset. 105309124Sdim#define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n) 106309124Sdim#define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n) 107292920Sdim 108309124Sdim#include "crt/device_runtime.h" 109292920Sdim#include "crt/host_runtime.h" 110292920Sdim// device_runtime.h defines __cxa_* macros that will conflict with 111292920Sdim// cxxabi.h. 112292920Sdim// FIXME: redefine these as __device__ functions. 113292920Sdim#undef __cxa_vec_ctor 114292920Sdim#undef __cxa_vec_cctor 115292920Sdim#undef __cxa_vec_dtor 116314564Sdim#undef __cxa_vec_new 117292920Sdim#undef __cxa_vec_new2 118292920Sdim#undef __cxa_vec_new3 119292920Sdim#undef __cxa_vec_delete2 120292920Sdim#undef __cxa_vec_delete 121292920Sdim#undef __cxa_vec_delete3 122292920Sdim#undef __cxa_pure_virtual 123292920Sdim 124314564Sdim// math_functions.hpp expects this host function be defined on MacOS, but it 125314564Sdim// ends up not being there because of the games we play here. Just define it 126314564Sdim// ourselves; it's simple enough. 127314564Sdim#ifdef __APPLE__ 128314564Sdiminline __host__ double __signbitd(double x) { 129314564Sdim return std::signbit(x); 130314564Sdim} 131314564Sdim#endif 132314564Sdim 133292920Sdim// We need decls for functions in CUDA's libdevice with __device__ 134292920Sdim// attribute only. Alas they come either as __host__ __device__ or 135292920Sdim// with no attributes at all. To work around that, define __CUDA_RTC__ 136292920Sdim// which produces HD variant and undef __host__ which gives us desided 137292920Sdim// decls with __device__ attribute. 138292920Sdim#pragma push_macro("__host__") 139292920Sdim#define __host__ 140292920Sdim#define __CUDACC_RTC__ 141292920Sdim#include "device_functions_decls.h" 142292920Sdim#undef __CUDACC_RTC__ 143292920Sdim 144292920Sdim// Temporarily poison __host__ macro to ensure it's not used by any of 145292920Sdim// the headers we're about to include. 146292920Sdim#define __host__ UNEXPECTED_HOST_ATTRIBUTE 147292920Sdim 148314564Sdim// CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values. 149314564Sdim// Previous versions used to check whether they are defined or not. 150314564Sdim// CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it 151314564Sdim// here to detect the switch. 152314564Sdim 153314564Sdim#if defined(CU_DEVICE_INVALID) 154314564Sdim#if !defined(__USE_FAST_MATH__) 155314564Sdim#define __USE_FAST_MATH__ 0 156314564Sdim#endif 157314564Sdim 158314564Sdim#if !defined(__CUDA_PREC_DIV) 159314564Sdim#define __CUDA_PREC_DIV 0 160314564Sdim#endif 161314564Sdim#endif 162314564Sdim 163292920Sdim// device_functions.hpp and math_functions*.hpp use 'static 164292920Sdim// __forceinline__' (with no __device__) for definitions of device 165292920Sdim// functions. Temporarily redefine __forceinline__ to include 166292920Sdim// __device__. 167292920Sdim#pragma push_macro("__forceinline__") 168292920Sdim#define __forceinline__ __device__ __inline__ __attribute__((always_inline)) 169292920Sdim#include "device_functions.hpp" 170309124Sdim 171309124Sdim// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we 172309124Sdim// get the slow-but-accurate or fast-but-inaccurate versions of functions like 173309124Sdim// sin and exp. This is controlled in clang by -fcuda-approx-transcendentals. 174309124Sdim// 175309124Sdim// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs. 176309124Sdim// slow divides), so we need to scope our define carefully here. 177309124Sdim#pragma push_macro("__USE_FAST_MATH__") 178309124Sdim#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__) 179314564Sdim#define __USE_FAST_MATH__ 1 180309124Sdim#endif 181292920Sdim#include "math_functions.hpp" 182309124Sdim#pragma pop_macro("__USE_FAST_MATH__") 183309124Sdim 184292920Sdim#include "math_functions_dbl_ptx3.hpp" 185292920Sdim#pragma pop_macro("__forceinline__") 186292920Sdim 187292920Sdim// Pull in host-only functions that are only available when neither 188292920Sdim// __CUDACC__ nor __CUDABE__ are defined. 189292920Sdim#undef __MATH_FUNCTIONS_HPP__ 190292920Sdim#undef __CUDABE__ 191292920Sdim#include "math_functions.hpp" 192292920Sdim// Alas, additional overloads for these functions are hard to get to. 193292920Sdim// Considering that we only need these overloads for a few functions, 194292920Sdim// we can provide them here. 195309124Sdimstatic inline float rsqrt(float __a) { return rsqrtf(__a); } 196309124Sdimstatic inline float rcbrt(float __a) { return rcbrtf(__a); } 197309124Sdimstatic inline float sinpi(float __a) { return sinpif(__a); } 198309124Sdimstatic inline float cospi(float __a) { return cospif(__a); } 199309124Sdimstatic inline void sincospi(float __a, float *__b, float *__c) { 200309124Sdim return sincospif(__a, __b, __c); 201292920Sdim} 202309124Sdimstatic inline float erfcinv(float __a) { return erfcinvf(__a); } 203309124Sdimstatic inline float normcdfinv(float __a) { return normcdfinvf(__a); } 204309124Sdimstatic inline float normcdf(float __a) { return normcdff(__a); } 205309124Sdimstatic inline float erfcx(float __a) { return erfcxf(__a); } 206292920Sdim 207292920Sdim// For some reason single-argument variant is not always declared by 208292920Sdim// CUDA headers. Alas, device_functions.hpp included below needs it. 209309124Sdimstatic inline __device__ void __brkpt(int __c) { __brkpt(); } 210292920Sdim 211292920Sdim// Now include *.hpp with definitions of various GPU functions. Alas, 212292920Sdim// a lot of thins get declared/defined with __host__ attribute which 213292920Sdim// we don't want and we have to define it out. We also have to include 214292920Sdim// {device,math}_functions.hpp again in order to extract the other 215292920Sdim// branch of #if/else inside. 216292920Sdim 217292920Sdim#define __host__ 218292920Sdim#undef __CUDABE__ 219292920Sdim#define __CUDACC__ 220292920Sdim#undef __DEVICE_FUNCTIONS_HPP__ 221309124Sdim#include "device_atomic_functions.hpp" 222292920Sdim#include "device_functions.hpp" 223292920Sdim#include "sm_20_atomic_functions.hpp" 224309124Sdim#include "sm_20_intrinsics.hpp" 225292920Sdim#include "sm_32_atomic_functions.hpp" 226309124Sdim 227309124Sdim// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the 228309124Sdim// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to 229309124Sdim// define them using builtins so that the optimizer can reason about and across 230309124Sdim// these instructions. In particular, using intrinsics for ldg gets us the 231309124Sdim// [addr+imm] addressing mode, which, although it doesn't actually exist in the 232309124Sdim// hardware, seems to generate faster machine code because ptxas can more easily 233309124Sdim// reason about our code. 234309124Sdim 235314564Sdim#if CUDA_VERSION >= 8000 236314564Sdim#include "sm_60_atomic_functions.hpp" 237314564Sdim#include "sm_61_intrinsics.hpp" 238314564Sdim#endif 239314564Sdim 240292920Sdim#undef __MATH_FUNCTIONS_HPP__ 241309124Sdim 242309124Sdim// math_functions.hpp defines ::signbit as a __host__ __device__ function. This 243309124Sdim// conflicts with libstdc++'s constexpr ::signbit, so we have to rename 244309124Sdim// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's 245309124Sdim// conditional on __GNUC__. :) 246309124Sdim#pragma push_macro("signbit") 247309124Sdim#pragma push_macro("__GNUC__") 248309124Sdim#undef __GNUC__ 249309124Sdim#define signbit __ignored_cuda_signbit 250292920Sdim#include "math_functions.hpp" 251309124Sdim#pragma pop_macro("__GNUC__") 252309124Sdim#pragma pop_macro("signbit") 253309124Sdim 254292920Sdim#pragma pop_macro("__host__") 255292920Sdim 256292920Sdim#include "texture_indirect_functions.h" 257292920Sdim 258292920Sdim// Restore state of __CUDA_ARCH__ and __THROW we had on entry. 259292920Sdim#pragma pop_macro("__CUDA_ARCH__") 260292920Sdim#pragma pop_macro("__THROW") 261292920Sdim 262292920Sdim// Set up compiler macros expected to be seen during compilation. 263292920Sdim#undef __CUDABE__ 264292920Sdim#define __CUDACC__ 265292920Sdim 266309124Sdimextern "C" { 267309124Sdim// Device-side CUDA system calls. 268309124Sdim// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls 269309124Sdim// We need these declarations and wrappers for device-side 270309124Sdim// malloc/free/printf calls to work without relying on 271309124Sdim// -fcuda-disable-target-call-checks option. 272309124Sdim__device__ int vprintf(const char *, const char *); 273309124Sdim__device__ void free(void *) __attribute((nothrow)); 274309124Sdim__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); 275309124Sdim__device__ void __assertfail(const char *__message, const char *__file, 276309124Sdim unsigned __line, const char *__function, 277309124Sdim size_t __charSize) __attribute__((noreturn)); 278309124Sdim 279309124Sdim// In order for standard assert() macro on linux to work we need to 280309124Sdim// provide device-side __assert_fail() 281309124Sdim__device__ static inline void __assert_fail(const char *__message, 282309124Sdim const char *__file, unsigned __line, 283309124Sdim const char *__function) { 284309124Sdim __assertfail(__message, __file, __line, __function, sizeof(char)); 285292920Sdim} 286292920Sdim 287309124Sdim// Clang will convert printf into vprintf, but we still need 288309124Sdim// device-side declaration for it. 289309124Sdim__device__ int printf(const char *, ...); 290309124Sdim} // extern "C" 291309124Sdim 292309124Sdim// We also need device-side std::malloc and std::free. 293309124Sdimnamespace std { 294309124Sdim__device__ static inline void free(void *__ptr) { ::free(__ptr); } 295309124Sdim__device__ static inline void *malloc(size_t __size) { 296309124Sdim return ::malloc(__size); 297309124Sdim} 298309124Sdim} // namespace std 299309124Sdim 300314564Sdim// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to 301314564Sdim// come after we've pulled in the definition of uint3 and dim3. 302309124Sdim 303309124Sdim__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { 304309124Sdim uint3 ret; 305309124Sdim ret.x = x; 306309124Sdim ret.y = y; 307309124Sdim ret.z = z; 308309124Sdim return ret; 309309124Sdim} 310309124Sdim 311309124Sdim__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { 312309124Sdim uint3 ret; 313309124Sdim ret.x = x; 314309124Sdim ret.y = y; 315309124Sdim ret.z = z; 316309124Sdim return ret; 317309124Sdim} 318309124Sdim 319309124Sdim__device__ inline __cuda_builtin_blockDim_t::operator dim3() const { 320309124Sdim return dim3(x, y, z); 321309124Sdim} 322309124Sdim 323309124Sdim__device__ inline __cuda_builtin_gridDim_t::operator dim3() const { 324309124Sdim return dim3(x, y, z); 325309124Sdim} 326309124Sdim 327309124Sdim#include <__clang_cuda_cmath.h> 328309124Sdim#include <__clang_cuda_intrinsics.h> 329314564Sdim#include <__clang_cuda_complex_builtins.h> 330309124Sdim 331309124Sdim// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host 332309124Sdim// mode, giving them their "proper" types of dim3 and uint3. This is 333314564Sdim// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as 334314564Sdim// hack, force-include the header (nvcc doesn't include it by default) but 335314564Sdim// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are 336314564Sdim// only used here for the redeclarations of blockDim and threadIdx.) 337309124Sdim#pragma push_macro("dim3") 338309124Sdim#pragma push_macro("uint3") 339309124Sdim#define dim3 __cuda_builtin_blockDim_t 340309124Sdim#define uint3 __cuda_builtin_threadIdx_t 341309124Sdim#include "curand_mtgp32_kernel.h" 342309124Sdim#pragma pop_macro("dim3") 343309124Sdim#pragma pop_macro("uint3") 344309124Sdim#pragma pop_macro("__USE_FAST_MATH__") 345309124Sdim 346292920Sdim#endif // __CUDA__ 347292920Sdim#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ 348