diff options
Diffstat (limited to 'contrib/llvm-project/clang/lib/Headers/__clang_cuda_runtime_wrapper.h')
| -rw-r--r-- | contrib/llvm-project/clang/lib/Headers/__clang_cuda_runtime_wrapper.h | 503 |
1 files changed, 503 insertions, 0 deletions
diff --git a/contrib/llvm-project/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/contrib/llvm-project/clang/lib/Headers/__clang_cuda_runtime_wrapper.h new file mode 100644 index 000000000000..d369c86fe106 --- /dev/null +++ b/contrib/llvm-project/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -0,0 +1,503 @@ +/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +/* + * WARNING: This header is intended to be directly -include'd by + * the compiler and is not supposed to be included by users. + * + * CUDA headers are implemented in a way that currently makes it + * impossible for user code to #include directly when compiling with + * Clang. They present different view of CUDA-supplied functions + * depending on where in NVCC's compilation pipeline the headers are + * included. Neither of these modes provides function definitions with + * correct attributes, so we use preprocessor to force the headers + * into a form that Clang can use. + * + * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's + * this file during every CUDA compilation. + */ + +#ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__ +#define __CLANG_CUDA_RUNTIME_WRAPPER_H__ + +#if defined(__CUDA__) && defined(__clang__) + +// Include some forward declares that must come before cmath. +#include <__clang_cuda_math_forward_declares.h> + +// Define __CUDACC__ early as libstdc++ standard headers with GNU extensions +// enabled depend on it to avoid using __float128, which is unsupported in +// CUDA. +#define __CUDACC__ + +// Include some standard headers to avoid CUDA headers including them +// while some required macros (like __THROW) are in a weird state. +#include <cmath> +#include <cstdlib> +#include <stdlib.h> +#include <string.h> +#undef __CUDACC__ + +// Preserve common macros that will be changed below by us or by CUDA +// headers. +#pragma push_macro("__THROW") +#pragma push_macro("__CUDA_ARCH__") + +// WARNING: Preprocessor hacks below are based on specific details of +// CUDA-7.x headers and are not expected to work with any other +// version of CUDA headers. +#include "cuda.h" +#if !defined(CUDA_VERSION) +#error "cuda.h did not define CUDA_VERSION" +#elif CUDA_VERSION < 7000 +#error "Unsupported CUDA version!" +#endif + +#pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") +#if CUDA_VERSION >= 10000 +#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__ +#endif + +// Make largest subset of device functions available during host +// compilation. +#ifndef __CUDA_ARCH__ +#define __CUDA_ARCH__ 9999 +#endif + +#include "__clang_cuda_builtin_vars.h" + +// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above +// has taken care of builtin variables declared in the file. +#define __DEVICE_LAUNCH_PARAMETERS_H__ + +// {math,device}_functions.h only have declarations of the +// functions. We don't need them as we're going to pull in their +// definitions from .hpp files. +#define __DEVICE_FUNCTIONS_H__ +#define __MATH_FUNCTIONS_H__ +#define __COMMON_FUNCTIONS_H__ +// device_functions_decls is replaced by __clang_cuda_device_functions.h +// included below. +#define __DEVICE_FUNCTIONS_DECLS_H__ + +#undef __CUDACC__ +#if CUDA_VERSION < 9000 +#define __CUDABE__ +#else +#define __CUDACC__ +#define __CUDA_LIBDEVICE__ +#endif +// Disables definitions of device-side runtime support stubs in +// cuda_device_runtime_api.h +#include "host_defines.h" +#undef __CUDACC__ +#include "driver_types.h" +#include "host_config.h" + +// Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in +// cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the +// functional equivalent of what we need. +#pragma push_macro("nv_weak") +#define nv_weak weak +#undef __CUDABE__ +#undef __CUDA_LIBDEVICE__ +#define __CUDACC__ +#include "cuda_runtime.h" + +#pragma pop_macro("nv_weak") +#undef __CUDACC__ +#define __CUDABE__ + +// CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does +// not have at the moment. Emulate them with a builtin memcpy/memset. +#define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n) +#define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n) + +#if CUDA_VERSION < 9000 +#include "crt/device_runtime.h" +#endif +#include "crt/host_runtime.h" +// device_runtime.h defines __cxa_* macros that will conflict with +// cxxabi.h. +// FIXME: redefine these as __device__ functions. +#undef __cxa_vec_ctor +#undef __cxa_vec_cctor +#undef __cxa_vec_dtor +#undef __cxa_vec_new +#undef __cxa_vec_new2 +#undef __cxa_vec_new3 +#undef __cxa_vec_delete2 +#undef __cxa_vec_delete +#undef __cxa_vec_delete3 +#undef __cxa_pure_virtual + +// math_functions.hpp expects this host function be defined on MacOS, but it +// ends up not being there because of the games we play here. Just define it +// ourselves; it's simple enough. +#ifdef __APPLE__ +inline __host__ double __signbitd(double x) { + return std::signbit(x); +} +#endif + +// CUDA 9.1 no longer provides declarations for libdevice functions, so we need +// to provide our own. +#include <__clang_cuda_libdevice_declares.h> + +// Wrappers for many device-side standard library functions, incl. math +// functions, became compiler builtins in CUDA-9 and have been removed from the +// CUDA headers. Clang now provides its own implementation of the wrappers. +#if CUDA_VERSION >= 9000 +#include <__clang_cuda_device_functions.h> +#include <__clang_cuda_math.h> +#endif + +// __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's +// counterpart does not do it, so we need to make it empty here to keep +// following CUDA includes happy. +#undef __THROW +#define __THROW + +// CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values. +// Previous versions used to check whether they are defined or not. +// CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it +// here to detect the switch. + +#if defined(CU_DEVICE_INVALID) +#if !defined(__USE_FAST_MATH__) +#define __USE_FAST_MATH__ 0 +#endif + +#if !defined(__CUDA_PREC_DIV) +#define __CUDA_PREC_DIV 0 +#endif +#endif + +// Temporarily poison __host__ macro to ensure it's not used by any of +// the headers we're about to include. +#pragma push_macro("__host__") +#define __host__ UNEXPECTED_HOST_ATTRIBUTE + +// device_functions.hpp and math_functions*.hpp use 'static +// __forceinline__' (with no __device__) for definitions of device +// functions. Temporarily redefine __forceinline__ to include +// __device__. +#pragma push_macro("__forceinline__") +#define __forceinline__ __device__ __inline__ __attribute__((always_inline)) +#if CUDA_VERSION < 9000 +#include "device_functions.hpp" +#endif + +// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we +// get the slow-but-accurate or fast-but-inaccurate versions of functions like +// sin and exp. This is controlled in clang by -fgpu-approx-transcendentals. +// +// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs. +// slow divides), so we need to scope our define carefully here. +#pragma push_macro("__USE_FAST_MATH__") +#if defined(__CLANG_GPU_APPROX_TRANSCENDENTALS__) +#define __USE_FAST_MATH__ 1 +#endif + +#if CUDA_VERSION >= 9000 +#include "crt/math_functions.hpp" +#else +#include "math_functions.hpp" +#endif + +#pragma pop_macro("__USE_FAST_MATH__") + +#if CUDA_VERSION < 9000 +#include "math_functions_dbl_ptx3.hpp" +#endif +#pragma pop_macro("__forceinline__") + +// Pull in host-only functions that are only available when neither +// __CUDACC__ nor __CUDABE__ are defined. +#undef __MATH_FUNCTIONS_HPP__ +#undef __CUDABE__ +#if CUDA_VERSION < 9000 +#include "math_functions.hpp" +#endif +// Alas, additional overloads for these functions are hard to get to. +// Considering that we only need these overloads for a few functions, +// we can provide them here. +static inline float rsqrt(float __a) { return rsqrtf(__a); } +static inline float rcbrt(float __a) { return rcbrtf(__a); } +static inline float sinpi(float __a) { return sinpif(__a); } +static inline float cospi(float __a) { return cospif(__a); } +static inline void sincospi(float __a, float *__b, float *__c) { + return sincospif(__a, __b, __c); +} +static inline float erfcinv(float __a) { return erfcinvf(__a); } +static inline float normcdfinv(float __a) { return normcdfinvf(__a); } +static inline float normcdf(float __a) { return normcdff(__a); } +static inline float erfcx(float __a) { return erfcxf(__a); } + +#if CUDA_VERSION < 9000 +// For some reason single-argument variant is not always declared by +// CUDA headers. Alas, device_functions.hpp included below needs it. +static inline __device__ void __brkpt(int __c) { __brkpt(); } +#endif + +// Now include *.hpp with definitions of various GPU functions. Alas, +// a lot of thins get declared/defined with __host__ attribute which +// we don't want and we have to define it out. We also have to include +// {device,math}_functions.hpp again in order to extract the other +// branch of #if/else inside. +#define __host__ +#undef __CUDABE__ +#define __CUDACC__ +#if CUDA_VERSION >= 9000 +// Some atomic functions became compiler builtins in CUDA-9 , so we need their +// declarations. +#include "device_atomic_functions.h" +#endif +#undef __DEVICE_FUNCTIONS_HPP__ +#include "device_atomic_functions.hpp" +#if CUDA_VERSION >= 9000 +#include "crt/device_functions.hpp" +#include "crt/device_double_functions.hpp" +#else +#include "device_functions.hpp" +#define __CUDABE__ +#include "device_double_functions.h" +#undef __CUDABE__ +#endif +#include "sm_20_atomic_functions.hpp" +// Predicate functions used in `__builtin_assume` need to have no side effect. +// However, sm_20_intrinsics.hpp doesn't define them with neither pure nor +// const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define +// them as pure ones. +#pragma push_macro("__isGlobal") +#pragma push_macro("__isShared") +#pragma push_macro("__isConstant") +#pragma push_macro("__isLocal") +#define __isGlobal __ignored_cuda___isGlobal +#define __isShared __ignored_cuda___isShared +#define __isConstant __ignored_cuda___isConstant +#define __isLocal __ignored_cuda___isLocal +#include "sm_20_intrinsics.hpp" +#pragma pop_macro("__isGlobal") +#pragma pop_macro("__isShared") +#pragma pop_macro("__isConstant") +#pragma pop_macro("__isLocal") +#pragma push_macro("__DEVICE__") +#define __DEVICE__ static __device__ __forceinline__ __attribute__((const)) +__DEVICE__ unsigned int __isGlobal(const void *p) { + return __nvvm_isspacep_global(p); +} +__DEVICE__ unsigned int __isShared(const void *p) { + return __nvvm_isspacep_shared(p); +} +__DEVICE__ unsigned int __isConstant(const void *p) { + return __nvvm_isspacep_const(p); +} +__DEVICE__ unsigned int __isLocal(const void *p) { + return __nvvm_isspacep_local(p); +} +#pragma pop_macro("__DEVICE__") +#include "sm_32_atomic_functions.hpp" + +// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the +// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to +// define them using builtins so that the optimizer can reason about and across +// these instructions. In particular, using intrinsics for ldg gets us the +// [addr+imm] addressing mode, which, although it doesn't actually exist in the +// hardware, seems to generate faster machine code because ptxas can more easily +// reason about our code. + +#if CUDA_VERSION >= 8000 +#pragma push_macro("__CUDA_ARCH__") +#undef __CUDA_ARCH__ +#include "sm_60_atomic_functions.hpp" +#include "sm_61_intrinsics.hpp" +#pragma pop_macro("__CUDA_ARCH__") +#endif + +#undef __MATH_FUNCTIONS_HPP__ + +// math_functions.hpp defines ::signbit as a __host__ __device__ function. This +// conflicts with libstdc++'s constexpr ::signbit, so we have to rename +// math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's +// conditional on __GNUC__. :) +#pragma push_macro("signbit") +#pragma push_macro("__GNUC__") +#undef __GNUC__ +#define signbit __ignored_cuda_signbit + +// CUDA-9 omits device-side definitions of some math functions if it sees +// include guard from math.h wrapper from libstdc++. We have to undo the header +// guard temporarily to get the definitions we need. +#pragma push_macro("_GLIBCXX_MATH_H") +#pragma push_macro("_LIBCPP_VERSION") +#if CUDA_VERSION >= 9000 +#undef _GLIBCXX_MATH_H +// We also need to undo another guard that checks for libc++ 3.8+ +#ifdef _LIBCPP_VERSION +#define _LIBCPP_VERSION 3700 +#endif +#endif + +#if CUDA_VERSION >= 9000 +#include "crt/math_functions.hpp" +#else +#include "math_functions.hpp" +#endif +#pragma pop_macro("_GLIBCXX_MATH_H") +#pragma pop_macro("_LIBCPP_VERSION") +#pragma pop_macro("__GNUC__") +#pragma pop_macro("signbit") + +#pragma pop_macro("__host__") + +// __clang_cuda_texture_intrinsics.h must be included first in order to provide +// implementation for __nv_tex_surf_handler that CUDA's headers depend on. +// The implementation requires c++11 and only works with CUDA-9 or newer. +#if __cplusplus >= 201103L && CUDA_VERSION >= 9000 +// clang-format off +#include <__clang_cuda_texture_intrinsics.h> +// clang-format on +#else +#if CUDA_VERSION >= 9000 +// Provide a hint that texture support needs C++11. +template <typename T> struct __nv_tex_needs_cxx11 { + const static bool value = false; +}; +template <class T> +__host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr, + cudaTextureObject_t obj, + float x) { + _Static_assert(__nv_tex_needs_cxx11<T>::value, + "Texture support requires C++11"); +} +#else +// Textures in CUDA-8 and older are not supported by clang.There's no +// convenient way to intercept texture use in these versions, so we can't +// produce a meaningful error. The source code that attempts to use textures +// will continue to fail as it does now. +#endif // CUDA_VERSION +#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000 +#include "texture_fetch_functions.h" +#include "texture_indirect_functions.h" + +// Restore state of __CUDA_ARCH__ and __THROW we had on entry. +#pragma pop_macro("__CUDA_ARCH__") +#pragma pop_macro("__THROW") + +// Set up compiler macros expected to be seen during compilation. +#undef __CUDABE__ +#define __CUDACC__ + +extern "C" { +// Device-side CUDA system calls. +// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls +// We need these declarations and wrappers for device-side +// malloc/free/printf calls to work without relying on +// -fcuda-disable-target-call-checks option. +__device__ int vprintf(const char *, const char *); +__device__ void free(void *) __attribute((nothrow)); +__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); + +// __assertfail() used to have a `noreturn` attribute. Unfortunately that +// contributed to triggering the longstanding bug in ptxas when assert was used +// in sufficiently convoluted code. See +// https://bugs.llvm.org/show_bug.cgi?id=27738 for the details. +__device__ void __assertfail(const char *__message, const char *__file, + unsigned __line, const char *__function, + size_t __charSize); + +// In order for standard assert() macro on linux to work we need to +// provide device-side __assert_fail() +__device__ static inline void __assert_fail(const char *__message, + const char *__file, unsigned __line, + const char *__function) { + __assertfail(__message, __file, __line, __function, sizeof(char)); +} + +// Clang will convert printf into vprintf, but we still need +// device-side declaration for it. +__device__ int printf(const char *, ...); +} // extern "C" + +// We also need device-side std::malloc and std::free. +namespace std { +__device__ static inline void free(void *__ptr) { ::free(__ptr); } +__device__ static inline void *malloc(size_t __size) { + return ::malloc(__size); +} +} // namespace std + +// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to +// come after we've pulled in the definition of uint3 and dim3. + +__device__ inline __cuda_builtin_threadIdx_t::operator dim3() const { + return dim3(x, y, z); +} + +__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { + return {x, y, z}; +} + +__device__ inline __cuda_builtin_blockIdx_t::operator dim3() const { + return dim3(x, y, z); +} + +__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { + return {x, y, z}; +} + +__device__ inline __cuda_builtin_blockDim_t::operator dim3() const { + return dim3(x, y, z); +} + +__device__ inline __cuda_builtin_blockDim_t::operator uint3() const { + return {x, y, z}; +} + +__device__ inline __cuda_builtin_gridDim_t::operator dim3() const { + return dim3(x, y, z); +} + +__device__ inline __cuda_builtin_gridDim_t::operator uint3() const { + return {x, y, z}; +} + +#include <__clang_cuda_cmath.h> +#include <__clang_cuda_intrinsics.h> +#include <__clang_cuda_complex_builtins.h> + +// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host +// mode, giving them their "proper" types of dim3 and uint3. This is +// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as +// hack, force-include the header (nvcc doesn't include it by default) but +// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are +// only used here for the redeclarations of blockDim and threadIdx.) +#pragma push_macro("dim3") +#pragma push_macro("uint3") +#define dim3 __cuda_builtin_blockDim_t +#define uint3 __cuda_builtin_threadIdx_t +#include "curand_mtgp32_kernel.h" +#pragma pop_macro("dim3") +#pragma pop_macro("uint3") +#pragma pop_macro("__USE_FAST_MATH__") +#pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") + +// CUDA runtime uses this undocumented function to access kernel launch +// configuration. The declaration is in crt/device_functions.h but that file +// includes a lot of other stuff we don't want. Instead, we'll provide our own +// declaration for it here. +#if CUDA_VERSION >= 9020 +extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, + size_t sharedMem = 0, + void *stream = 0); +#endif + +#endif // __CUDA__ +#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ |