clang  8.0.0svn
__clang_cuda_runtime_wrapper.h
Go to the documentation of this file.
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 
45 // Include some forward declares that must come before cmath.
47 
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 > 10000
66 #error "Unsupported CUDA version!"
67 #endif
68 
69 #pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
70 #if CUDA_VERSION >= 10000
71 #define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
72 #endif
73 
74 // Make largest subset of device functions available during host
75 // compilation -- SM_35 for the time being.
76 #ifndef __CUDA_ARCH__
77 #define __CUDA_ARCH__ 350
78 #endif
79 
81 
82 // No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above
83 // has taken care of builtin variables declared in the file.
84 #define __DEVICE_LAUNCH_PARAMETERS_H__
85 
86 // {math,device}_functions.h only have declarations of the
87 // functions. We don't need them as we're going to pull in their
88 // definitions from .hpp files.
89 #define __DEVICE_FUNCTIONS_H__
90 #define __MATH_FUNCTIONS_H__
91 #define __COMMON_FUNCTIONS_H__
92 // device_functions_decls is replaced by __clang_cuda_device_functions.h
93 // included below.
94 #define __DEVICE_FUNCTIONS_DECLS_H__
95 
96 #undef __CUDACC__
97 #if CUDA_VERSION < 9000
98 #define __CUDABE__
99 #else
100 #define __CUDA_LIBDEVICE__
101 #endif
102 // Disables definitions of device-side runtime support stubs in
103 // cuda_device_runtime_api.h
104 #include "driver_types.h"
105 #include "host_config.h"
106 #include "host_defines.h"
107 
108 // Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in
109 // cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the
110 // functional equivalent of what we need.
111 #pragma push_macro("nv_weak")
112 #define nv_weak weak
113 #undef __CUDABE__
114 #undef __CUDA_LIBDEVICE__
115 #define __CUDACC__
116 #include "cuda_runtime.h"
117 
118 #pragma pop_macro("nv_weak")
119 #undef __CUDACC__
120 #define __CUDABE__
121 
122 // CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
123 // not have at the moment. Emulate them with a builtin memcpy/memset.
124 #define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n)
125 #define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n)
126 
127 #if CUDA_VERSION < 9000
128 #include "crt/device_runtime.h"
129 #endif
130 #include "crt/host_runtime.h"
131 // device_runtime.h defines __cxa_* macros that will conflict with
132 // cxxabi.h.
133 // FIXME: redefine these as __device__ functions.
134 #undef __cxa_vec_ctor
135 #undef __cxa_vec_cctor
136 #undef __cxa_vec_dtor
137 #undef __cxa_vec_new
138 #undef __cxa_vec_new2
139 #undef __cxa_vec_new3
140 #undef __cxa_vec_delete2
141 #undef __cxa_vec_delete
142 #undef __cxa_vec_delete3
143 #undef __cxa_pure_virtual
144 
145 // math_functions.hpp expects this host function be defined on MacOS, but it
146 // ends up not being there because of the games we play here. Just define it
147 // ourselves; it's simple enough.
148 #ifdef __APPLE__
149 inline __host__ double __signbitd(double x) {
150  return std::signbit(x);
151 }
152 #endif
153 
154 // CUDA 9.1 no longer provides declarations for libdevice functions, so we need
155 // to provide our own.
157 
158 // Wrappers for many device-side standard library functions became compiler
159 // builtins in CUDA-9 and have been removed from the CUDA headers. Clang now
160 // provides its own implementation of the wrappers.
161 #if CUDA_VERSION >= 9000
163 #endif
164 
165 // __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's
166 // counterpart does not do it, so we need to make it empty here to keep
167 // following CUDA includes happy.
168 #undef __THROW
169 #define __THROW
170 
171 // CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values.
172 // Previous versions used to check whether they are defined or not.
173 // CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it
174 // here to detect the switch.
175 
176 #if defined(CU_DEVICE_INVALID)
177 #if !defined(__USE_FAST_MATH__)
178 #define __USE_FAST_MATH__ 0
179 #endif
180 
181 #if !defined(__CUDA_PREC_DIV)
182 #define __CUDA_PREC_DIV 0
183 #endif
184 #endif
185 
186 // Temporarily poison __host__ macro to ensure it's not used by any of
187 // the headers we're about to include.
188 #pragma push_macro("__host__")
189 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
190 
191 // device_functions.hpp and math_functions*.hpp use 'static
192 // __forceinline__' (with no __device__) for definitions of device
193 // functions. Temporarily redefine __forceinline__ to include
194 // __device__.
195 #pragma push_macro("__forceinline__")
196 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
197 #if CUDA_VERSION < 9000
198 #include "device_functions.hpp"
199 #endif
200 
201 // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
202 // get the slow-but-accurate or fast-but-inaccurate versions of functions like
203 // sin and exp. This is controlled in clang by -fcuda-approx-transcendentals.
204 //
205 // device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
206 // slow divides), so we need to scope our define carefully here.
207 #pragma push_macro("__USE_FAST_MATH__")
208 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
209 #define __USE_FAST_MATH__ 1
210 #endif
211 
212 #if CUDA_VERSION >= 9000
213 // CUDA-9.2 needs host-side memcpy for some host functions in
214 // device_functions.hpp
215 #if CUDA_VERSION >= 9020
216 #include <string.h>
217 #endif
218 #include "crt/math_functions.hpp"
219 #else
220 #include "math_functions.hpp"
221 #endif
222 
223 #pragma pop_macro("__USE_FAST_MATH__")
224 
225 #if CUDA_VERSION < 9000
226 #include "math_functions_dbl_ptx3.hpp"
227 #endif
228 #pragma pop_macro("__forceinline__")
229 
230 // Pull in host-only functions that are only available when neither
231 // __CUDACC__ nor __CUDABE__ are defined.
232 #undef __MATH_FUNCTIONS_HPP__
233 #undef __CUDABE__
234 #if CUDA_VERSION < 9000
235 #include "math_functions.hpp"
236 #endif
237 // Alas, additional overloads for these functions are hard to get to.
238 // Considering that we only need these overloads for a few functions,
239 // we can provide them here.
240 static inline float rsqrt(float __a) { return rsqrtf(__a); }
241 static inline float rcbrt(float __a) { return rcbrtf(__a); }
242 static inline float sinpi(float __a) { return sinpif(__a); }
243 static inline float cospi(float __a) { return cospif(__a); }
244 static inline void sincospi(float __a, float *__b, float *__c) {
245  return sincospif(__a, __b, __c);
246 }
247 static inline float erfcinv(float __a) { return erfcinvf(__a); }
248 static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
249 static inline float normcdf(float __a) { return normcdff(__a); }
250 static inline float erfcx(float __a) { return erfcxf(__a); }
251 
252 #if CUDA_VERSION < 9000
253 // For some reason single-argument variant is not always declared by
254 // CUDA headers. Alas, device_functions.hpp included below needs it.
255 static inline __device__ void __brkpt(int __c) { __brkpt(); }
256 #endif
257 
258 // Now include *.hpp with definitions of various GPU functions. Alas,
259 // a lot of thins get declared/defined with __host__ attribute which
260 // we don't want and we have to define it out. We also have to include
261 // {device,math}_functions.hpp again in order to extract the other
262 // branch of #if/else inside.
263 #define __host__
264 #undef __CUDABE__
265 #define __CUDACC__
266 #if CUDA_VERSION >= 9000
267 // Some atomic functions became compiler builtins in CUDA-9 , so we need their
268 // declarations.
269 #include "device_atomic_functions.h"
270 #endif
271 #undef __DEVICE_FUNCTIONS_HPP__
272 #include "device_atomic_functions.hpp"
273 #if CUDA_VERSION >= 9000
274 #include "crt/device_functions.hpp"
275 #include "crt/device_double_functions.hpp"
276 #else
277 #include "device_functions.hpp"
278 #define __CUDABE__
279 #include "device_double_functions.h"
280 #undef __CUDABE__
281 #endif
282 #include "sm_20_atomic_functions.hpp"
283 #include "sm_20_intrinsics.hpp"
284 #include "sm_32_atomic_functions.hpp"
285 
286 // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the
287 // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
288 // define them using builtins so that the optimizer can reason about and across
289 // these instructions. In particular, using intrinsics for ldg gets us the
290 // [addr+imm] addressing mode, which, although it doesn't actually exist in the
291 // hardware, seems to generate faster machine code because ptxas can more easily
292 // reason about our code.
293 
294 #if CUDA_VERSION >= 8000
295 #pragma push_macro("__CUDA_ARCH__")
296 #undef __CUDA_ARCH__
297 #include "sm_60_atomic_functions.hpp"
298 #include "sm_61_intrinsics.hpp"
299 #pragma pop_macro("__CUDA_ARCH__")
300 #endif
301 
302 #undef __MATH_FUNCTIONS_HPP__
303 
304 // math_functions.hpp defines ::signbit as a __host__ __device__ function. This
305 // conflicts with libstdc++'s constexpr ::signbit, so we have to rename
306 // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's
307 // conditional on __GNUC__. :)
308 #pragma push_macro("signbit")
309 #pragma push_macro("__GNUC__")
310 #undef __GNUC__
311 #define signbit __ignored_cuda_signbit
312 
313 // CUDA-9 omits device-side definitions of some math functions if it sees
314 // include guard from math.h wrapper from libstdc++. We have to undo the header
315 // guard temporarily to get the definitions we need.
316 #pragma push_macro("_GLIBCXX_MATH_H")
317 #pragma push_macro("_LIBCPP_VERSION")
318 #if CUDA_VERSION >= 9000
319 #undef _GLIBCXX_MATH_H
320 // We also need to undo another guard that checks for libc++ 3.8+
321 #ifdef _LIBCPP_VERSION
322 #define _LIBCPP_VERSION 3700
323 #endif
324 #endif
325 
326 #if CUDA_VERSION >= 9000
327 #include "crt/math_functions.hpp"
328 #else
329 #include "math_functions.hpp"
330 #endif
331 #pragma pop_macro("_GLIBCXX_MATH_H")
332 #pragma pop_macro("_LIBCPP_VERSION")
333 #pragma pop_macro("__GNUC__")
334 #pragma pop_macro("signbit")
335 
336 #pragma pop_macro("__host__")
337 
338 #include "texture_indirect_functions.h"
339 
340 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
341 #pragma pop_macro("__CUDA_ARCH__")
342 #pragma pop_macro("__THROW")
343 
344 // Set up compiler macros expected to be seen during compilation.
345 #undef __CUDABE__
346 #define __CUDACC__
347 
348 extern "C" {
349 // Device-side CUDA system calls.
350 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
351 // We need these declarations and wrappers for device-side
352 // malloc/free/printf calls to work without relying on
353 // -fcuda-disable-target-call-checks option.
354 __device__ int vprintf(const char *, const char *);
355 __device__ void free(void *) __attribute((nothrow));
356 __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
357 __device__ void __assertfail(const char *__message, const char *__file,
358  unsigned __line, const char *__function,
359  size_t __charSize) __attribute__((noreturn));
360 
361 // In order for standard assert() macro on linux to work we need to
362 // provide device-side __assert_fail()
363 __device__ static inline void __assert_fail(const char *__message,
364  const char *__file, unsigned __line,
365  const char *__function) {
366  __assertfail(__message, __file, __line, __function, sizeof(char));
367 }
368 
369 // Clang will convert printf into vprintf, but we still need
370 // device-side declaration for it.
371 __device__ int printf(const char *, ...);
372 } // extern "C"
373 
374 // We also need device-side std::malloc and std::free.
375 namespace std {
376 __device__ static inline void free(void *__ptr) { ::free(__ptr); }
377 __device__ static inline void *malloc(size_t __size) {
378  return ::malloc(__size);
379 }
380 } // namespace std
381 
382 // Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to
383 // come after we've pulled in the definition of uint3 and dim3.
384 
385 __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
386  uint3 ret;
387  ret.x = x;
388  ret.y = y;
389  ret.z = z;
390  return ret;
391 }
392 
393 __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
394  uint3 ret;
395  ret.x = x;
396  ret.y = y;
397  ret.z = z;
398  return ret;
399 }
400 
401 __device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
402  return dim3(x, y, z);
403 }
404 
405 __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
406  return dim3(x, y, z);
407 }
408 
409 #include <__clang_cuda_cmath.h>
410 #include <__clang_cuda_intrinsics.h>
412 
413 // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
414 // mode, giving them their "proper" types of dim3 and uint3. This is
415 // incompatible with the types we give in __clang_cuda_builtin_vars.h. As as
416 // hack, force-include the header (nvcc doesn't include it by default) but
417 // redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are
418 // only used here for the redeclarations of blockDim and threadIdx.)
419 #pragma push_macro("dim3")
420 #pragma push_macro("uint3")
421 #define dim3 __cuda_builtin_blockDim_t
422 #define uint3 __cuda_builtin_threadIdx_t
423 #include "curand_mtgp32_kernel.h"
424 #pragma pop_macro("dim3")
425 #pragma pop_macro("uint3")
426 #pragma pop_macro("__USE_FAST_MATH__")
427 #pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
428 
429 #endif // __CUDA__
430 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
__DEVICE__ bool signbit(float __x)
Test for sign bit.
__DEVICE__ float normcdff(float __a)
__DEVICE__ void sincospif(float __a, float *__sptr, float *__cptr)
__DEVICE__ double cospi(double __a)
__DEVICE__ double sinpi(double __a)
__DEVICE__ double normcdfinv(double __a)
Definition: Format.h:2075
__DEVICE__ float rsqrtf(float __a)
__DEVICE__ double rsqrt(double __a)
static __inline__ void int __a
Definition: emmintrin.h:4204
__DEVICE__ float cospif(float __a)
__DEVICE__ float erfcinvf(float __a)
char __v64qi __attribute__((__vector_size__(64)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
Definition: avx512fintrin.h:30
__DEVICE__ float rcbrtf(float __a)
__DEVICE__ void sincospi(double __a, double *__sptr, double *__cptr)
__DEVICE__ double erfcinv(double __a)
static __inline__ vector float vector float __b
Definition: altivec.h:534
#define noreturn
Definition: stdnoreturn.h:27
__DEVICE__ void __brkpt()
__DEVICE__ int __signbitd(double __a)
int printf(__constant const char *st,...) __attribute__((format(printf
__DEVICE__ double erfcx(double __a)
__DEVICE__ float sinpif(float __a)
__DEVICE__ float normcdfinvf(float __a)
__DEVICE__ double normcdf(double __a)
__DEVICE__ double rcbrt(double __a)
__DEVICE__ float erfcxf(float __a)
static __inline__ vector float vector float vector float __c
Definition: altivec.h:4199