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 > 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
146// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
147// get the slow-but-accurate or fast-but-inaccurate versions of functions like
148// sin and exp.  This is controlled in clang by -fcuda-approx-transcendentals.
149//
150// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
151// slow divides), so we need to scope our define carefully here.
152#pragma push_macro("__USE_FAST_MATH__")
153#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
154#define __USE_FAST_MATH__
155#endif
156#include "math_functions.hpp"
157#pragma pop_macro("__USE_FAST_MATH__")
158
159#include "math_functions_dbl_ptx3.hpp"
160#pragma pop_macro("__forceinline__")
161
162// Pull in host-only functions that are only available when neither
163// __CUDACC__ nor __CUDABE__ are defined.
164#undef __MATH_FUNCTIONS_HPP__
165#undef __CUDABE__
166#include "math_functions.hpp"
167// Alas, additional overloads for these functions are hard to get to.
168// Considering that we only need these overloads for a few functions,
169// we can provide them here.
170static inline float rsqrt(float __a) { return rsqrtf(__a); }
171static inline float rcbrt(float __a) { return rcbrtf(__a); }
172static inline float sinpi(float __a) { return sinpif(__a); }
173static inline float cospi(float __a) { return cospif(__a); }
174static inline void sincospi(float __a, float *__b, float *__c) {
175  return sincospif(__a, __b, __c);
176}
177static inline float erfcinv(float __a) { return erfcinvf(__a); }
178static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
179static inline float normcdf(float __a) { return normcdff(__a); }
180static inline float erfcx(float __a) { return erfcxf(__a); }
181
182// For some reason single-argument variant is not always declared by
183// CUDA headers. Alas, device_functions.hpp included below needs it.
184static inline __device__ void __brkpt(int __c) { __brkpt(); }
185
186// Now include *.hpp with definitions of various GPU functions.  Alas,
187// a lot of thins get declared/defined with __host__ attribute which
188// we don't want and we have to define it out. We also have to include
189// {device,math}_functions.hpp again in order to extract the other
190// branch of #if/else inside.
191
192#define __host__
193#undef __CUDABE__
194#define __CUDACC__
195#undef __DEVICE_FUNCTIONS_HPP__
196#include "device_atomic_functions.hpp"
197#include "device_functions.hpp"
198#include "sm_20_atomic_functions.hpp"
199#include "sm_20_intrinsics.hpp"
200#include "sm_32_atomic_functions.hpp"
201
202// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h.  These define the
203// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
204// define them using builtins so that the optimizer can reason about and across
205// these instructions.  In particular, using intrinsics for ldg gets us the
206// [addr+imm] addressing mode, which, although it doesn't actually exist in the
207// hardware, seems to generate faster machine code because ptxas can more easily
208// reason about our code.
209
210#undef __MATH_FUNCTIONS_HPP__
211
212// math_functions.hpp defines ::signbit as a __host__ __device__ function.  This
213// conflicts with libstdc++'s constexpr ::signbit, so we have to rename
214// math_function.hpp's ::signbit.  It's guarded by #undef signbit, but that's
215// conditional on __GNUC__.  :)
216#pragma push_macro("signbit")
217#pragma push_macro("__GNUC__")
218#undef __GNUC__
219#define signbit __ignored_cuda_signbit
220#include "math_functions.hpp"
221#pragma pop_macro("__GNUC__")
222#pragma pop_macro("signbit")
223
224#pragma pop_macro("__host__")
225
226#include "texture_indirect_functions.h"
227
228// Restore state of __CUDA_ARCH__ and __THROW we had on entry.
229#pragma pop_macro("__CUDA_ARCH__")
230#pragma pop_macro("__THROW")
231
232// Set up compiler macros expected to be seen during compilation.
233#undef __CUDABE__
234#define __CUDACC__
235
236extern "C" {
237// Device-side CUDA system calls.
238// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
239// We need these declarations and wrappers for device-side
240// malloc/free/printf calls to work without relying on
241// -fcuda-disable-target-call-checks option.
242__device__ int vprintf(const char *, const char *);
243__device__ void free(void *) __attribute((nothrow));
244__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
245__device__ void __assertfail(const char *__message, const char *__file,
246                             unsigned __line, const char *__function,
247                             size_t __charSize) __attribute__((noreturn));
248
249// In order for standard assert() macro on linux to work we need to
250// provide device-side __assert_fail()
251__device__ static inline void __assert_fail(const char *__message,
252                                            const char *__file, unsigned __line,
253                                            const char *__function) {
254  __assertfail(__message, __file, __line, __function, sizeof(char));
255}
256
257// Clang will convert printf into vprintf, but we still need
258// device-side declaration for it.
259__device__ int printf(const char *, ...);
260} // extern "C"
261
262// We also need device-side std::malloc and std::free.
263namespace std {
264__device__ static inline void free(void *__ptr) { ::free(__ptr); }
265__device__ static inline void *malloc(size_t __size) {
266  return ::malloc(__size);
267}
268} // namespace std
269
270// Out-of-line implementations from cuda_builtin_vars.h.  These need to come
271// after we've pulled in the definition of uint3 and dim3.
272
273__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
274  uint3 ret;
275  ret.x = x;
276  ret.y = y;
277  ret.z = z;
278  return ret;
279}
280
281__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
282  uint3 ret;
283  ret.x = x;
284  ret.y = y;
285  ret.z = z;
286  return ret;
287}
288
289__device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
290  return dim3(x, y, z);
291}
292
293__device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
294  return dim3(x, y, z);
295}
296
297#include <__clang_cuda_cmath.h>
298#include <__clang_cuda_intrinsics.h>
299
300// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
301// mode, giving them their "proper" types of dim3 and uint3.  This is
302// incompatible with the types we give in cuda_builtin_vars.h.  As as hack,
303// force-include the header (nvcc doesn't include it by default) but redefine
304// dim3 and uint3 to our builtin types.  (Thankfully dim3 and uint3 are only
305// used here for the redeclarations of blockDim and threadIdx.)
306#pragma push_macro("dim3")
307#pragma push_macro("uint3")
308#define dim3 __cuda_builtin_blockDim_t
309#define uint3 __cuda_builtin_threadIdx_t
310#include "curand_mtgp32_kernel.h"
311#pragma pop_macro("dim3")
312#pragma pop_macro("uint3")
313#pragma pop_macro("__USE_FAST_MATH__")
314
315#endif // __CUDA__
316#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
317