clang  14.0.0git
__clang_cuda_runtime_wrapper.h
Go to the documentation of this file.
1 /*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
2  *
3  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4  * See https://llvm.org/LICENSE.txt for license information.
5  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6  *
7  *===-----------------------------------------------------------------------===
8  */
9 
10 /*
11  * WARNING: This header is intended to be directly -include'd by
12  * the compiler and is not supposed to be included by users.
13  *
14  * CUDA headers are implemented in a way that currently makes it
15  * impossible for user code to #include directly when compiling with
16  * Clang. They present different view of CUDA-supplied functions
17  * depending on where in NVCC's compilation pipeline the headers are
18  * included. Neither of these modes provides function definitions with
19  * correct attributes, so we use preprocessor to force the headers
20  * into a form that Clang can use.
21  *
22  * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's
23  * this file during every CUDA compilation.
24  */
25 
26 #ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__
27 #define __CLANG_CUDA_RUNTIME_WRAPPER_H__
28 
29 #if defined(__CUDA__) && defined(__clang__)
30 
31 // Include some forward declares that must come before cmath.
33 
34 // Define __CUDACC__ early as libstdc++ standard headers with GNU extensions
35 // enabled depend on it to avoid using __float128, which is unsupported in
36 // CUDA.
37 #define __CUDACC__
38 
39 // Include some standard headers to avoid CUDA headers including them
40 // while some required macros (like __THROW) are in a weird state.
41 #include <cmath>
42 #include <cstdlib>
43 #include <stdlib.h>
44 #include <string.h>
45 #undef __CUDACC__
46 
47 // Preserve common macros that will be changed below by us or by CUDA
48 // headers.
49 #pragma push_macro("__THROW")
50 #pragma push_macro("__CUDA_ARCH__")
51 
52 // WARNING: Preprocessor hacks below are based on specific details of
53 // CUDA-7.x headers and are not expected to work with any other
54 // version of CUDA headers.
55 #include "cuda.h"
56 #if !defined(CUDA_VERSION)
57 #error "cuda.h did not define CUDA_VERSION"
58 #elif CUDA_VERSION < 7000
59 #error "Unsupported CUDA version!"
60 #endif
61 
62 #pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
63 #if CUDA_VERSION >= 10000
64 #define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
65 #endif
66 
67 // Make largest subset of device functions available during host
68 // compilation.
69 #ifndef __CUDA_ARCH__
70 #define __CUDA_ARCH__ 9999
71 #endif
72 
74 
75 // No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above
76 // has taken care of builtin variables declared in the file.
77 #define __DEVICE_LAUNCH_PARAMETERS_H__
78 
79 // {math,device}_functions.h only have declarations of the
80 // functions. We don't need them as we're going to pull in their
81 // definitions from .hpp files.
82 #define __DEVICE_FUNCTIONS_H__
83 #define __MATH_FUNCTIONS_H__
84 #define __COMMON_FUNCTIONS_H__
85 // device_functions_decls is replaced by __clang_cuda_device_functions.h
86 // included below.
87 #define __DEVICE_FUNCTIONS_DECLS_H__
88 
89 #undef __CUDACC__
90 #if CUDA_VERSION < 9000
91 #define __CUDABE__
92 #else
93 #define __CUDACC__
94 #define __CUDA_LIBDEVICE__
95 #endif
96 // Disables definitions of device-side runtime support stubs in
97 // cuda_device_runtime_api.h
98 #include "host_defines.h"
99 #undef __CUDACC__
100 #include "driver_types.h"
101 #include "host_config.h"
102 
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")
107 #define nv_weak weak
108 #undef __CUDABE__
109 #undef __CUDA_LIBDEVICE__
110 #define __CUDACC__
111 #include "cuda_runtime.h"
112 
113 #pragma pop_macro("nv_weak")
114 #undef __CUDACC__
115 #define __CUDABE__
116 
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)
121 
122 #if CUDA_VERSION < 9000
123 #include "crt/device_runtime.h"
124 #endif
125 #include "crt/host_runtime.h"
126 // device_runtime.h defines __cxa_* macros that will conflict with
127 // cxxabi.h.
128 // FIXME: redefine these as __device__ functions.
129 #undef __cxa_vec_ctor
130 #undef __cxa_vec_cctor
131 #undef __cxa_vec_dtor
132 #undef __cxa_vec_new
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
139 
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.
143 #ifdef __APPLE__
144 inline __host__ double __signbitd(double x) {
145  return std::signbit(x);
146 }
147 #endif
148 
149 // CUDA 9.1 no longer provides declarations for libdevice functions, so we need
150 // to provide our own.
152 
153 // Wrappers for many device-side standard library functions, incl. math
154 // functions, became compiler builtins in CUDA-9 and have been removed from the
155 // CUDA headers. Clang now provides its own implementation of the wrappers.
156 #if CUDA_VERSION >= 9000
158 #include <__clang_cuda_math.h>
159 #endif
160 
161 // __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's
162 // counterpart does not do it, so we need to make it empty here to keep
163 // following CUDA includes happy.
164 #undef __THROW
165 #define __THROW
166 
167 // CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values.
168 // Previous versions used to check whether they are defined or not.
169 // CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it
170 // here to detect the switch.
171 
172 #if defined(CU_DEVICE_INVALID)
173 #if !defined(__USE_FAST_MATH__)
174 #define __USE_FAST_MATH__ 0
175 #endif
176 
177 #if !defined(__CUDA_PREC_DIV)
178 #define __CUDA_PREC_DIV 0
179 #endif
180 #endif
181 
182 // Temporarily poison __host__ macro to ensure it's not used by any of
183 // the headers we're about to include.
184 #pragma push_macro("__host__")
185 #define __host__ UNEXPECTED_HOST_ATTRIBUTE
186 
187 // device_functions.hpp and math_functions*.hpp use 'static
188 // __forceinline__' (with no __device__) for definitions of device
189 // functions. Temporarily redefine __forceinline__ to include
190 // __device__.
191 #pragma push_macro("__forceinline__")
192 #define __forceinline__ __device__ __inline__ __attribute__((always_inline))
193 #if CUDA_VERSION < 9000
194 #include "device_functions.hpp"
195 #endif
196 
197 // math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
198 // get the slow-but-accurate or fast-but-inaccurate versions of functions like
199 // sin and exp. This is controlled in clang by -fcuda-approx-transcendentals.
200 //
201 // device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
202 // slow divides), so we need to scope our define carefully here.
203 #pragma push_macro("__USE_FAST_MATH__")
204 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
205 #define __USE_FAST_MATH__ 1
206 #endif
207 
208 #if CUDA_VERSION >= 9000
209 #include "crt/math_functions.hpp"
210 #else
211 #include "math_functions.hpp"
212 #endif
213 
214 #pragma pop_macro("__USE_FAST_MATH__")
215 
216 #if CUDA_VERSION < 9000
217 #include "math_functions_dbl_ptx3.hpp"
218 #endif
219 #pragma pop_macro("__forceinline__")
220 
221 // Pull in host-only functions that are only available when neither
222 // __CUDACC__ nor __CUDABE__ are defined.
223 #undef __MATH_FUNCTIONS_HPP__
224 #undef __CUDABE__
225 #if CUDA_VERSION < 9000
226 #include "math_functions.hpp"
227 #endif
228 // Alas, additional overloads for these functions are hard to get to.
229 // Considering that we only need these overloads for a few functions,
230 // we can provide them here.
231 static inline float rsqrt(float __a) { return rsqrtf(__a); }
232 static inline float rcbrt(float __a) { return rcbrtf(__a); }
233 static inline float sinpi(float __a) { return sinpif(__a); }
234 static inline float cospi(float __a) { return cospif(__a); }
235 static inline void sincospi(float __a, float *__b, float *__c) {
236  return sincospif(__a, __b, __c);
237 }
238 static inline float erfcinv(float __a) { return erfcinvf(__a); }
239 static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
240 static inline float normcdf(float __a) { return normcdff(__a); }
241 static inline float erfcx(float __a) { return erfcxf(__a); }
242 
243 #if CUDA_VERSION < 9000
244 // For some reason single-argument variant is not always declared by
245 // CUDA headers. Alas, device_functions.hpp included below needs it.
246 static inline __device__ void __brkpt(int __c) { __brkpt(); }
247 #endif
248 
249 // Now include *.hpp with definitions of various GPU functions. Alas,
250 // a lot of thins get declared/defined with __host__ attribute which
251 // we don't want and we have to define it out. We also have to include
252 // {device,math}_functions.hpp again in order to extract the other
253 // branch of #if/else inside.
254 #define __host__
255 #undef __CUDABE__
256 #define __CUDACC__
257 #if CUDA_VERSION >= 9000
258 // Some atomic functions became compiler builtins in CUDA-9 , so we need their
259 // declarations.
260 #include "device_atomic_functions.h"
261 #endif
262 #undef __DEVICE_FUNCTIONS_HPP__
263 #include "device_atomic_functions.hpp"
264 #if CUDA_VERSION >= 9000
265 #include "crt/device_functions.hpp"
266 #include "crt/device_double_functions.hpp"
267 #else
268 #include "device_functions.hpp"
269 #define __CUDABE__
270 #include "device_double_functions.h"
271 #undef __CUDABE__
272 #endif
273 #include "sm_20_atomic_functions.hpp"
274 // Predicate functions used in `__builtin_assume` need to have no side effect.
275 // However, sm_20_intrinsics.hpp doesn't define them with neither pure nor
276 // const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define
277 // them as pure ones.
278 #pragma push_macro("__isGlobal")
279 #pragma push_macro("__isShared")
280 #pragma push_macro("__isConstant")
281 #pragma push_macro("__isLocal")
282 #define __isGlobal __ignored_cuda___isGlobal
283 #define __isShared __ignored_cuda___isShared
284 #define __isConstant __ignored_cuda___isConstant
285 #define __isLocal __ignored_cuda___isLocal
286 #include "sm_20_intrinsics.hpp"
287 #pragma pop_macro("__isGlobal")
288 #pragma pop_macro("__isShared")
289 #pragma pop_macro("__isConstant")
290 #pragma pop_macro("__isLocal")
291 #pragma push_macro("__DEVICE__")
292 #define __DEVICE__ static __device__ __forceinline__ __attribute__((const))
293 __DEVICE__ unsigned int __isGlobal(const void *p) {
294  return __nvvm_isspacep_global(p);
295 }
296 __DEVICE__ unsigned int __isShared(const void *p) {
297  return __nvvm_isspacep_shared(p);
298 }
299 __DEVICE__ unsigned int __isConstant(const void *p) {
300  return __nvvm_isspacep_const(p);
301 }
302 __DEVICE__ unsigned int __isLocal(const void *p) {
303  return __nvvm_isspacep_local(p);
304 }
305 #pragma pop_macro("__DEVICE__")
306 #include "sm_32_atomic_functions.hpp"
307 
308 // Don't include sm_30_intrinsics.h and sm_32_intrinsics.h. These define the
309 // __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
310 // define them using builtins so that the optimizer can reason about and across
311 // these instructions. In particular, using intrinsics for ldg gets us the
312 // [addr+imm] addressing mode, which, although it doesn't actually exist in the
313 // hardware, seems to generate faster machine code because ptxas can more easily
314 // reason about our code.
315 
316 #if CUDA_VERSION >= 8000
317 #pragma push_macro("__CUDA_ARCH__")
318 #undef __CUDA_ARCH__
319 #include "sm_60_atomic_functions.hpp"
320 #include "sm_61_intrinsics.hpp"
321 #pragma pop_macro("__CUDA_ARCH__")
322 #endif
323 
324 #undef __MATH_FUNCTIONS_HPP__
325 
326 // math_functions.hpp defines ::signbit as a __host__ __device__ function. This
327 // conflicts with libstdc++'s constexpr ::signbit, so we have to rename
328 // math_function.hpp's ::signbit. It's guarded by #undef signbit, but that's
329 // conditional on __GNUC__. :)
330 #pragma push_macro("signbit")
331 #pragma push_macro("__GNUC__")
332 #undef __GNUC__
333 #define signbit __ignored_cuda_signbit
334 
335 // CUDA-9 omits device-side definitions of some math functions if it sees
336 // include guard from math.h wrapper from libstdc++. We have to undo the header
337 // guard temporarily to get the definitions we need.
338 #pragma push_macro("_GLIBCXX_MATH_H")
339 #pragma push_macro("_LIBCPP_VERSION")
340 #if CUDA_VERSION >= 9000
341 #undef _GLIBCXX_MATH_H
342 // We also need to undo another guard that checks for libc++ 3.8+
343 #ifdef _LIBCPP_VERSION
344 #define _LIBCPP_VERSION 3700
345 #endif
346 #endif
347 
348 #if CUDA_VERSION >= 9000
349 #include "crt/math_functions.hpp"
350 #else
351 #include "math_functions.hpp"
352 #endif
353 #pragma pop_macro("_GLIBCXX_MATH_H")
354 #pragma pop_macro("_LIBCPP_VERSION")
355 #pragma pop_macro("__GNUC__")
356 #pragma pop_macro("signbit")
357 
358 #pragma pop_macro("__host__")
359 
360 // __clang_cuda_texture_intrinsics.h must be included first in order to provide
361 // implementation for __nv_tex_surf_handler that CUDA's headers depend on.
362 // The implementation requires c++11 and only works with CUDA-9 or newer.
363 #if __cplusplus >= 201103L && CUDA_VERSION >= 9000
364 // clang-format off
366 // clang-format on
367 #else
368 #if CUDA_VERSION >= 9000
369 // Provide a hint that texture support needs C++11.
370 template <typename T> struct __nv_tex_needs_cxx11 {
371  const static bool value = false;
372 };
373 template <class T>
374 __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
375  cudaTextureObject_t obj,
376  float x) {
377  _Static_assert(__nv_tex_needs_cxx11<T>::value,
378  "Texture support requires C++11");
379 }
380 #else
381 // Textures in CUDA-8 and older are not supported by clang.There's no
382 // convenient way to intercept texture use in these versions, so we can't
383 // produce a meaningful error. The source code that attempts to use textures
384 // will continue to fail as it does now.
385 #endif // CUDA_VERSION
386 #endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
387 #include "texture_fetch_functions.h"
388 #include "texture_indirect_functions.h"
389 
390 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
391 #pragma pop_macro("__CUDA_ARCH__")
392 #pragma pop_macro("__THROW")
393 
394 // Set up compiler macros expected to be seen during compilation.
395 #undef __CUDABE__
396 #define __CUDACC__
397 
398 extern "C" {
399 // Device-side CUDA system calls.
400 // http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
401 // We need these declarations and wrappers for device-side
402 // malloc/free/printf calls to work without relying on
403 // -fcuda-disable-target-call-checks option.
404 __device__ int vprintf(const char *, const char *);
405 __device__ void free(void *) __attribute((nothrow));
406 __device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));
407 
408 // __assertfail() used to have a `noreturn` attribute. Unfortunately that
409 // contributed to triggering the longstanding bug in ptxas when assert was used
410 // in sufficiently convoluted code. See
411 // https://bugs.llvm.org/show_bug.cgi?id=27738 for the details.
412 __device__ void __assertfail(const char *__message, const char *__file,
413  unsigned __line, const char *__function,
414  size_t __charSize);
415 
416 // In order for standard assert() macro on linux to work we need to
417 // provide device-side __assert_fail()
418 __device__ static inline void __assert_fail(const char *__message,
419  const char *__file, unsigned __line,
420  const char *__function) {
421  __assertfail(__message, __file, __line, __function, sizeof(char));
422 }
423 
424 // Clang will convert printf into vprintf, but we still need
425 // device-side declaration for it.
426 __device__ int printf(const char *, ...);
427 } // extern "C"
428 
429 // We also need device-side std::malloc and std::free.
430 namespace std {
431 __device__ static inline void free(void *__ptr) { ::free(__ptr); }
432 __device__ static inline void *malloc(size_t __size) {
433  return ::malloc(__size);
434 }
435 } // namespace std
436 
437 // Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to
438 // come after we've pulled in the definition of uint3 and dim3.
439 
440 __device__ inline __cuda_builtin_threadIdx_t::operator dim3() const {
441  return dim3(x, y, z);
442 }
443 
444 __device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
445  return {x, y, z};
446 }
447 
448 __device__ inline __cuda_builtin_blockIdx_t::operator dim3() const {
449  return dim3(x, y, z);
450 }
451 
452 __device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
453  return {x, y, z};
454 }
455 
456 __device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
457  return dim3(x, y, z);
458 }
459 
460 __device__ inline __cuda_builtin_blockDim_t::operator uint3() const {
461  return {x, y, z};
462 }
463 
464 __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
465  return dim3(x, y, z);
466 }
467 
468 __device__ inline __cuda_builtin_gridDim_t::operator uint3() const {
469  return {x, y, z};
470 }
471 
472 #include <__clang_cuda_cmath.h>
473 #include <__clang_cuda_intrinsics.h>
475 
476 // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
477 // mode, giving them their "proper" types of dim3 and uint3. This is
478 // incompatible with the types we give in __clang_cuda_builtin_vars.h. As as
479 // hack, force-include the header (nvcc doesn't include it by default) but
480 // redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are
481 // only used here for the redeclarations of blockDim and threadIdx.)
482 #pragma push_macro("dim3")
483 #pragma push_macro("uint3")
484 #define dim3 __cuda_builtin_blockDim_t
485 #define uint3 __cuda_builtin_threadIdx_t
486 #include "curand_mtgp32_kernel.h"
487 #pragma pop_macro("dim3")
488 #pragma pop_macro("uint3")
489 #pragma pop_macro("__USE_FAST_MATH__")
490 #pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
491 
492 // CUDA runtime uses this undocumented function to access kernel launch
493 // configuration. The declaration is in crt/device_functions.h but that file
494 // includes a lot of other stuff we don't want. Instead, we'll provide our own
495 // declaration for it here.
496 #if CUDA_VERSION >= 9020
497 extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
498  size_t sharedMem = 0,
499  void *stream = 0);
500 #endif
501 
502 #endif // __CUDA__
503 #endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__
erfcinv
__DEVICE__ double erfcinv(double __a)
Definition: __clang_cuda_math.h:97
signbit
__DEVICE__ bool signbit(float __x)
Test for sign bit.
Definition: __clang_cuda_cmath.h:170
rsqrtf
__DEVICE__ float rsqrtf(float __a)
Definition: __clang_cuda_math.h:276
cospi
__DEVICE__ double cospi(double __a)
Definition: __clang_cuda_math.h:88
__a
static __inline__ void int __a
Definition: emmintrin.h:4189
normcdf
__DEVICE__ double normcdf(double __a)
Definition: __clang_cuda_math.h:223
erfcxf
__DEVICE__ float erfcxf(float __a)
Definition: __clang_cuda_math.h:100
__nv_tex_surf_handler
#define __nv_tex_surf_handler(__op, __ptr,...)
Definition: __clang_cuda_texture_intrinsics.h:21
erfcinvf
__DEVICE__ float erfcinvf(float __a)
Definition: __clang_cuda_math.h:98
__clang_cuda_complex_builtins.h
__clang_cuda_cmath.h
__clang_cuda_libdevice_declares.h
rsqrt
__DEVICE__ double rsqrt(double __a)
Definition: __clang_cuda_math.h:275
blockDim
__CUDA_BUILTIN_VAR __cuda_builtin_blockDim_t blockDim
Definition: __clang_cuda_builtin_vars.h:108
sinpi
__DEVICE__ double sinpi(double __a)
Definition: __clang_cuda_math.h:311
rcbrt
__DEVICE__ double rcbrt(double __a)
Definition: __clang_cuda_math.h:234
rcbrtf
__DEVICE__ float rcbrtf(float __a)
Definition: __clang_cuda_math.h:235
__clang_cuda_texture_intrinsics.h
__signbitd
__DEVICE__ int __signbitd(double __a)
Definition: __clang_cuda_device_functions.h:518
__DEVICE__
#define __DEVICE__
Definition: __clang_cuda_cmath.h:38
x
IRgen optimization opportunities The common pattern of short x
Definition: README.txt:7
erfcx
__DEVICE__ double erfcx(double __a)
Definition: __clang_cuda_math.h:99
normcdfinvf
__DEVICE__ float normcdfinvf(float __a)
Definition: __clang_cuda_math.h:226
__clang_cuda_math.h
sincospi
__DEVICE_VOID__ void sincospi(double __a, double *__s, double *__c)
Definition: __clang_cuda_math.h:300
sinpif
__DEVICE__ float sinpif(float __a)
Definition: __clang_cuda_math.h:312
normcdff
__DEVICE__ float normcdff(float __a)
Definition: __clang_cuda_math.h:224
__clang_cuda_math_forward_declares.h
gridDim
__CUDA_BUILTIN_VAR __cuda_builtin_gridDim_t gridDim
Definition: __clang_cuda_builtin_vars.h:109
sincospif
__DEVICE_VOID__ void sincospif(float __a, float *__s, float *__c)
Definition: __clang_cuda_math.h:303
__clang_cuda_device_functions.h
std
Definition: Format.h:4275
__clang_cuda_intrinsics.h
__attribute__
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
Definition: __clang_hip_libdevice_declares.h:311
__b
static __inline__ vector float vector float __b
Definition: altivec.h:570
normcdfinv
__DEVICE__ double normcdfinv(double __a)
Definition: __clang_cuda_math.h:225
printf
int printf(__constant const char *st,...) __attribute__((format(printf
cospif
__DEVICE__ float cospif(float __a)
Definition: __clang_cuda_math.h:89
__c
static __inline__ vector float vector float vector float __c
Definition: altivec.h:4755
clang::transformer::name
RangeSelector name(std::string ID)
Given a node with a "name", (like NamedDecl, DeclRefExpr, CxxCtorInitializer, and TypeLoc) selects th...
Definition: RangeSelector.cpp:200
__clang_cuda_builtin_vars.h