Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 1 | /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------=== |
| 2 | * |
Logan Chien | df4f766 | 2019-09-04 16:45:23 -0700 | [diff] [blame] | 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-exception |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 6 | * |
| 7 | *===-----------------------------------------------------------------------=== |
| 8 | */ |
| 9 | |
| 10 | /* |
| 11 | * WARNING: This header is intended to be directly -include'd by |
| 12 | * the compiler and is not supposed to be included by users. |
| 13 | * |
| 14 | * CUDA headers are implemented in a way that currently makes it |
| 15 | * impossible for user code to #include directly when compiling with |
| 16 | * Clang. They present different view of CUDA-supplied functions |
| 17 | * depending on where in NVCC's compilation pipeline the headers are |
| 18 | * included. Neither of these modes provides function definitions with |
| 19 | * correct attributes, so we use preprocessor to force the headers |
| 20 | * into a form that Clang can use. |
| 21 | * |
| 22 | * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's |
| 23 | * 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 | |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 34 | // Define __CUDACC__ early as libstdc++ standard headers with GNU extensions |
| 35 | // enabled depend on it to avoid using __float128, which is unsupported in |
| 36 | // CUDA. |
| 37 | #define __CUDACC__ |
| 38 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 39 | // Include some standard headers to avoid CUDA headers including them |
| 40 | // while some required macros (like __THROW) are in a weird state. |
| 41 | #include <cmath> |
| 42 | #include <cstdlib> |
| 43 | #include <stdlib.h> |
Pirama Arumuga Nainar | ec8c89d | 2022-02-23 09:26:16 -0800 | [diff] [blame] | 44 | #include <string.h> |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 45 | #undef __CUDACC__ |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 46 | |
| 47 | // Preserve common macros that will be changed below by us or by CUDA |
| 48 | // headers. |
| 49 | #pragma push_macro("__THROW") |
| 50 | #pragma push_macro("__CUDA_ARCH__") |
| 51 | |
| 52 | // WARNING: Preprocessor hacks below are based on specific details of |
| 53 | // CUDA-7.x headers and are not expected to work with any other |
| 54 | // version of CUDA headers. |
| 55 | #include "cuda.h" |
| 56 | #if !defined(CUDA_VERSION) |
| 57 | #error "cuda.h did not define CUDA_VERSION" |
Sasha Smundak | 33d5ddd | 2020-05-04 13:37:26 -0700 | [diff] [blame] | 58 | #elif CUDA_VERSION < 7000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 59 | #error "Unsupported CUDA version!" |
| 60 | #endif |
| 61 | |
Logan Chien | b0c8402 | 2018-11-09 16:19:54 +0800 | [diff] [blame] | 62 | #pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") |
| 63 | #if CUDA_VERSION >= 10000 |
| 64 | #define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__ |
| 65 | #endif |
| 66 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 67 | // Make largest subset of device functions available during host |
Pirama Arumuga Nainar | ec8c89d | 2022-02-23 09:26:16 -0800 | [diff] [blame] | 68 | // compilation. |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 69 | #ifndef __CUDA_ARCH__ |
Pirama Arumuga Nainar | ec8c89d | 2022-02-23 09:26:16 -0800 | [diff] [blame] | 70 | #define __CUDA_ARCH__ 9999 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 71 | #endif |
| 72 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 73 | #include "__clang_cuda_builtin_vars.h" |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 74 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 75 | // No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 76 | // 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 the |
| 80 | // functions. We don't need them as we're going to pull in their |
| 81 | // definitions from .hpp files. |
| 82 | #define __DEVICE_FUNCTIONS_H__ |
| 83 | #define __MATH_FUNCTIONS_H__ |
| 84 | #define __COMMON_FUNCTIONS_H__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 85 | // device_functions_decls is replaced by __clang_cuda_device_functions.h |
| 86 | // included below. |
| 87 | #define __DEVICE_FUNCTIONS_DECLS_H__ |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 88 | |
| 89 | #undef __CUDACC__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 90 | #if CUDA_VERSION < 9000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 91 | #define __CUDABE__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 92 | #else |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 93 | #define __CUDACC__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 94 | #define __CUDA_LIBDEVICE__ |
| 95 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 96 | // Disables definitions of device-side runtime support stubs in |
| 97 | // cuda_device_runtime_api.h |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 98 | #include "host_defines.h" |
| 99 | #undef __CUDACC__ |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 100 | #include "driver_types.h" |
| 101 | #include "host_config.h" |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 102 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 103 | // Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in |
| 104 | // cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the |
| 105 | // functional equivalent of what we need. |
| 106 | #pragma push_macro("nv_weak") |
| 107 | #define nv_weak weak |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 108 | #undef __CUDABE__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 109 | #undef __CUDA_LIBDEVICE__ |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 110 | #define __CUDACC__ |
| 111 | #include "cuda_runtime.h" |
| 112 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 113 | #pragma pop_macro("nv_weak") |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 114 | #undef __CUDACC__ |
| 115 | #define __CUDABE__ |
| 116 | |
| 117 | // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does |
| 118 | // 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 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 122 | #if CUDA_VERSION < 9000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 123 | #include "crt/device_runtime.h" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 124 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 125 | #include "crt/host_runtime.h" |
| 126 | // device_runtime.h defines __cxa_* macros that will conflict with |
| 127 | // cxxabi.h. |
| 128 | // FIXME: redefine these as __device__ functions. |
| 129 | #undef __cxa_vec_ctor |
| 130 | #undef __cxa_vec_cctor |
| 131 | #undef __cxa_vec_dtor |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 132 | #undef __cxa_vec_new |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 133 | #undef __cxa_vec_new2 |
| 134 | #undef __cxa_vec_new3 |
| 135 | #undef __cxa_vec_delete2 |
| 136 | #undef __cxa_vec_delete |
| 137 | #undef __cxa_vec_delete3 |
| 138 | #undef __cxa_pure_virtual |
| 139 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 140 | // math_functions.hpp expects this host function be defined on MacOS, but it |
| 141 | // ends up not being there because of the games we play here. Just define it |
| 142 | // ourselves; it's simple enough. |
| 143 | #ifdef __APPLE__ |
| 144 | inline __host__ double __signbitd(double x) { |
| 145 | return std::signbit(x); |
| 146 | } |
| 147 | #endif |
| 148 | |
| 149 | // CUDA 9.1 no longer provides declarations for libdevice functions, so we need |
| 150 | // to provide our own. |
| 151 | #include <__clang_cuda_libdevice_declares.h> |
| 152 | |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 153 | // Wrappers for many device-side standard library functions, incl. math |
| 154 | // functions, became compiler builtins in CUDA-9 and have been removed from the |
| 155 | // CUDA headers. Clang now provides its own implementation of the wrappers. |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 156 | #if CUDA_VERSION >= 9000 |
| 157 | #include <__clang_cuda_device_functions.h> |
Sasha Smundak | 0fc590b | 2020-10-07 08:11:59 -0700 | [diff] [blame] | 158 | #include <__clang_cuda_math.h> |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 159 | #endif |
| 160 | |
| 161 | // __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's |
| 162 | // counterpart does not do it, so we need to make it empty here to keep |
| 163 | // following CUDA includes happy. |
| 164 | #undef __THROW |
| 165 | #define __THROW |
| 166 | |
| 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 it |
| 170 | // here to detect the switch. |
| 171 | |
| 172 | #if defined(CU_DEVICE_INVALID) |
| 173 | #if !defined(__USE_FAST_MATH__) |
| 174 | #define __USE_FAST_MATH__ 0 |
| 175 | #endif |
| 176 | |
| 177 | #if !defined(__CUDA_PREC_DIV) |
| 178 | #define __CUDA_PREC_DIV 0 |
| 179 | #endif |
| 180 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 181 | |
| 182 | // Temporarily poison __host__ macro to ensure it's not used by any of |
| 183 | // the headers we're about to include. |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 184 | #pragma push_macro("__host__") |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 185 | #define __host__ UNEXPECTED_HOST_ATTRIBUTE |
| 186 | |
| 187 | // device_functions.hpp and math_functions*.hpp use 'static |
| 188 | // __forceinline__' (with no __device__) for definitions of device |
| 189 | // functions. Temporarily redefine __forceinline__ to include |
| 190 | // __device__. |
| 191 | #pragma push_macro("__forceinline__") |
| 192 | #define __forceinline__ __device__ __inline__ __attribute__((always_inline)) |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 193 | #if CUDA_VERSION < 9000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 194 | #include "device_functions.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 195 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 196 | |
| 197 | // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we |
| 198 | // get the slow-but-accurate or fast-but-inaccurate versions of functions like |
| 199 | // sin and exp. This is controlled in clang by -fcuda-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_CUDA_APPROX_TRANSCENDENTALS__) |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 205 | #define __USE_FAST_MATH__ 1 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 206 | #endif |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 207 | |
| 208 | #if CUDA_VERSION >= 9000 |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 209 | #include "crt/math_functions.hpp" |
| 210 | #else |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 211 | #include "math_functions.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 212 | #endif |
| 213 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 214 | #pragma pop_macro("__USE_FAST_MATH__") |
| 215 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 216 | #if CUDA_VERSION < 9000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 217 | #include "math_functions_dbl_ptx3.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 218 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 219 | #pragma pop_macro("__forceinline__") |
| 220 | |
| 221 | // Pull in host-only functions that are only available when neither |
| 222 | // __CUDACC__ nor __CUDABE__ are defined. |
| 223 | #undef __MATH_FUNCTIONS_HPP__ |
| 224 | #undef __CUDABE__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 225 | #if CUDA_VERSION < 9000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 226 | #include "math_functions.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 227 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 228 | // 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. |
| 231 | static inline float rsqrt(float __a) { return rsqrtf(__a); } |
| 232 | static inline float rcbrt(float __a) { return rcbrtf(__a); } |
| 233 | static inline float sinpi(float __a) { return sinpif(__a); } |
| 234 | static inline float cospi(float __a) { return cospif(__a); } |
| 235 | static inline void sincospi(float __a, float *__b, float *__c) { |
| 236 | return sincospif(__a, __b, __c); |
| 237 | } |
| 238 | static inline float erfcinv(float __a) { return erfcinvf(__a); } |
| 239 | static inline float normcdfinv(float __a) { return normcdfinvf(__a); } |
| 240 | static inline float normcdf(float __a) { return normcdff(__a); } |
| 241 | static inline float erfcx(float __a) { return erfcxf(__a); } |
| 242 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 243 | #if CUDA_VERSION < 9000 |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 244 | // For some reason single-argument variant is not always declared by |
| 245 | // CUDA headers. Alas, device_functions.hpp included below needs it. |
| 246 | static inline __device__ void __brkpt(int __c) { __brkpt(); } |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 247 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 248 | |
| 249 | // Now include *.hpp with definitions of various GPU functions. Alas, |
| 250 | // a lot of thins get declared/defined with __host__ attribute which |
| 251 | // we don't want and we have to define it out. We also have to include |
| 252 | // {device,math}_functions.hpp again in order to extract the other |
| 253 | // branch of #if/else inside. |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 254 | #define __host__ |
| 255 | #undef __CUDABE__ |
| 256 | #define __CUDACC__ |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 257 | #if CUDA_VERSION >= 9000 |
| 258 | // Some atomic functions became compiler builtins in CUDA-9 , so we need their |
| 259 | // declarations. |
| 260 | #include "device_atomic_functions.h" |
| 261 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 262 | #undef __DEVICE_FUNCTIONS_HPP__ |
| 263 | #include "device_atomic_functions.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 264 | #if CUDA_VERSION >= 9000 |
| 265 | #include "crt/device_functions.hpp" |
| 266 | #include "crt/device_double_functions.hpp" |
| 267 | #else |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 268 | #include "device_functions.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 269 | #define __CUDABE__ |
| 270 | #include "device_double_functions.h" |
| 271 | #undef __CUDABE__ |
| 272 | #endif |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 273 | #include "sm_20_atomic_functions.hpp" |
Pirama Arumuga Nainar | ec8c89d | 2022-02-23 09:26:16 -0800 | [diff] [blame] | 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 nor |
| 276 | // const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define |
| 277 | // 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___isGlobal |
| 283 | #define __isShared __ignored_cuda___isShared |
| 284 | #define __isConstant __ignored_cuda___isConstant |
| 285 | #define __isLocal __ignored_cuda___isLocal |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 286 | #include "sm_20_intrinsics.hpp" |
Pirama Arumuga Nainar | ec8c89d | 2022-02-23 09:26:16 -0800 | [diff] [blame] | 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__") |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 306 | #include "sm_32_atomic_functions.hpp" |
| 307 | |
| 308 | // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the |
| 309 | // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to |
| 310 | // define them using builtins so that the optimizer can reason about and across |
| 311 | // these instructions. In particular, using intrinsics for ldg gets us the |
| 312 | // [addr+imm] addressing mode, which, although it doesn't actually exist in the |
| 313 | // hardware, seems to generate faster machine code because ptxas can more easily |
| 314 | // reason about our code. |
| 315 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 316 | #if CUDA_VERSION >= 8000 |
| 317 | #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 | #endif |
| 323 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 324 | #undef __MATH_FUNCTIONS_HPP__ |
| 325 | |
| 326 | // math_functions.hpp defines ::signbit as a __host__ __device__ function. This |
| 327 | // conflicts with libstdc++'s constexpr ::signbit, so we have to rename |
| 328 | // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's |
| 329 | // conditional on __GNUC__. :) |
| 330 | #pragma push_macro("signbit") |
| 331 | #pragma push_macro("__GNUC__") |
| 332 | #undef __GNUC__ |
| 333 | #define signbit __ignored_cuda_signbit |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 334 | |
| 335 | // CUDA-9 omits device-side definitions of some math functions if it sees |
| 336 | // include guard from math.h wrapper from libstdc++. We have to undo the header |
| 337 | // 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 >= 9000 |
| 341 | #undef _GLIBCXX_MATH_H |
| 342 | // We also need to undo another guard that checks for libc++ 3.8+ |
| 343 | #ifdef _LIBCPP_VERSION |
| 344 | #define _LIBCPP_VERSION 3700 |
| 345 | #endif |
| 346 | #endif |
| 347 | |
| 348 | #if CUDA_VERSION >= 9000 |
| 349 | #include "crt/math_functions.hpp" |
| 350 | #else |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 351 | #include "math_functions.hpp" |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 352 | #endif |
| 353 | #pragma pop_macro("_GLIBCXX_MATH_H") |
| 354 | #pragma pop_macro("_LIBCPP_VERSION") |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 355 | #pragma pop_macro("__GNUC__") |
| 356 | #pragma pop_macro("signbit") |
| 357 | |
| 358 | #pragma pop_macro("__host__") |
| 359 | |
Pirama Arumuga Nainar | ec8c89d | 2022-02-23 09:26:16 -0800 | [diff] [blame] | 360 | // __clang_cuda_texture_intrinsics.h must be included first in order to provide |
| 361 | // 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 >= 9000 |
| 364 | // clang-format off |
| 365 | #include <__clang_cuda_texture_intrinsics.h> |
| 366 | // clang-format on |
| 367 | #else |
| 368 | #if CUDA_VERSION >= 9000 |
| 369 | // Provide a hint that texture support needs C++11. |
| 370 | template <typename T> struct __nv_tex_needs_cxx11 { |
| 371 | const static bool value = false; |
| 372 | }; |
| 373 | template <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 | #else |
| 381 | // Textures in CUDA-8 and older are not supported by clang.There's no |
| 382 | // convenient way to intercept texture use in these versions, so we can't |
| 383 | // produce a meaningful error. The source code that attempts to use textures |
| 384 | // will continue to fail as it does now. |
| 385 | #endif // CUDA_VERSION |
| 386 | #endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000 |
| 387 | #include "texture_fetch_functions.h" |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 388 | #include "texture_indirect_functions.h" |
| 389 | |
| 390 | // Restore state of __CUDA_ARCH__ and __THROW we had on entry. |
| 391 | #pragma pop_macro("__CUDA_ARCH__") |
| 392 | #pragma pop_macro("__THROW") |
| 393 | |
| 394 | // Set up compiler macros expected to be seen during compilation. |
| 395 | #undef __CUDABE__ |
| 396 | #define __CUDACC__ |
| 397 | |
| 398 | extern "C" { |
| 399 | // Device-side CUDA system calls. |
| 400 | // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls |
| 401 | // We need these declarations and wrappers for device-side |
| 402 | // malloc/free/printf calls to work without relying on |
| 403 | // -fcuda-disable-target-call-checks option. |
| 404 | __device__ int vprintf(const char *, const char *); |
| 405 | __device__ void free(void *) __attribute((nothrow)); |
| 406 | __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 407 | |
| 408 | // __assertfail() used to have a `noreturn` attribute. Unfortunately that |
| 409 | // contributed to triggering the longstanding bug in ptxas when assert was used |
| 410 | // in sufficiently convoluted code. See |
| 411 | // https://bugs.llvm.org/show_bug.cgi?id=27738 for the details. |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 412 | __device__ void __assertfail(const char *__message, const char *__file, |
| 413 | unsigned __line, const char *__function, |
Pirama Arumuga Nainar | 7e1f839 | 2021-08-16 17:30:48 -0700 | [diff] [blame] | 414 | size_t __charSize); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 415 | |
| 416 | // In order for standard assert() macro on linux to work we need to |
| 417 | // provide device-side __assert_fail() |
| 418 | __device__ static inline void __assert_fail(const char *__message, |
| 419 | const char *__file, unsigned __line, |
| 420 | const char *__function) { |
| 421 | __assertfail(__message, __file, __line, __function, sizeof(char)); |
| 422 | } |
| 423 | |
| 424 | // Clang will convert printf into vprintf, but we still need |
| 425 | // device-side declaration for it. |
| 426 | __device__ int printf(const char *, ...); |
| 427 | } // extern "C" |
| 428 | |
| 429 | // We also need device-side std::malloc and std::free. |
| 430 | namespace std { |
| 431 | __device__ static inline void free(void *__ptr) { ::free(__ptr); } |
| 432 | __device__ static inline void *malloc(size_t __size) { |
| 433 | return ::malloc(__size); |
| 434 | } |
| 435 | } // namespace std |
| 436 | |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 437 | // Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to |
| 438 | // come after we've pulled in the definition of uint3 and dim3. |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 439 | |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 440 | __device__ inline __cuda_builtin_threadIdx_t::operator dim3() const { |
| 441 | return dim3(x, y, z); |
| 442 | } |
| 443 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 444 | __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 445 | return {x, y, z}; |
| 446 | } |
| 447 | |
| 448 | __device__ inline __cuda_builtin_blockIdx_t::operator dim3() const { |
| 449 | return dim3(x, y, z); |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 450 | } |
| 451 | |
| 452 | __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 453 | return {x, y, z}; |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 454 | } |
| 455 | |
| 456 | __device__ inline __cuda_builtin_blockDim_t::operator dim3() const { |
| 457 | return dim3(x, y, z); |
| 458 | } |
| 459 | |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 460 | __device__ inline __cuda_builtin_blockDim_t::operator uint3() const { |
| 461 | return {x, y, z}; |
| 462 | } |
| 463 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 464 | __device__ inline __cuda_builtin_gridDim_t::operator dim3() const { |
| 465 | return dim3(x, y, z); |
| 466 | } |
| 467 | |
Sasha Smundak | 4b1f33a | 2021-01-11 15:05:07 -0800 | [diff] [blame] | 468 | __device__ inline __cuda_builtin_gridDim_t::operator uint3() const { |
| 469 | return {x, y, z}; |
| 470 | } |
| 471 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 472 | #include <__clang_cuda_cmath.h> |
| 473 | #include <__clang_cuda_intrinsics.h> |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 474 | #include <__clang_cuda_complex_builtins.h> |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 475 | |
| 476 | // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host |
| 477 | // mode, giving them their "proper" types of dim3 and uint3. This is |
Logan Chien | 55afb0a | 2018-10-15 10:42:14 +0800 | [diff] [blame] | 478 | // incompatible with the types we give in __clang_cuda_builtin_vars.h. As as |
| 479 | // hack, force-include the header (nvcc doesn't include it by default) but |
| 480 | // redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are |
| 481 | // only used here for the redeclarations of blockDim and threadIdx.) |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 482 | #pragma push_macro("dim3") |
| 483 | #pragma push_macro("uint3") |
| 484 | #define dim3 __cuda_builtin_blockDim_t |
| 485 | #define uint3 __cuda_builtin_threadIdx_t |
| 486 | #include "curand_mtgp32_kernel.h" |
| 487 | #pragma pop_macro("dim3") |
| 488 | #pragma pop_macro("uint3") |
| 489 | #pragma pop_macro("__USE_FAST_MATH__") |
Logan Chien | b0c8402 | 2018-11-09 16:19:54 +0800 | [diff] [blame] | 490 | #pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__") |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 491 | |
Logan Chien | dbcf412 | 2019-03-21 10:50:25 +0800 | [diff] [blame] | 492 | // CUDA runtime uses this undocumented function to access kernel launch |
| 493 | // configuration. The declaration is in crt/device_functions.h but that file |
| 494 | // includes a lot of other stuff we don't want. Instead, we'll provide our own |
| 495 | // declaration for it here. |
| 496 | #if CUDA_VERSION >= 9020 |
| 497 | extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim, |
| 498 | size_t sharedMem = 0, |
| 499 | void *stream = 0); |
| 500 | #endif |
| 501 | |
Logan Chien | 2833ffb | 2018-10-09 10:03:24 +0800 | [diff] [blame] | 502 | #endif // __CUDA__ |
| 503 | #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ |