Clang Project

clang_source_code/lib/Headers/__clang_cuda_runtime_wrapper.h
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.
46#include <__clang_cuda_math_forward_declares.h>
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 > 10010
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
80#include "__clang_cuda_builtin_vars.h"
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__
149inline __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.
156#include <__clang_cuda_libdevice_declares.h>
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
162#include <__clang_cuda_device_functions.h>
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.
240static inline float rsqrt(float __a) { return rsqrtf(__a); }
241static inline float rcbrt(float __a) { return rcbrtf(__a); }
242static inline float sinpi(float __a) { return sinpif(__a); }
243static inline float cospi(float __a) { return cospif(__a); }
244static inline void sincospi(float __a, float *__b, float *__c) {
245  return sincospif(__a, __b, __c);
246}
247static inline float erfcinv(float __a) { return erfcinvf(__a); }
248static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
249static inline float normcdf(float __a) { return normcdff(__a); }
250static 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.
255static 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
348extern "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.
375namespace 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>
411#include <__clang_cuda_complex_builtins.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// CUDA runtime uses this undocumented function to access kernel launch
430// configuration. The declaration is in crt/device_functions.h but that file
431// includes a lot of other stuff we don't want. Instead, we'll provide our own
432// declaration for it here.
433#if CUDA_VERSION >= 9020
434extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
435                                                size_t sharedMem = 0,
436                                                void *stream = 0);
437#endif
438
439#endif // __CUDA__
440#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
441
CodeThinker, Enjoy your reading code