Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 1 | /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------=== |
| 2 | * |
| 3 | * Permission is hereby granted, free of charge, to any person obtaining a copy |
| 4 | * of this software and associated documentation files (the "Software"), to deal |
| 5 | * in the Software without restriction, including without limitation the rights |
| 6 | * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell |
| 7 | * copies of the Software, and to permit persons to whom the Software is |
| 8 | * furnished to do so, subject to the following conditions: |
| 9 | * |
| 10 | * The above copyright notice and this permission notice shall be included in |
| 11 | * all copies or substantial portions of the Software. |
| 12 | * |
| 13 | * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR |
| 14 | * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, |
| 15 | * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE |
| 16 | * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER |
| 17 | * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, |
| 18 | * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN |
| 19 | * THE SOFTWARE. |
| 20 | * |
| 21 | *===-----------------------------------------------------------------------=== |
| 22 | */ |
| 23 | |
| 24 | /* |
| 25 | * WARNING: This header is intended to be directly -include'd by |
| 26 | * the compiler and is not supposed to be included by users. |
| 27 | * |
| 28 | * CUDA headers are implemented in a way that currently makes it |
| 29 | * impossible for user code to #include directly when compiling with |
| 30 | * Clang. They present different view of CUDA-supplied functions |
| 31 | * depending on where in NVCC's compilation pipeline the headers are |
| 32 | * included. Neither of these modes provides function definitions with |
| 33 | * correct attributes, so we use preprocessor to force the headers |
| 34 | * into a form that Clang can use. |
| 35 | * |
| 36 | * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's |
| 37 | * this file during every CUDA compilation. |
| 38 | */ |
| 39 | |
| 40 | #ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__ |
| 41 | #define __CLANG_CUDA_RUNTIME_WRAPPER_H__ |
| 42 | |
| 43 | #if defined(__CUDA__) && defined(__clang__) |
| 44 | |
Ben Murdoch | 61f157c | 2016-09-16 13:49:30 +0100 | [diff] [blame] | 45 | // Include some forward declares that must come before cmath. |
| 46 | #include <__clang_cuda_math_forward_declares.h> |
| 47 | |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 48 | // Include some standard headers to avoid CUDA headers including them |
| 49 | // while some required macros (like __THROW) are in a weird state. |
| 50 | #include <cmath> |
| 51 | #include <cstdlib> |
| 52 | #include <stdlib.h> |
| 53 | |
| 54 | // Preserve common macros that will be changed below by us or by CUDA |
| 55 | // headers. |
| 56 | #pragma push_macro("__THROW") |
| 57 | #pragma push_macro("__CUDA_ARCH__") |
| 58 | |
| 59 | // WARNING: Preprocessor hacks below are based on specific details of |
| 60 | // CUDA-7.x headers and are not expected to work with any other |
| 61 | // version of CUDA headers. |
| 62 | #include "cuda.h" |
| 63 | #if !defined(CUDA_VERSION) |
| 64 | #error "cuda.h did not define CUDA_VERSION" |
| 65 | #elif CUDA_VERSION < 7000 || CUDA_VERSION > 7050 |
| 66 | #error "Unsupported CUDA version!" |
| 67 | #endif |
| 68 | |
| 69 | // Make largest subset of device functions available during host |
| 70 | // compilation -- SM_35 for the time being. |
| 71 | #ifndef __CUDA_ARCH__ |
| 72 | #define __CUDA_ARCH__ 350 |
| 73 | #endif |
| 74 | |
| 75 | #include "cuda_builtin_vars.h" |
| 76 | |
| 77 | // No need for device_launch_parameters.h as cuda_builtin_vars.h above |
| 78 | // has taken care of builtin variables declared in the file. |
| 79 | #define __DEVICE_LAUNCH_PARAMETERS_H__ |
| 80 | |
| 81 | // {math,device}_functions.h only have declarations of the |
| 82 | // functions. We don't need them as we're going to pull in their |
| 83 | // definitions from .hpp files. |
| 84 | #define __DEVICE_FUNCTIONS_H__ |
| 85 | #define __MATH_FUNCTIONS_H__ |
| 86 | #define __COMMON_FUNCTIONS_H__ |
| 87 | |
| 88 | #undef __CUDACC__ |
| 89 | #define __CUDABE__ |
| 90 | // Disables definitions of device-side runtime support stubs in |
| 91 | // cuda_device_runtime_api.h |
| 92 | #include "driver_types.h" |
| 93 | #include "host_config.h" |
| 94 | #include "host_defines.h" |
| 95 | |
| 96 | #undef __CUDABE__ |
| 97 | #define __CUDACC__ |
| 98 | #include "cuda_runtime.h" |
| 99 | |
| 100 | #undef __CUDACC__ |
| 101 | #define __CUDABE__ |
| 102 | |
| 103 | // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does |
| 104 | // not have at the moment. Emulate them with a builtin memcpy/memset. |
| 105 | #define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n) |
| 106 | #define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n) |
| 107 | |
| 108 | #include "crt/device_runtime.h" |
| 109 | #include "crt/host_runtime.h" |
| 110 | // device_runtime.h defines __cxa_* macros that will conflict with |
| 111 | // cxxabi.h. |
| 112 | // FIXME: redefine these as __device__ functions. |
| 113 | #undef __cxa_vec_ctor |
| 114 | #undef __cxa_vec_cctor |
| 115 | #undef __cxa_vec_dtor |
| 116 | #undef __cxa_vec_new2 |
| 117 | #undef __cxa_vec_new3 |
| 118 | #undef __cxa_vec_delete2 |
| 119 | #undef __cxa_vec_delete |
| 120 | #undef __cxa_vec_delete3 |
| 121 | #undef __cxa_pure_virtual |
| 122 | |
| 123 | // We need decls for functions in CUDA's libdevice with __device__ |
| 124 | // attribute only. Alas they come either as __host__ __device__ or |
| 125 | // with no attributes at all. To work around that, define __CUDA_RTC__ |
| 126 | // which produces HD variant and undef __host__ which gives us desided |
| 127 | // decls with __device__ attribute. |
| 128 | #pragma push_macro("__host__") |
| 129 | #define __host__ |
| 130 | #define __CUDACC_RTC__ |
| 131 | #include "device_functions_decls.h" |
| 132 | #undef __CUDACC_RTC__ |
| 133 | |
| 134 | // Temporarily poison __host__ macro to ensure it's not used by any of |
| 135 | // the headers we're about to include. |
| 136 | #define __host__ UNEXPECTED_HOST_ATTRIBUTE |
| 137 | |
| 138 | // device_functions.hpp and math_functions*.hpp use 'static |
| 139 | // __forceinline__' (with no __device__) for definitions of device |
| 140 | // functions. Temporarily redefine __forceinline__ to include |
| 141 | // __device__. |
| 142 | #pragma push_macro("__forceinline__") |
| 143 | #define __forceinline__ __device__ __inline__ __attribute__((always_inline)) |
| 144 | #include "device_functions.hpp" |
| 145 | #include "math_functions.hpp" |
| 146 | #include "math_functions_dbl_ptx3.hpp" |
| 147 | #pragma pop_macro("__forceinline__") |
| 148 | |
| 149 | // Pull in host-only functions that are only available when neither |
| 150 | // __CUDACC__ nor __CUDABE__ are defined. |
| 151 | #undef __MATH_FUNCTIONS_HPP__ |
| 152 | #undef __CUDABE__ |
| 153 | #include "math_functions.hpp" |
| 154 | // Alas, additional overloads for these functions are hard to get to. |
| 155 | // Considering that we only need these overloads for a few functions, |
| 156 | // we can provide them here. |
| 157 | static inline float rsqrt(float __a) { return rsqrtf(__a); } |
| 158 | static inline float rcbrt(float __a) { return rcbrtf(__a); } |
| 159 | static inline float sinpi(float __a) { return sinpif(__a); } |
| 160 | static inline float cospi(float __a) { return cospif(__a); } |
| 161 | static inline void sincospi(float __a, float *__b, float *__c) { |
| 162 | return sincospif(__a, __b, __c); |
| 163 | } |
| 164 | static inline float erfcinv(float __a) { return erfcinvf(__a); } |
| 165 | static inline float normcdfinv(float __a) { return normcdfinvf(__a); } |
| 166 | static inline float normcdf(float __a) { return normcdff(__a); } |
| 167 | static inline float erfcx(float __a) { return erfcxf(__a); } |
| 168 | |
| 169 | // For some reason single-argument variant is not always declared by |
| 170 | // CUDA headers. Alas, device_functions.hpp included below needs it. |
| 171 | static inline __device__ void __brkpt(int __c) { __brkpt(); } |
| 172 | |
| 173 | // Now include *.hpp with definitions of various GPU functions. Alas, |
| 174 | // a lot of thins get declared/defined with __host__ attribute which |
| 175 | // we don't want and we have to define it out. We also have to include |
| 176 | // {device,math}_functions.hpp again in order to extract the other |
| 177 | // branch of #if/else inside. |
| 178 | |
| 179 | #define __host__ |
| 180 | #undef __CUDABE__ |
| 181 | #define __CUDACC__ |
| 182 | #undef __DEVICE_FUNCTIONS_HPP__ |
| 183 | #include "device_atomic_functions.hpp" |
| 184 | #include "device_functions.hpp" |
| 185 | #include "sm_20_atomic_functions.hpp" |
| 186 | #include "sm_20_intrinsics.hpp" |
| 187 | #include "sm_32_atomic_functions.hpp" |
| 188 | // sm_30_intrinsics.h has declarations that use default argument, so |
| 189 | // we have to include it and it will in turn include .hpp |
| 190 | #include "sm_30_intrinsics.h" |
| 191 | #include "sm_32_intrinsics.hpp" |
Ben Murdoch | 61f157c | 2016-09-16 13:49:30 +0100 | [diff] [blame] | 192 | |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 193 | #undef __MATH_FUNCTIONS_HPP__ |
Ben Murdoch | 61f157c | 2016-09-16 13:49:30 +0100 | [diff] [blame] | 194 | |
| 195 | // math_functions.hpp defines ::signbit as a __host__ __device__ function. This |
| 196 | // conflicts with libstdc++'s constexpr ::signbit, so we have to rename |
| 197 | // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's |
| 198 | // conditional on __GNUC__. :) |
| 199 | #pragma push_macro("signbit") |
| 200 | #pragma push_macro("__GNUC__") |
| 201 | #undef __GNUC__ |
| 202 | #define signbit __ignored_cuda_signbit |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 203 | #include "math_functions.hpp" |
Ben Murdoch | 61f157c | 2016-09-16 13:49:30 +0100 | [diff] [blame] | 204 | #pragma pop_macro("__GNUC__") |
| 205 | #pragma pop_macro("signbit") |
| 206 | |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 207 | #pragma pop_macro("__host__") |
| 208 | |
| 209 | #include "texture_indirect_functions.h" |
| 210 | |
| 211 | // Restore state of __CUDA_ARCH__ and __THROW we had on entry. |
| 212 | #pragma pop_macro("__CUDA_ARCH__") |
| 213 | #pragma pop_macro("__THROW") |
| 214 | |
| 215 | // Set up compiler macros expected to be seen during compilation. |
| 216 | #undef __CUDABE__ |
| 217 | #define __CUDACC__ |
| 218 | |
Ben Murdoch | 097c5b2 | 2016-05-18 11:27:45 +0100 | [diff] [blame] | 219 | extern "C" { |
| 220 | // Device-side CUDA system calls. |
| 221 | // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls |
| 222 | // We need these declarations and wrappers for device-side |
| 223 | // malloc/free/printf calls to work without relying on |
| 224 | // -fcuda-disable-target-call-checks option. |
| 225 | __device__ int vprintf(const char *, const char *); |
| 226 | __device__ void free(void *) __attribute((nothrow)); |
| 227 | __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc)); |
| 228 | __device__ void __assertfail(const char *__message, const char *__file, |
| 229 | unsigned __line, const char *__function, |
| 230 | size_t __charSize) __attribute__((noreturn)); |
| 231 | |
| 232 | // In order for standard assert() macro on linux to work we need to |
| 233 | // provide device-side __assert_fail() |
| 234 | __device__ static inline void __assert_fail(const char *__message, |
| 235 | const char *__file, unsigned __line, |
| 236 | const char *__function) { |
| 237 | __assertfail(__message, __file, __line, __function, sizeof(char)); |
| 238 | } |
| 239 | |
| 240 | // Clang will convert printf into vprintf, but we still need |
| 241 | // device-side declaration for it. |
| 242 | __device__ int printf(const char *, ...); |
| 243 | } // extern "C" |
| 244 | |
| 245 | // We also need device-side std::malloc and std::free. |
| 246 | namespace std { |
| 247 | __device__ static inline void free(void *__ptr) { ::free(__ptr); } |
| 248 | __device__ static inline void *malloc(size_t __size) { |
| 249 | return ::malloc(__size); |
| 250 | } |
| 251 | } // namespace std |
| 252 | |
| 253 | // Out-of-line implementations from cuda_builtin_vars.h. These need to come |
| 254 | // after we've pulled in the definition of uint3 and dim3. |
| 255 | |
| 256 | __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const { |
| 257 | uint3 ret; |
| 258 | ret.x = x; |
| 259 | ret.y = y; |
| 260 | ret.z = z; |
| 261 | return ret; |
| 262 | } |
| 263 | |
| 264 | __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const { |
| 265 | uint3 ret; |
| 266 | ret.x = x; |
| 267 | ret.y = y; |
| 268 | ret.z = z; |
| 269 | return ret; |
| 270 | } |
| 271 | |
| 272 | __device__ inline __cuda_builtin_blockDim_t::operator dim3() const { |
| 273 | return dim3(x, y, z); |
| 274 | } |
| 275 | |
| 276 | __device__ inline __cuda_builtin_gridDim_t::operator dim3() const { |
| 277 | return dim3(x, y, z); |
| 278 | } |
| 279 | |
| 280 | #include <__clang_cuda_cmath.h> |
| 281 | |
| 282 | // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host |
| 283 | // mode, giving them their "proper" types of dim3 and uint3. This is |
| 284 | // incompatible with the types we give in cuda_builtin_vars.h. As as hack, |
| 285 | // force-include the header (nvcc doesn't include it by default) but redefine |
| 286 | // dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are only |
| 287 | // used here for the redeclarations of blockDim and threadIdx.) |
| 288 | #pragma push_macro("dim3") |
| 289 | #pragma push_macro("uint3") |
| 290 | #define dim3 __cuda_builtin_blockDim_t |
| 291 | #define uint3 __cuda_builtin_threadIdx_t |
| 292 | #include "curand_mtgp32_kernel.h" |
| 293 | #pragma pop_macro("dim3") |
| 294 | #pragma pop_macro("uint3") |
| 295 | |
| 296 | #endif // __CUDA__ |
| 297 | #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__ |