510 lines · c
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-exception6 *7 *===-----------------------------------------------------------------------===8 */9 10/*11 * WARNING: This header is intended to be directly -include'd by12 * the compiler and is not supposed to be included by users.13 *14 * CUDA headers are implemented in a way that currently makes it15 * impossible for user code to #include directly when compiling with16 * Clang. They present different view of CUDA-supplied functions17 * depending on where in NVCC's compilation pipeline the headers are18 * included. Neither of these modes provides function definitions with19 * correct attributes, so we use preprocessor to force the headers20 * into a form that Clang can use.21 *22 * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's23 * 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 extensions35// enabled depend on it to avoid using __float128, which is unsupported in36// CUDA.37#define __CUDACC__38 39// Include some standard headers to avoid CUDA headers including them40// 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 CUDA48// headers.49#pragma push_macro("__THROW")50#pragma push_macro("__CUDA_ARCH__")51 52// WARNING: Preprocessor hacks below are based on specific details of53// CUDA-7.x headers and are not expected to work with any other54// 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 < 700059#error "Unsupported CUDA version!"60#endif61 62#pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")63#if CUDA_VERSION >= 1000064#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__65#endif66 67// Make largest subset of device functions available during host68// compilation.69#ifndef __CUDA_ARCH__70#define __CUDA_ARCH__ 999971#endif72 73#include "__clang_cuda_builtin_vars.h"74 75// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above76// 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 the80// functions. We don't need them as we're going to pull in their81// 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.h86// included below.87#define __DEVICE_FUNCTIONS_DECLS_H__88 89#undef __CUDACC__90#if CUDA_VERSION < 900091#define __CUDABE__92#else93#define __CUDACC__94#define __CUDA_LIBDEVICE__95#endif96// Disables definitions of device-side runtime support stubs in97// cuda_device_runtime_api.h98#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)) in104// cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the105// functional equivalent of what we need.106#pragma push_macro("nv_weak")107#define nv_weak weak108#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 does118// 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 < 9000123#include "crt/device_runtime.h"124#endif125#include "crt/host_runtime.h"126// device_runtime.h defines __cxa_* macros that will conflict with127// cxxabi.h.128// FIXME: redefine these as __device__ functions.129#undef __cxa_vec_ctor130#undef __cxa_vec_cctor131#undef __cxa_vec_dtor132#undef __cxa_vec_new133#undef __cxa_vec_new2134#undef __cxa_vec_new3135#undef __cxa_vec_delete2136#undef __cxa_vec_delete137#undef __cxa_vec_delete3138#undef __cxa_pure_virtual139 140// math_functions.hpp expects this host function be defined on MacOS, but it141// ends up not being there because of the games we play here. Just define it142// ourselves; it's simple enough.143#ifdef __APPLE__144inline __host__ double __signbitd(double x) {145 return std::signbit(x);146}147#endif148 149// CUDA 9.1 no longer provides declarations for libdevice functions, so we need150// to provide our own.151#include <__clang_cuda_libdevice_declares.h>152 153// Wrappers for many device-side standard library functions, incl. math154// functions, became compiler builtins in CUDA-9 and have been removed from the155// CUDA headers. Clang now provides its own implementation of the wrappers.156#if CUDA_VERSION >= 9000157#include <__clang_cuda_device_functions.h>158#include <__clang_cuda_math.h>159#endif160 161// __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's162// counterpart does not do it, so we need to make it empty here to keep163// following CUDA includes happy.164#undef __THROW165#define __THROW166 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 it170// here to detect the switch.171 172#if defined(CU_DEVICE_INVALID)173#if !defined(__USE_FAST_MATH__)174#define __USE_FAST_MATH__ 0175#endif176 177#if !defined(__CUDA_PREC_DIV)178#define __CUDA_PREC_DIV 0179#endif180#endif181 182// Temporarily poison __host__ macro to ensure it's not used by any of183// the headers we're about to include.184#pragma push_macro("__host__")185#define __host__ UNEXPECTED_HOST_ATTRIBUTE186 187// device_functions.hpp and math_functions*.hpp use 'static188// __forceinline__' (with no __device__) for definitions of device189// functions. Temporarily redefine __forceinline__ to include190// __device__.191#pragma push_macro("__forceinline__")192#define __forceinline__ __device__ __inline__ __attribute__((always_inline))193#if CUDA_VERSION < 9000194#include "device_functions.hpp"195#endif196 197// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we198// get the slow-but-accurate or fast-but-inaccurate versions of functions like199// 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__ 1206#endif207 208#if CUDA_VERSION >= 9000209#include "crt/math_functions.hpp"210#else211#include "math_functions.hpp"212#endif213 214#pragma pop_macro("__USE_FAST_MATH__")215 216#if CUDA_VERSION < 9000217#include "math_functions_dbl_ptx3.hpp"218#endif219#pragma pop_macro("__forceinline__")220 221// Pull in host-only functions that are only available when neither222// __CUDACC__ nor __CUDABE__ are defined.223#undef __MATH_FUNCTIONS_HPP__224#undef __CUDABE__225#if CUDA_VERSION < 9000226#include "math_functions.hpp"227#endif228// 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.231static inline float rsqrt(float __a) { return rsqrtf(__a); }232static inline float rcbrt(float __a) { return rcbrtf(__a); }233static inline float sinpi(float __a) { return sinpif(__a); }234static inline float cospi(float __a) { return cospif(__a); }235static inline void sincospi(float __a, float *__b, float *__c) {236 return sincospif(__a, __b, __c);237}238static inline float erfcinv(float __a) { return erfcinvf(__a); }239static inline float normcdfinv(float __a) { return normcdfinvf(__a); }240static inline float normcdf(float __a) { return normcdff(__a); }241static inline float erfcx(float __a) { return erfcxf(__a); }242 243#if CUDA_VERSION < 9000244// For some reason single-argument variant is not always declared by245// CUDA headers. Alas, device_functions.hpp included below needs it.246static inline __device__ void __brkpt(int __c) { __brkpt(); }247#endif248 249// Now include *.hpp with definitions of various GPU functions. Alas,250// a lot of thins get declared/defined with __host__ attribute which251// we don't want and we have to define it out. We also have to include252// {device,math}_functions.hpp again in order to extract the other253// branch of #if/else inside.254#define __host__255#undef __CUDABE__256#define __CUDACC__257#if CUDA_VERSION >= 9000258// Some atomic functions became compiler builtins in CUDA-9 , so we need their259// declarations.260#include "device_atomic_functions.h"261#endif262#undef __DEVICE_FUNCTIONS_HPP__263#include "device_atomic_functions.hpp"264#if CUDA_VERSION >= 9000265#include "crt/device_functions.hpp"266#include "crt/device_double_functions.hpp"267#else268#include "device_functions.hpp"269#define __CUDABE__270#include "device_double_functions.h"271#undef __CUDABE__272#endif273#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 nor276// const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define277// 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___isGlobal283#define __isShared __ignored_cuda___isShared284#define __isConstant __ignored_cuda___isConstant285#define __isLocal __ignored_cuda___isLocal286#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 the309// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to310// define them using builtins so that the optimizer can reason about and across311// these instructions. In particular, using intrinsics for ldg gets us the312// [addr+imm] addressing mode, which, although it doesn't actually exist in the313// hardware, seems to generate faster machine code because ptxas can more easily314// reason about our code.315 316#if CUDA_VERSION >= 8000317#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#endif323 324#undef __MATH_FUNCTIONS_HPP__325 326// math_functions.hpp defines ::signbit as a __host__ __device__ function. This327// conflicts with libstdc++'s constexpr ::signbit, so we have to rename328// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's329// conditional on __GNUC__. :)330#pragma push_macro("signbit")331#pragma push_macro("__GNUC__")332#undef __GNUC__333#define signbit __ignored_cuda_signbit334 335// CUDA-9 omits device-side definitions of some math functions if it sees336// include guard from math.h wrapper from libstdc++. We have to undo the header337// 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 >= 9000341#undef _GLIBCXX_MATH_H342// We also need to undo another guard that checks for libc++ 3.8+343#ifdef _LIBCPP_VERSION344#define _LIBCPP_VERSION 3700345#endif346#endif347 348#if CUDA_VERSION >= 9000349#include "crt/math_functions.hpp"350#else351#include "math_functions.hpp"352#endif353#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 provide361// 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 >= 9000364// clang-format off365#include <__clang_cuda_texture_intrinsics.h>366// clang-format on367#else368#if CUDA_VERSION >= 9000369// Provide a hint that texture support needs C++11.370template <typename T> struct __nv_tex_needs_cxx11 {371 const static bool value = false;372};373template <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#else381// Textures in CUDA-8 and older are not supported by clang.There's no382// convenient way to intercept texture use in these versions, so we can't383// produce a meaningful error. The source code that attempts to use textures384// will continue to fail as it does now.385#endif // CUDA_VERSION386#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000387#include "surface_indirect_functions.h"388#if CUDA_VERSION < 13000389// Direct texture fetch functions had been deprecated since CUDA-11.390// The file in CUDA-12 only carried unused texture types, and is no longer391// needed.392#include "texture_fetch_functions.h"393#endif // CUDA_VERSION < 13000394#include "texture_indirect_functions.h"395 396// Restore state of __CUDA_ARCH__ and __THROW we had on entry.397#pragma pop_macro("__CUDA_ARCH__")398#pragma pop_macro("__THROW")399 400// Set up compiler macros expected to be seen during compilation.401#undef __CUDABE__402#define __CUDACC__403 404extern "C" {405// Device-side CUDA system calls.406// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls407// We need these declarations and wrappers for device-side408// malloc/free/printf calls to work without relying on409// -fcuda-disable-target-call-checks option.410__device__ int vprintf(const char *, const char *);411__device__ void free(void *) __attribute((nothrow));412__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));413 414// __assertfail() used to have a `noreturn` attribute. Unfortunately that415// contributed to triggering the longstanding bug in ptxas when assert was used416// in sufficiently convoluted code. See417// https://bugs.llvm.org/show_bug.cgi?id=27738 for the details.418__device__ void __assertfail(const char *__message, const char *__file,419 unsigned __line, const char *__function,420 size_t __charSize);421 422// In order for standard assert() macro on linux to work we need to423// provide device-side __assert_fail()424__device__ static inline void __assert_fail(const char *__message,425 const char *__file, unsigned __line,426 const char *__function) {427 __assertfail(__message, __file, __line, __function, sizeof(char));428}429 430// Clang will convert printf into vprintf, but we still need431// device-side declaration for it.432__device__ int printf(const char *, ...);433} // extern "C"434 435// We also need device-side std::malloc and std::free.436namespace std {437__device__ static inline void free(void *__ptr) { ::free(__ptr); }438__device__ static inline void *malloc(size_t __size) {439 return ::malloc(__size);440}441} // namespace std442 443// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to444// come after we've pulled in the definition of uint3 and dim3.445 446__device__ inline __cuda_builtin_threadIdx_t::operator dim3() const {447 return dim3(x, y, z);448}449 450__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {451 return {x, y, z};452}453 454__device__ inline __cuda_builtin_blockIdx_t::operator dim3() const {455 return dim3(x, y, z);456}457 458__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {459 return {x, y, z};460}461 462__device__ inline __cuda_builtin_blockDim_t::operator dim3() const {463 return dim3(x, y, z);464}465 466__device__ inline __cuda_builtin_blockDim_t::operator uint3() const {467 return {x, y, z};468}469 470__device__ inline __cuda_builtin_gridDim_t::operator dim3() const {471 return dim3(x, y, z);472}473 474__device__ inline __cuda_builtin_gridDim_t::operator uint3() const {475 return {x, y, z};476}477 478#include <__clang_cuda_cmath.h>479#include <__clang_cuda_intrinsics.h>480#include <__clang_cuda_complex_builtins.h>481 482// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host483// mode, giving them their "proper" types of dim3 and uint3. This is484// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as485// hack, force-include the header (nvcc doesn't include it by default) but486// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are487// only used here for the redeclarations of blockDim and threadIdx.)488#pragma push_macro("dim3")489#pragma push_macro("uint3")490#define dim3 __cuda_builtin_blockDim_t491#define uint3 __cuda_builtin_threadIdx_t492#include "curand_mtgp32_kernel.h"493#pragma pop_macro("dim3")494#pragma pop_macro("uint3")495#pragma pop_macro("__USE_FAST_MATH__")496#pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")497 498// CUDA runtime uses this undocumented function to access kernel launch499// configuration. The declaration is in crt/device_functions.h but that file500// includes a lot of other stuff we don't want. Instead, we'll provide our own501// declaration for it here.502#if CUDA_VERSION >= 9020503extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,504 size_t sharedMem = 0,505 void *stream = 0);506#endif507 508#endif // __CUDA__509#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__510