1 /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
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:
10 * The above copyright notice and this permission notice shall be included in
11 * all copies or substantial portions of the Software.
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
21 *===-----------------------------------------------------------------------===
25 * WARNING: This header is intended to be directly -include'd by
26 * the compiler and is not supposed to be included by users.
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.
36 * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's
37 * this file during every CUDA compilation.
40 #ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__
41 #define __CLANG_CUDA_RUNTIME_WRAPPER_H__
43 #if defined(__CUDA__) && defined(__clang__)
45 // Include some forward declares that must come before cmath.
46 #include <__clang_cuda_math_forward_declares.h>
48 // Include some standard headers to avoid CUDA headers including them
49 // while some required macros (like __THROW) are in a weird state.
54 // Preserve common macros that will be changed below by us or by CUDA
56 #pragma push_macro("__THROW")
57 #pragma push_macro("__CUDA_ARCH__")
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.
63 #if !defined(CUDA_VERSION)
64 #error "cuda.h did not define CUDA_VERSION"
65 #elif CUDA_VERSION < 7000 || CUDA_VERSION > 9020
66 #error "Unsupported CUDA version!"
69 // Make largest subset of device functions available during host
70 // compilation -- SM_35 for the time being.
72 #define __CUDA_ARCH__ 350
75 #include "__clang_cuda_builtin_vars.h"
77 // No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above
78 // has taken care of builtin variables declared in the file.
79 #define __DEVICE_LAUNCH_PARAMETERS_H__
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 // device_functions_decls is replaced by __clang_cuda_device_functions.h
89 #define __DEVICE_FUNCTIONS_DECLS_H__
92 #if CUDA_VERSION < 9000
95 #define __CUDA_LIBDEVICE__
97 // Disables definitions of device-side runtime support stubs in
98 // cuda_device_runtime_api.h
99 #include "driver_types.h"
100 #include "host_config.h"
101 #include "host_defines.h"
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")
109 #undef __CUDA_LIBDEVICE__
111 #include "cuda_runtime.h"
113 #pragma pop_macro("nv_weak")
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)
122 #if CUDA_VERSION < 9000
123 #include "crt/device_runtime.h"
125 #include "crt/host_runtime.h"
126 // device_runtime.h defines __cxa_* macros that will conflict with
128 // FIXME: redefine these as __device__ functions.
129 #undef __cxa_vec_ctor
130 #undef __cxa_vec_cctor
131 #undef __cxa_vec_dtor
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
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.
144 inline __host__ double __signbitd(double x) {
145 return std::signbit(x);
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>
153 // Wrappers for many device-side standard library functions became compiler
154 // builtins in CUDA-9 and have been removed from the CUDA headers. Clang now
155 // provides its own implementation of the wrappers.
156 #if CUDA_VERSION >= 9000
157 #include <__clang_cuda_device_functions.h>
160 // __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's
161 // counterpart does not do it, so we need to make it empty here to keep
162 // following CUDA includes happy.
166 // CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values.
167 // Previous versions used to check whether they are defined or not.
168 // CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it
169 // here to detect the switch.
171 #if defined(CU_DEVICE_INVALID)
172 #if !defined(__USE_FAST_MATH__)
173 #define __USE_FAST_MATH__ 0
176 #if !defined(__CUDA_PREC_DIV)
177 #define __CUDA_PREC_DIV 0
181 // Temporarily poison __host__ macro to ensure it's not used by any of
182 // the headers we're about to include.
183 #pragma push_macro("__host__")
184 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
186 // device_functions.hpp and math_functions*.hpp use 'static
187 // __forceinline__' (with no __device__) for definitions of device
188 // functions. Temporarily redefine __forceinline__ to include
190 #pragma push_macro("__forceinline__")
191 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
192 #if CUDA_VERSION < 9000
193 #include "device_functions.hpp"
196 // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
197 // get the slow-but-accurate or fast-but-inaccurate versions of functions like
198 // sin and exp. This is controlled in clang by -fcuda-approx-transcendentals.
200 // device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
201 // slow divides), so we need to scope our define carefully here.
202 #pragma push_macro("__USE_FAST_MATH__")
203 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
204 #define __USE_FAST_MATH__ 1
207 #if CUDA_VERSION >= 9000
208 // CUDA-9.2 needs host-side memcpy for some host functions in
209 // device_functions.hpp
210 #if CUDA_VERSION >= 9020
213 #include "crt/math_functions.hpp"
215 #include "math_functions.hpp"
218 #pragma pop_macro("__USE_FAST_MATH__")
220 #if CUDA_VERSION < 9000
221 #include "math_functions_dbl_ptx3.hpp"
223 #pragma pop_macro("__forceinline__")
225 // Pull in host-only functions that are only available when neither
226 // __CUDACC__ nor __CUDABE__ are defined.
227 #undef __MATH_FUNCTIONS_HPP__
229 #if CUDA_VERSION < 9000
230 #include "math_functions.hpp"
232 // Alas, additional overloads for these functions are hard to get to.
233 // Considering that we only need these overloads for a few functions,
234 // we can provide them here.
235 static inline float rsqrt(float __a) { return rsqrtf(__a); }
236 static inline float rcbrt(float __a) { return rcbrtf(__a); }
237 static inline float sinpi(float __a) { return sinpif(__a); }
238 static inline float cospi(float __a) { return cospif(__a); }
239 static inline void sincospi(float __a, float *__b, float *__c) {
240 return sincospif(__a, __b, __c);
242 static inline float erfcinv(float __a) { return erfcinvf(__a); }
243 static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
244 static inline float normcdf(float __a) { return normcdff(__a); }
245 static inline float erfcx(float __a) { return erfcxf(__a); }
247 #if CUDA_VERSION < 9000
248 // For some reason single-argument variant is not always declared by
249 // CUDA headers. Alas, device_functions.hpp included below needs it.
250 static inline __device__ void __brkpt(int __c) { __brkpt(); }
253 // Now include *.hpp with definitions of various GPU functions. Alas,
254 // a lot of thins get declared/defined with __host__ attribute which
255 // we don't want and we have to define it out. We also have to include
256 // {device,math}_functions.hpp again in order to extract the other
257 // branch of #if/else inside.
261 #if CUDA_VERSION >= 9000
262 // Some atomic functions became compiler builtins in CUDA-9 , so we need their
264 #include "device_atomic_functions.h"
266 #undef __DEVICE_FUNCTIONS_HPP__
267 #include "device_atomic_functions.hpp"
268 #if CUDA_VERSION >= 9000
269 #include "crt/device_functions.hpp"
270 #include "crt/device_double_functions.hpp"
272 #include "device_functions.hpp"
274 #include "device_double_functions.h"
277 #include "sm_20_atomic_functions.hpp"
278 #include "sm_20_intrinsics.hpp"
279 #include "sm_32_atomic_functions.hpp"
281 // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the
282 // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
283 // define them using builtins so that the optimizer can reason about and across
284 // these instructions. In particular, using intrinsics for ldg gets us the
285 // [addr+imm] addressing mode, which, although it doesn't actually exist in the
286 // hardware, seems to generate faster machine code because ptxas can more easily
287 // reason about our code.
289 #if CUDA_VERSION >= 8000
290 #pragma push_macro("__CUDA_ARCH__")
292 #include "sm_60_atomic_functions.hpp"
293 #include "sm_61_intrinsics.hpp"
294 #pragma pop_macro("__CUDA_ARCH__")
297 #undef __MATH_FUNCTIONS_HPP__
299 // math_functions.hpp defines ::signbit as a __host__ __device__ function. This
300 // conflicts with libstdc++'s constexpr ::signbit, so we have to rename
301 // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's
302 // conditional on __GNUC__. :)
303 #pragma push_macro("signbit")
304 #pragma push_macro("__GNUC__")
306 #define signbit __ignored_cuda_signbit
308 // CUDA-9 omits device-side definitions of some math functions if it sees
309 // include guard from math.h wrapper from libstdc++. We have to undo the header
310 // guard temporarily to get the definitions we need.
311 #pragma push_macro("_GLIBCXX_MATH_H")
312 #pragma push_macro("_LIBCPP_VERSION")
313 #if CUDA_VERSION >= 9000
314 #undef _GLIBCXX_MATH_H
315 // We also need to undo another guard that checks for libc++ 3.8+
316 #ifdef _LIBCPP_VERSION
317 #define _LIBCPP_VERSION 3700
321 #if CUDA_VERSION >= 9000
322 #include "crt/math_functions.hpp"
324 #include "math_functions.hpp"
326 #pragma pop_macro("_GLIBCXX_MATH_H")
327 #pragma pop_macro("_LIBCPP_VERSION")
328 #pragma pop_macro("__GNUC__")
329 #pragma pop_macro("signbit")
331 #pragma pop_macro("__host__")
333 #include "texture_indirect_functions.h"
335 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
336 #pragma pop_macro("__CUDA_ARCH__")
337 #pragma pop_macro("__THROW")
339 // Set up compiler macros expected to be seen during compilation.
344 // Device-side CUDA system calls.
345 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
346 // We need these declarations and wrappers for device-side
347 // malloc/free/printf calls to work without relying on
348 // -fcuda-disable-target-call-checks option.
349 __device__ int vprintf(const char *, const char *);
350 __device__ void free(void *) __attribute((nothrow));
351 __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
352 __device__ void __assertfail(const char *__message, const char *__file,
353 unsigned __line, const char *__function,
354 size_t __charSize) __attribute__((noreturn));
356 // In order for standard assert() macro on linux to work we need to
357 // provide device-side __assert_fail()
358 __device__ static inline void __assert_fail(const char *__message,
359 const char *__file, unsigned __line,
360 const char *__function) {
361 __assertfail(__message, __file, __line, __function, sizeof(char));
364 // Clang will convert printf into vprintf, but we still need
365 // device-side declaration for it.
366 __device__ int printf(const char *, ...);
369 // We also need device-side std::malloc and std::free.
371 __device__ static inline void free(void *__ptr) { ::free(__ptr); }
372 __device__ static inline void *malloc(size_t __size) {
373 return ::malloc(__size);
377 // Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to
378 // come after we've pulled in the definition of uint3 and dim3.
380 __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
388 __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
396 __device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
397 return dim3(x, y, z);
400 __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
401 return dim3(x, y, z);
404 #include <__clang_cuda_cmath.h>
405 #include <__clang_cuda_intrinsics.h>
406 #include <__clang_cuda_complex_builtins.h>
408 // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
409 // mode, giving them their "proper" types of dim3 and uint3. This is
410 // incompatible with the types we give in __clang_cuda_builtin_vars.h. As as
411 // hack, force-include the header (nvcc doesn't include it by default) but
412 // redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are
413 // only used here for the redeclarations of blockDim and threadIdx.)
414 #pragma push_macro("dim3")
415 #pragma push_macro("uint3")
416 #define dim3 __cuda_builtin_blockDim_t
417 #define uint3 __cuda_builtin_threadIdx_t
418 #include "curand_mtgp32_kernel.h"
419 #pragma pop_macro("dim3")
420 #pragma pop_macro("uint3")
421 #pragma pop_macro("__USE_FAST_MATH__")
424 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__