1292920Sdim/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------=== 2292920Sdim * 3353358Sdim * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4353358Sdim * See https://llvm.org/LICENSE.txt for license information. 5353358Sdim * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6292920Sdim * 7292920Sdim *===-----------------------------------------------------------------------=== 8292920Sdim */ 9292920Sdim 10292920Sdim/* 11292920Sdim * WARNING: This header is intended to be directly -include'd by 12292920Sdim * the compiler and is not supposed to be included by users. 13292920Sdim * 14292920Sdim * CUDA headers are implemented in a way that currently makes it 15292920Sdim * impossible for user code to #include directly when compiling with 16292920Sdim * Clang. They present different view of CUDA-supplied functions 17292920Sdim * depending on where in NVCC's compilation pipeline the headers are 18292920Sdim * included. Neither of these modes provides function definitions with 19292920Sdim * correct attributes, so we use preprocessor to force the headers 20292920Sdim * into a form that Clang can use. 21292920Sdim * 22292920Sdim * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's 23292920Sdim * this file during every CUDA compilation. 24292920Sdim */ 25292920Sdim 26292920Sdim#ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__ 27292920Sdim#define __CLANG_CUDA_RUNTIME_WRAPPER_H__ 28292920Sdim 29292920Sdim#if defined(__CUDA__) && defined(__clang__) 30292920Sdim 31309124Sdim// Include some forward declares that must come before cmath. 32309124Sdim#include <__clang_cuda_math_forward_declares.h> 33309124Sdim 34292920Sdim// Include some standard headers to avoid CUDA headers including them 35292920Sdim// while some required macros (like __THROW) are in a weird state. 36309124Sdim#include <cmath> 37309124Sdim#include <cstdlib> 38292920Sdim#include <stdlib.h> 39292920Sdim 40292920Sdim// Preserve common macros that will be changed below by us or by CUDA 41292920Sdim// headers. 42292920Sdim#pragma push_macro("__THROW") 43292920Sdim#pragma push_macro("__CUDA_ARCH__") 44292920Sdim 45292920Sdim// WARNING: Preprocessor hacks below are based on specific details of 46292920Sdim// CUDA-7.x headers and are not expected to work with any other 47292920Sdim// version of CUDA headers. 48292920Sdim#include "cuda.h" 49292920Sdim#if !defined(CUDA_VERSION) 50292920Sdim#error "cuda.h did not define CUDA_VERSION" 51360784Sdim#elif CUDA_VERSION < 7000 52292920Sdim#error "Unsupported CUDA version!" 53292920Sdim#endif 54292920Sdim 55344779Sdim#pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") 56344779Sdim#if CUDA_VERSION >= 10000 57344779Sdim#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__ 58344779Sdim#endif 59344779Sdim 60292920Sdim// Make largest subset of device functions available during host 61292920Sdim// compilation -- SM_35 for the time being. 62292920Sdim#ifndef __CUDA_ARCH__ 63292920Sdim#define __CUDA_ARCH__ 350 64292920Sdim#endif 65292920Sdim 66314564Sdim#include "__clang_cuda_builtin_vars.h" 67292920Sdim 68314564Sdim// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above 69292920Sdim// has taken care of builtin variables declared in the file. 70292920Sdim#define __DEVICE_LAUNCH_PARAMETERS_H__ 71292920Sdim 72292920Sdim// {math,device}_functions.h only have declarations of the 73292920Sdim// functions. We don't need them as we're going to pull in their 74292920Sdim// definitions from .hpp files. 75292920Sdim#define __DEVICE_FUNCTIONS_H__ 76292920Sdim#define __MATH_FUNCTIONS_H__ 77309124Sdim#define __COMMON_FUNCTIONS_H__ 78341825Sdim// device_functions_decls is replaced by __clang_cuda_device_functions.h 79341825Sdim// included below. 80341825Sdim#define __DEVICE_FUNCTIONS_DECLS_H__ 81292920Sdim 82292920Sdim#undef __CUDACC__ 83327952Sdim#if CUDA_VERSION < 9000 84292920Sdim#define __CUDABE__ 85327952Sdim#else 86327952Sdim#define __CUDA_LIBDEVICE__ 87327952Sdim#endif 88292920Sdim// Disables definitions of device-side runtime support stubs in 89292920Sdim// cuda_device_runtime_api.h 90309124Sdim#include "driver_types.h" 91292920Sdim#include "host_config.h" 92292920Sdim#include "host_defines.h" 93292920Sdim 94341825Sdim// Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in 95341825Sdim// cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the 96341825Sdim// functional equivalent of what we need. 97341825Sdim#pragma push_macro("nv_weak") 98341825Sdim#define nv_weak weak 99292920Sdim#undef __CUDABE__ 100327952Sdim#undef __CUDA_LIBDEVICE__ 101292920Sdim#define __CUDACC__ 102292920Sdim#include "cuda_runtime.h" 103292920Sdim 104341825Sdim#pragma pop_macro("nv_weak") 105292920Sdim#undef __CUDACC__ 106292920Sdim#define __CUDABE__ 107292920Sdim 108292920Sdim// CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does 109292920Sdim// not have at the moment. Emulate them with a builtin memcpy/memset. 110309124Sdim#define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n) 111309124Sdim#define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n) 112292920Sdim 113327952Sdim#if CUDA_VERSION < 9000 114309124Sdim#include "crt/device_runtime.h" 115327952Sdim#endif 116292920Sdim#include "crt/host_runtime.h" 117292920Sdim// device_runtime.h defines __cxa_* macros that will conflict with 118292920Sdim// cxxabi.h. 119292920Sdim// FIXME: redefine these as __device__ functions. 120292920Sdim#undef __cxa_vec_ctor 121292920Sdim#undef __cxa_vec_cctor 122292920Sdim#undef __cxa_vec_dtor 123314564Sdim#undef __cxa_vec_new 124292920Sdim#undef __cxa_vec_new2 125292920Sdim#undef __cxa_vec_new3 126292920Sdim#undef __cxa_vec_delete2 127292920Sdim#undef __cxa_vec_delete 128292920Sdim#undef __cxa_vec_delete3 129292920Sdim#undef __cxa_pure_virtual 130292920Sdim 131314564Sdim// math_functions.hpp expects this host function be defined on MacOS, but it 132314564Sdim// ends up not being there because of the games we play here. Just define it 133314564Sdim// ourselves; it's simple enough. 134314564Sdim#ifdef __APPLE__ 135314564Sdiminline __host__ double __signbitd(double x) { 136314564Sdim return std::signbit(x); 137314564Sdim} 138314564Sdim#endif 139314564Sdim 140341825Sdim// CUDA 9.1 no longer provides declarations for libdevice functions, so we need 141341825Sdim// to provide our own. 142341825Sdim#include <__clang_cuda_libdevice_declares.h> 143292920Sdim 144341825Sdim// Wrappers for many device-side standard library functions became compiler 145341825Sdim// builtins in CUDA-9 and have been removed from the CUDA headers. Clang now 146341825Sdim// provides its own implementation of the wrappers. 147341825Sdim#if CUDA_VERSION >= 9000 148341825Sdim#include <__clang_cuda_device_functions.h> 149341825Sdim#endif 150292920Sdim 151341825Sdim// __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's 152341825Sdim// counterpart does not do it, so we need to make it empty here to keep 153341825Sdim// following CUDA includes happy. 154341825Sdim#undef __THROW 155341825Sdim#define __THROW 156341825Sdim 157314564Sdim// CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values. 158314564Sdim// Previous versions used to check whether they are defined or not. 159314564Sdim// CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it 160314564Sdim// here to detect the switch. 161314564Sdim 162314564Sdim#if defined(CU_DEVICE_INVALID) 163314564Sdim#if !defined(__USE_FAST_MATH__) 164314564Sdim#define __USE_FAST_MATH__ 0 165314564Sdim#endif 166314564Sdim 167314564Sdim#if !defined(__CUDA_PREC_DIV) 168314564Sdim#define __CUDA_PREC_DIV 0 169314564Sdim#endif 170314564Sdim#endif 171314564Sdim 172341825Sdim// Temporarily poison __host__ macro to ensure it's not used by any of 173341825Sdim// the headers we're about to include. 174341825Sdim#pragma push_macro("__host__") 175341825Sdim#define __host__ UNEXPECTED_HOST_ATTRIBUTE 176341825Sdim 177292920Sdim// device_functions.hpp and math_functions*.hpp use 'static 178292920Sdim// __forceinline__' (with no __device__) for definitions of device 179292920Sdim// functions. Temporarily redefine __forceinline__ to include 180292920Sdim// __device__. 181292920Sdim#pragma push_macro("__forceinline__") 182292920Sdim#define __forceinline__ __device__ __inline__ __attribute__((always_inline)) 183341825Sdim#if CUDA_VERSION < 9000 184341825Sdim#include "device_functions.hpp" 185327952Sdim#endif 186327952Sdim 187309124Sdim// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we 188309124Sdim// get the slow-but-accurate or fast-but-inaccurate versions of functions like 189309124Sdim// sin and exp. This is controlled in clang by -fcuda-approx-transcendentals. 190309124Sdim// 191309124Sdim// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs. 192309124Sdim// slow divides), so we need to scope our define carefully here. 193309124Sdim#pragma push_macro("__USE_FAST_MATH__") 194309124Sdim#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__) 195314564Sdim#define __USE_FAST_MATH__ 1 196309124Sdim#endif 197341825Sdim 198341825Sdim#if CUDA_VERSION >= 9000 199341825Sdim// CUDA-9.2 needs host-side memcpy for some host functions in 200341825Sdim// device_functions.hpp 201341825Sdim#if CUDA_VERSION >= 9020 202341825Sdim#include <string.h> 203341825Sdim#endif 204341825Sdim#include "crt/math_functions.hpp" 205341825Sdim#else 206292920Sdim#include "math_functions.hpp" 207341825Sdim#endif 208341825Sdim 209309124Sdim#pragma pop_macro("__USE_FAST_MATH__") 210309124Sdim 211341825Sdim#if CUDA_VERSION < 9000 212292920Sdim#include "math_functions_dbl_ptx3.hpp" 213341825Sdim#endif 214292920Sdim#pragma pop_macro("__forceinline__") 215292920Sdim 216292920Sdim// Pull in host-only functions that are only available when neither 217292920Sdim// __CUDACC__ nor __CUDABE__ are defined. 218292920Sdim#undef __MATH_FUNCTIONS_HPP__ 219292920Sdim#undef __CUDABE__ 220341825Sdim#if CUDA_VERSION < 9000 221292920Sdim#include "math_functions.hpp" 222341825Sdim#endif 223292920Sdim// Alas, additional overloads for these functions are hard to get to. 224292920Sdim// Considering that we only need these overloads for a few functions, 225292920Sdim// we can provide them here. 226309124Sdimstatic inline float rsqrt(float __a) { return rsqrtf(__a); } 227309124Sdimstatic inline float rcbrt(float __a) { return rcbrtf(__a); } 228309124Sdimstatic inline float sinpi(float __a) { return sinpif(__a); } 229309124Sdimstatic inline float cospi(float __a) { return cospif(__a); } 230309124Sdimstatic inline void sincospi(float __a, float *__b, float *__c) { 231309124Sdim return sincospif(__a, __b, __c); 232292920Sdim} 233309124Sdimstatic inline float erfcinv(float __a) { return erfcinvf(__a); } 234309124Sdimstatic inline float normcdfinv(float __a) { return normcdfinvf(__a); } 235309124Sdimstatic inline float normcdf(float __a) { return normcdff(__a); } 236309124Sdimstatic inline float erfcx(float __a) { return erfcxf(__a); } 237292920Sdim 238341825Sdim#if CUDA_VERSION < 9000 239292920Sdim// For some reason single-argument variant is not always declared by 240292920Sdim// CUDA headers. Alas, device_functions.hpp included below needs it. 241309124Sdimstatic inline __device__ void __brkpt(int __c) { __brkpt(); } 242341825Sdim#endif 243292920Sdim 244292920Sdim// Now include *.hpp with definitions of various GPU functions. Alas, 245292920Sdim// a lot of thins get declared/defined with __host__ attribute which 246292920Sdim// we don't want and we have to define it out. We also have to include 247292920Sdim// {device,math}_functions.hpp again in order to extract the other 248292920Sdim// branch of #if/else inside. 249292920Sdim#define __host__ 250292920Sdim#undef __CUDABE__ 251292920Sdim#define __CUDACC__ 252341825Sdim#if CUDA_VERSION >= 9000 253341825Sdim// Some atomic functions became compiler builtins in CUDA-9 , so we need their 254341825Sdim// declarations. 255341825Sdim#include "device_atomic_functions.h" 256341825Sdim#endif 257292920Sdim#undef __DEVICE_FUNCTIONS_HPP__ 258309124Sdim#include "device_atomic_functions.hpp" 259341825Sdim#if CUDA_VERSION >= 9000 260341825Sdim#include "crt/device_functions.hpp" 261341825Sdim#include "crt/device_double_functions.hpp" 262341825Sdim#else 263292920Sdim#include "device_functions.hpp" 264341825Sdim#define __CUDABE__ 265341825Sdim#include "device_double_functions.h" 266341825Sdim#undef __CUDABE__ 267341825Sdim#endif 268292920Sdim#include "sm_20_atomic_functions.hpp" 269309124Sdim#include "sm_20_intrinsics.hpp" 270292920Sdim#include "sm_32_atomic_functions.hpp" 271309124Sdim 272309124Sdim// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the 273309124Sdim// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to 274309124Sdim// define them using builtins so that the optimizer can reason about and across 275309124Sdim// these instructions. In particular, using intrinsics for ldg gets us the 276309124Sdim// [addr+imm] addressing mode, which, although it doesn't actually exist in the 277309124Sdim// hardware, seems to generate faster machine code because ptxas can more easily 278309124Sdim// reason about our code. 279309124Sdim 280314564Sdim#if CUDA_VERSION >= 8000 281341825Sdim#pragma push_macro("__CUDA_ARCH__") 282341825Sdim#undef __CUDA_ARCH__ 283314564Sdim#include "sm_60_atomic_functions.hpp" 284314564Sdim#include "sm_61_intrinsics.hpp" 285341825Sdim#pragma pop_macro("__CUDA_ARCH__") 286314564Sdim#endif 287314564Sdim 288292920Sdim#undef __MATH_FUNCTIONS_HPP__ 289309124Sdim 290309124Sdim// math_functions.hpp defines ::signbit as a __host__ __device__ function. This 291309124Sdim// conflicts with libstdc++'s constexpr ::signbit, so we have to rename 292309124Sdim// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's 293309124Sdim// conditional on __GNUC__. :) 294309124Sdim#pragma push_macro("signbit") 295309124Sdim#pragma push_macro("__GNUC__") 296309124Sdim#undef __GNUC__ 297309124Sdim#define signbit __ignored_cuda_signbit 298327952Sdim 299327952Sdim// CUDA-9 omits device-side definitions of some math functions if it sees 300327952Sdim// include guard from math.h wrapper from libstdc++. We have to undo the header 301327952Sdim// guard temporarily to get the definitions we need. 302327952Sdim#pragma push_macro("_GLIBCXX_MATH_H") 303327952Sdim#pragma push_macro("_LIBCPP_VERSION") 304327952Sdim#if CUDA_VERSION >= 9000 305327952Sdim#undef _GLIBCXX_MATH_H 306327952Sdim// We also need to undo another guard that checks for libc++ 3.8+ 307327952Sdim#ifdef _LIBCPP_VERSION 308327952Sdim#define _LIBCPP_VERSION 3700 309327952Sdim#endif 310327952Sdim#endif 311327952Sdim 312341825Sdim#if CUDA_VERSION >= 9000 313341825Sdim#include "crt/math_functions.hpp" 314341825Sdim#else 315292920Sdim#include "math_functions.hpp" 316341825Sdim#endif 317327952Sdim#pragma pop_macro("_GLIBCXX_MATH_H") 318327952Sdim#pragma pop_macro("_LIBCPP_VERSION") 319309124Sdim#pragma pop_macro("__GNUC__") 320309124Sdim#pragma pop_macro("signbit") 321309124Sdim 322292920Sdim#pragma pop_macro("__host__") 323292920Sdim 324292920Sdim#include "texture_indirect_functions.h" 325292920Sdim 326292920Sdim// Restore state of __CUDA_ARCH__ and __THROW we had on entry. 327292920Sdim#pragma pop_macro("__CUDA_ARCH__") 328292920Sdim#pragma pop_macro("__THROW") 329292920Sdim 330292920Sdim// Set up compiler macros expected to be seen during compilation. 331292920Sdim#undef __CUDABE__ 332292920Sdim#define __CUDACC__ 333292920Sdim 334309124Sdimextern "C" { 335309124Sdim// Device-side CUDA system calls. 336309124Sdim// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls 337309124Sdim// We need these declarations and wrappers for device-side 338309124Sdim// malloc/free/printf calls to work without relying on 339309124Sdim// -fcuda-disable-target-call-checks option. 340309124Sdim__device__ int vprintf(const char *, const char *); 341309124Sdim__device__ void free(void *) __attribute((nothrow)); 342309124Sdim__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); 343309124Sdim__device__ void __assertfail(const char *__message, const char *__file, 344309124Sdim unsigned __line, const char *__function, 345309124Sdim size_t __charSize) __attribute__((noreturn)); 346309124Sdim 347309124Sdim// In order for standard assert() macro on linux to work we need to 348309124Sdim// provide device-side __assert_fail() 349309124Sdim__device__ static inline void __assert_fail(const char *__message, 350309124Sdim const char *__file, unsigned __line, 351309124Sdim const char *__function) { 352309124Sdim __assertfail(__message, __file, __line, __function, sizeof(char)); 353292920Sdim} 354292920Sdim 355309124Sdim// Clang will convert printf into vprintf, but we still need 356309124Sdim// device-side declaration for it. 357309124Sdim__device__ int printf(const char *, ...); 358309124Sdim} // extern "C" 359309124Sdim 360309124Sdim// We also need device-side std::malloc and std::free. 361309124Sdimnamespace std { 362309124Sdim__device__ static inline void free(void *__ptr) { ::free(__ptr); } 363309124Sdim__device__ static inline void *malloc(size_t __size) { 364309124Sdim return ::malloc(__size); 365309124Sdim} 366309124Sdim} // namespace std 367309124Sdim 368314564Sdim// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to 369314564Sdim// come after we've pulled in the definition of uint3 and dim3. 370309124Sdim 371309124Sdim__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { 372309124Sdim uint3 ret; 373309124Sdim ret.x = x; 374309124Sdim ret.y = y; 375309124Sdim ret.z = z; 376309124Sdim return ret; 377309124Sdim} 378309124Sdim 379309124Sdim__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { 380309124Sdim uint3 ret; 381309124Sdim ret.x = x; 382309124Sdim ret.y = y; 383309124Sdim ret.z = z; 384309124Sdim return ret; 385309124Sdim} 386309124Sdim 387309124Sdim__device__ inline __cuda_builtin_blockDim_t::operator dim3() const { 388309124Sdim return dim3(x, y, z); 389309124Sdim} 390309124Sdim 391309124Sdim__device__ inline __cuda_builtin_gridDim_t::operator dim3() const { 392309124Sdim return dim3(x, y, z); 393309124Sdim} 394309124Sdim 395309124Sdim#include <__clang_cuda_cmath.h> 396309124Sdim#include <__clang_cuda_intrinsics.h> 397314564Sdim#include <__clang_cuda_complex_builtins.h> 398309124Sdim 399309124Sdim// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host 400309124Sdim// mode, giving them their "proper" types of dim3 and uint3. This is 401314564Sdim// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as 402314564Sdim// hack, force-include the header (nvcc doesn't include it by default) but 403314564Sdim// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are 404314564Sdim// only used here for the redeclarations of blockDim and threadIdx.) 405309124Sdim#pragma push_macro("dim3") 406309124Sdim#pragma push_macro("uint3") 407309124Sdim#define dim3 __cuda_builtin_blockDim_t 408309124Sdim#define uint3 __cuda_builtin_threadIdx_t 409309124Sdim#include "curand_mtgp32_kernel.h" 410309124Sdim#pragma pop_macro("dim3") 411309124Sdim#pragma pop_macro("uint3") 412309124Sdim#pragma pop_macro("__USE_FAST_MATH__") 413344779Sdim#pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") 414309124Sdim 415353358Sdim// CUDA runtime uses this undocumented function to access kernel launch 416353358Sdim// configuration. The declaration is in crt/device_functions.h but that file 417353358Sdim// includes a lot of other stuff we don't want. Instead, we'll provide our own 418353358Sdim// declaration for it here. 419353358Sdim#if CUDA_VERSION >= 9020 420353358Sdimextern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, 421353358Sdim size_t sharedMem = 0, 422353358Sdim void *stream = 0); 423353358Sdim#endif 424353358Sdim 425292920Sdim#endif // __CUDA__ 426292920Sdim#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ 427