Gecko [ access-piscine.fr ]

Name	Size	Permission
cuda_wrappers	[ DIR ]	drwxr-xr-x
llvm_libc_wrappers	[ DIR ]	drwxr-xr-x
openmp_wrappers	[ DIR ]	drwxr-xr-x
ppc_wrappers	[ DIR ]	drwxr-xr-x
__clang_cuda_builtin_vars.h	4.78 KB	-rw-r--r--
__clang_cuda_cmath.h	18.06 KB	-rw-r--r--
__clang_cuda_complex_builtins....	9.36 KB	-rw-r--r--
__clang_cuda_device_functions....	56.68 KB	-rw-r--r--
__clang_cuda_intrinsics.h	29.93 KB	-rw-r--r--
__clang_cuda_libdevice_declare...	21.87 KB	-rw-r--r--
__clang_cuda_math.h	15.99 KB	-rw-r--r--
__clang_cuda_math_forward_decl...	8.27 KB	-rw-r--r--
__clang_cuda_runtime_wrapper.h	17.61 KB	-rw-r--r--
__clang_cuda_texture_intrinsic...	31.86 KB	-rw-r--r--
__clang_hip_cmath.h	26.34 KB	-rw-r--r--
__clang_hip_libdevice_declares...	19.87 KB	-rw-r--r--
__clang_hip_math.h	31.96 KB	-rw-r--r--
__clang_hip_runtime_wrapper.h	4.65 KB	-rw-r--r--
__clang_hip_stdlib.h	1.19 KB	-rw-r--r--
__stddef_max_align_t.h	857 B	-rw-r--r--
__wmmintrin_aes.h	5.15 KB	-rw-r--r--
__wmmintrin_pclmul.h	1.99 KB	-rw-r--r--
adxintrin.h	7.37 KB	-rw-r--r--
altivec.h	697.32 KB	-rw-r--r--
ammintrin.h	7.54 KB	-rw-r--r--
amxcomplexintrin.h	6.81 KB	-rw-r--r--
amxfp16intrin.h	1.82 KB	-rw-r--r--
amxintrin.h	21.12 KB	-rw-r--r--
arm64intr.h	993 B	-rw-r--r--
arm_acle.h	25.66 KB	-rw-r--r--
arm_bf16.h	548 B	-rw-r--r--
arm_cde.h	32.67 KB	-rw-r--r--
arm_cmse.h	6.21 KB	-rw-r--r--
arm_fp16.h	16.92 KB	-rw-r--r--
arm_mve.h	1.48 MB	-rw-r--r--
arm_neon.h	2.45 MB	-rw-r--r--
arm_neon_sve_bridge.h	9.48 KB	-rw-r--r--
arm_sme_draft_spec_subject_to_...	60.2 KB	-rw-r--r--
arm_sve.h	1.51 MB	-rw-r--r--
armintr.h	843 B	-rw-r--r--
avx2intrin.h	186.96 KB	-rw-r--r--
avx512bf16intrin.h	10.51 KB	-rw-r--r--
avx512bitalgintrin.h	2.41 KB	-rw-r--r--
avx512bwintrin.h	75.33 KB	-rw-r--r--
avx512cdintrin.h	4.12 KB	-rw-r--r--
avx512dqintrin.h	58.75 KB	-rw-r--r--
avx512erintrin.h	11.83 KB	-rw-r--r--
avx512fintrin.h	382.64 KB	-rw-r--r--
avx512fp16intrin.h	156.63 KB	-rw-r--r--
avx512ifmaintrin.h	2.49 KB	-rw-r--r--
avx512ifmavlintrin.h	4.31 KB	-rw-r--r--
avx512pfintrin.h	4.53 KB	-rw-r--r--
avx512vbmi2intrin.h	13.17 KB	-rw-r--r--
avx512vbmiintrin.h	3.72 KB	-rw-r--r--
avx512vbmivlintrin.h	6.94 KB	-rw-r--r--
avx512vlbf16intrin.h	19.21 KB	-rw-r--r--
avx512vlbitalgintrin.h	4.23 KB	-rw-r--r--
avx512vlbwintrin.h	121.26 KB	-rw-r--r--
avx512vlcdintrin.h	7.66 KB	-rw-r--r--
avx512vldqintrin.h	46.41 KB	-rw-r--r--
avx512vlfp16intrin.h	85.51 KB	-rw-r--r--
avx512vlintrin.h	322.29 KB	-rw-r--r--
avx512vlvbmi2intrin.h	25.72 KB	-rw-r--r--
avx512vlvnniintrin.h	13.13 KB	-rw-r--r--
avx512vlvp2intersectintrin.h	4.44 KB	-rw-r--r--
avx512vnniintrin.h	4.21 KB	-rw-r--r--
avx512vp2intersectintrin.h	2.9 KB	-rw-r--r--
avx512vpopcntdqintrin.h	2 KB	-rw-r--r--
avx512vpopcntdqvlintrin.h	3.31 KB	-rw-r--r--
avxifmaintrin.h	5.75 KB	-rw-r--r--
avxintrin.h	195.41 KB	-rw-r--r--
avxneconvertintrin.h	14.09 KB	-rw-r--r--
avxvnniint16intrin.h	17.41 KB	-rw-r--r--
avxvnniint8intrin.h	18.67 KB	-rw-r--r--
avxvnniintrin.h	10.44 KB	-rw-r--r--
bmi2intrin.h	7.09 KB	-rw-r--r--
bmiintrin.h	14.12 KB	-rw-r--r--
builtins.h	741 B	-rw-r--r--
cet.h	1.49 KB	-rw-r--r--
cetintrin.h	3.27 KB	-rw-r--r--
cldemoteintrin.h	1.18 KB	-rw-r--r--
clflushoptintrin.h	1.17 KB	-rw-r--r--
clwbintrin.h	1.2 KB	-rw-r--r--
clzerointrin.h	1.19 KB	-rw-r--r--
cmpccxaddintrin.h	2.33 KB	-rw-r--r--
cpuid.h	11.01 KB	-rw-r--r--
crc32intrin.h	3.27 KB	-rw-r--r--
emmintrin.h	192.64 KB	-rw-r--r--
enqcmdintrin.h	2.12 KB	-rw-r--r--
f16cintrin.h	5.39 KB	-rw-r--r--
float.h	5.63 KB	-rw-r--r--
fma4intrin.h	6.82 KB	-rw-r--r--
fmaintrin.h	28.4 KB	-rw-r--r--
fxsrintrin.h	2.82 KB	-rw-r--r--
gfniintrin.h	7.57 KB	-rw-r--r--
hexagon_circ_brev_intrinsics.h	15.59 KB	-rw-r--r--
hexagon_protos.h	374.42 KB	-rw-r--r--
hexagon_types.h	130.33 KB	-rw-r--r--
hresetintrin.h	1.36 KB	-rw-r--r--
htmintrin.h	6.14 KB	-rw-r--r--
htmxlintrin.h	9.01 KB	-rw-r--r--
hvx_hexagon_protos.h	254.26 KB	-rw-r--r--
ia32intrin.h	12.72 KB	-rw-r--r--
immintrin.h	23.57 KB	-rw-r--r--
intrin.h	28.22 KB	-rw-r--r--
inttypes.h	2.26 KB	-rw-r--r--
invpcidintrin.h	764 B	-rw-r--r--
iso646.h	656 B	-rw-r--r--
keylockerintrin.h	17.98 KB	-rw-r--r--
larchintrin.h	7.8 KB	-rw-r--r--
limits.h	3.61 KB	-rw-r--r--
lwpintrin.h	5 KB	-rw-r--r--
lzcntintrin.h	3.18 KB	-rw-r--r--
mm3dnow.h	4.5 KB	-rw-r--r--
mm_malloc.h	1.88 KB	-rw-r--r--
mmintrin.h	55.98 KB	-rw-r--r--
module.modulemap	3.33 KB	-rw-r--r--
movdirintrin.h	1.57 KB	-rw-r--r--
msa.h	25.01 KB	-rw-r--r--
mwaitxintrin.h	2.19 KB	-rw-r--r--
nmmintrin.h	709 B	-rw-r--r--
opencl-c-base.h	30.38 KB	-rw-r--r--
opencl-c.h	874.39 KB	-rw-r--r--
pconfigintrin.h	1.19 KB	-rw-r--r--
pkuintrin.h	934 B	-rw-r--r--
pmmintrin.h	10.5 KB	-rw-r--r--
popcntintrin.h	1.82 KB	-rw-r--r--
prfchiintrin.h	2.02 KB	-rw-r--r--
prfchwintrin.h	2.06 KB	-rw-r--r--
ptwriteintrin.h	1.05 KB	-rw-r--r--
raointintrin.h	6.59 KB	-rw-r--r--
rdpruintrin.h	1.59 KB	-rw-r--r--
rdseedintrin.h	2.85 KB	-rw-r--r--
riscv_ntlh.h	855 B	-rw-r--r--
rtmintrin.h	1.25 KB	-rw-r--r--
s390intrin.h	604 B	-rw-r--r--
serializeintrin.h	881 B	-rw-r--r--
sgxintrin.h	1.77 KB	-rw-r--r--
sha512intrin.h	5.95 KB	-rw-r--r--
shaintrin.h	7.37 KB	-rw-r--r--
sifive_vector.h	522 B	-rw-r--r--
sm3intrin.h	7.29 KB	-rw-r--r--
sm4intrin.h	8.2 KB	-rw-r--r--
smmintrin.h	99.32 KB	-rw-r--r--
stdalign.h	911 B	-rw-r--r--
stdarg.h	1.66 KB	-rw-r--r--
stdatomic.h	8.3 KB	-rw-r--r--
stdbool.h	1.04 KB	-rw-r--r--
stddef.h	4.16 KB	-rw-r--r--
stdint.h	32.49 KB	-rw-r--r--
stdnoreturn.h	1.17 KB	-rw-r--r--
tbmintrin.h	3.15 KB	-rw-r--r--
tgmath.h	29.68 KB	-rw-r--r--
tmmintrin.h	29.51 KB	-rw-r--r--
tsxldtrkintrin.h	1.97 KB	-rw-r--r--
uintrintrin.h	4.96 KB	-rw-r--r--
unwind.h	11.21 KB	-rw-r--r--
vadefs.h	1.39 KB	-rw-r--r--
vaesintrin.h	2.46 KB	-rw-r--r--
varargs.h	477 B	-rw-r--r--
vecintrin.h	360.82 KB	-rw-r--r--
velintrin.h	2.1 KB	-rw-r--r--
velintrin_approx.h	3.54 KB	-rw-r--r--
velintrin_gen.h	69.06 KB	-rw-r--r--
vpclmulqdqintrin.h	1.06 KB	-rw-r--r--
waitpkgintrin.h	1.33 KB	-rw-r--r--
wasm_simd128.h	76.25 KB	-rw-r--r--
wbnoinvdintrin.h	749 B	-rw-r--r--
wmmintrin.h	659 B	-rw-r--r--
x86gprintrin.h	2.32 KB	-rw-r--r--
x86intrin.h	1.81 KB	-rw-r--r--
xmmintrin.h	106.73 KB	-rw-r--r--
xopintrin.h	19.96 KB	-rw-r--r--
xsavecintrin.h	2.51 KB	-rw-r--r--
xsaveintrin.h	1.64 KB	-rw-r--r--
xsaveoptintrin.h	1 KB	-rw-r--r--
xsavesintrin.h	1.24 KB	-rw-r--r--
xtestintrin.h	873 B	-rw-r--r--

Code Editor : __clang_cuda_runtime_wrapper.h

/*===---- __clang_cuda_runtime_wrapper.h - CUDA runtime support -------------===
 *
 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 * See https://llvm.org/LICENSE.txt for license information.
 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 *
 *===-----------------------------------------------------------------------===
 */

/*
 * WARNING: This header is intended to be directly -include'd by
 * the compiler and is not supposed to be included by users.
 *
 * CUDA headers are implemented in a way that currently makes it
 * impossible for user code to #include directly when compiling with
 * Clang. They present different view of CUDA-supplied functions
 * depending on where in NVCC's compilation pipeline the headers are
 * included. Neither of these modes provides function definitions with
 * correct attributes, so we use preprocessor to force the headers
 * into a form that Clang can use.
 *
 * Similarly to NVCC which -include's cuda_runtime.h, Clang -include's
 * this file during every CUDA compilation.
 */

#ifndef __CLANG_CUDA_RUNTIME_WRAPPER_H__
#define __CLANG_CUDA_RUNTIME_WRAPPER_H__

#if defined(__CUDA__) && defined(__clang__)

// Include some forward declares that must come before cmath.
#include <__clang_cuda_math_forward_declares.h>

// Define __CUDACC__ early as libstdc++ standard headers with GNU extensions
// enabled depend on it to avoid using __float128, which is unsupported in
// CUDA.
#define __CUDACC__

// Include some standard headers to avoid CUDA headers including them
// while some required macros (like __THROW) are in a weird state.
#include <cmath>
#include <cstdlib>
#include <stdlib.h>
#include <string.h>
#undef __CUDACC__

// Preserve common macros that will be changed below by us or by CUDA
// headers.
#pragma push_macro("__THROW")
#pragma push_macro("__CUDA_ARCH__")

// WARNING: Preprocessor hacks below are based on specific details of
// CUDA-7.x headers and are not expected to work with any other
// version of CUDA headers.
#include "cuda.h"
#if !defined(CUDA_VERSION)
#error "cuda.h did not define CUDA_VERSION"
#elif CUDA_VERSION < 7000
#error "Unsupported CUDA version!"
#endif

#pragma push_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")
#if CUDA_VERSION >= 10000
#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
#endif

// Make largest subset of device functions available during host
// compilation.
#ifndef __CUDA_ARCH__
#define __CUDA_ARCH__ 9999
#endif

#include "__clang_cuda_builtin_vars.h"

// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above
// has taken care of builtin variables declared in the file.
#define __DEVICE_LAUNCH_PARAMETERS_H__

// {math,device}_functions.h only have declarations of the
// functions. We don't need them as we're going to pull in their
// definitions from .hpp files.
#define __DEVICE_FUNCTIONS_H__
#define __MATH_FUNCTIONS_H__
#define __COMMON_FUNCTIONS_H__
// device_functions_decls is replaced by __clang_cuda_device_functions.h
// included below.
#define __DEVICE_FUNCTIONS_DECLS_H__

#undef __CUDACC__
#if CUDA_VERSION < 9000
#define __CUDABE__
#else
#define __CUDACC__
#define __CUDA_LIBDEVICE__
#endif
// Disables definitions of device-side runtime support stubs in
// cuda_device_runtime_api.h
#include "host_defines.h"
#undef __CUDACC__
#include "driver_types.h"
#include "host_config.h"

// Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in
// cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the
// functional equivalent of what we need.
#pragma push_macro("nv_weak")
#define nv_weak weak
#undef __CUDABE__
#undef __CUDA_LIBDEVICE__
#define __CUDACC__
#include "cuda_runtime.h"

#pragma pop_macro("nv_weak")
#undef __CUDACC__
#define __CUDABE__

// CUDA headers use __nvvm_memcpy and __nvvm_memset which Clang does
// not have at the moment. Emulate them with a builtin memcpy/memset.
#define __nvvm_memcpy(s, d, n, a) __builtin_memcpy(s, d, n)
#define __nvvm_memset(d, c, n, a) __builtin_memset(d, c, n)

#if CUDA_VERSION < 9000
#include "crt/device_runtime.h"
#endif
#include "crt/host_runtime.h"
// device_runtime.h defines __cxa_* macros that will conflict with
// cxxabi.h.
// FIXME: redefine these as __device__ functions.
#undef __cxa_vec_ctor
#undef __cxa_vec_cctor
#undef __cxa_vec_dtor
#undef __cxa_vec_new
#undef __cxa_vec_new2
#undef __cxa_vec_new3
#undef __cxa_vec_delete2
#undef __cxa_vec_delete
#undef __cxa_vec_delete3
#undef __cxa_pure_virtual

// math_functions.hpp expects this host function be defined on MacOS, but it
// ends up not being there because of the games we play here.  Just define it
// ourselves; it's simple enough.
#ifdef __APPLE__
inline __host__ double __signbitd(double x) {
  return std::signbit(x);
}
#endif

// CUDA 9.1 no longer provides declarations for libdevice functions, so we need
// to provide our own.
#include <__clang_cuda_libdevice_declares.h>

// Wrappers for many device-side standard library functions, incl. math
// functions, became compiler builtins in CUDA-9 and have been removed from the
// CUDA headers. Clang now provides its own implementation of the wrappers.
#if CUDA_VERSION >= 9000
#include <__clang_cuda_device_functions.h>
#include <__clang_cuda_math.h>
#endif

// __THROW is redefined to be empty by device_functions_decls.h in CUDA. Clang's
// counterpart does not do it, so we need to make it empty here to keep
// following CUDA includes happy.
#undef __THROW
#define __THROW

// CUDA 8.0.41 relies on __USE_FAST_MATH__ and __CUDA_PREC_DIV's values.
// Previous versions used to check whether they are defined or not.
// CU_DEVICE_INVALID macro is only defined in 8.0.41, so we use it
// here to detect the switch.

#if defined(CU_DEVICE_INVALID)
#if !defined(__USE_FAST_MATH__)
#define __USE_FAST_MATH__ 0
#endif

#if !defined(__CUDA_PREC_DIV)
#define __CUDA_PREC_DIV 0
#endif
#endif

// Temporarily poison __host__ macro to ensure it's not used by any of
// the headers we're about to include.
#pragma push_macro("__host__")
#define __host__ UNEXPECTED_HOST_ATTRIBUTE

// device_functions.hpp and math_functions*.hpp use 'static
// __forceinline__' (with no __device__) for definitions of device
// functions. Temporarily redefine __forceinline__ to include
// __device__.
#pragma push_macro("__forceinline__")
#define __forceinline__ __device__ __inline__ __attribute__((always_inline))
#if CUDA_VERSION < 9000
#include "device_functions.hpp"
#endif

// math_function.hpp uses the __USE_FAST_MATH__ macro to determine whether we
// get the slow-but-accurate or fast-but-inaccurate versions of functions like
// sin and exp.  This is controlled in clang by -fcuda-approx-transcendentals.
//
// device_functions.hpp uses __USE_FAST_MATH__ for a different purpose (fast vs.
// slow divides), so we need to scope our define carefully here.
#pragma push_macro("__USE_FAST_MATH__")
#if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
#define __USE_FAST_MATH__ 1
#endif

#if CUDA_VERSION >= 9000
#include "crt/math_functions.hpp"
#else
#include "math_functions.hpp"
#endif

#pragma pop_macro("__USE_FAST_MATH__")

#if CUDA_VERSION < 9000
#include "math_functions_dbl_ptx3.hpp"
#endif
#pragma pop_macro("__forceinline__")

// Pull in host-only functions that are only available when neither
// __CUDACC__ nor __CUDABE__ are defined.
#undef __MATH_FUNCTIONS_HPP__
#undef __CUDABE__
#if CUDA_VERSION < 9000
#include "math_functions.hpp"
#endif
// Alas, additional overloads for these functions are hard to get to.
// Considering that we only need these overloads for a few functions,
// we can provide them here.
static inline float rsqrt(float __a) { return rsqrtf(__a); }
static inline float rcbrt(float __a) { return rcbrtf(__a); }
static inline float sinpi(float __a) { return sinpif(__a); }
static inline float cospi(float __a) { return cospif(__a); }
static inline void sincospi(float __a, float *__b, float *__c) {
  return sincospif(__a, __b, __c);
}
static inline float erfcinv(float __a) { return erfcinvf(__a); }
static inline float normcdfinv(float __a) { return normcdfinvf(__a); }
static inline float normcdf(float __a) { return normcdff(__a); }
static inline float erfcx(float __a) { return erfcxf(__a); }

#if CUDA_VERSION < 9000
// For some reason single-argument variant is not always declared by
// CUDA headers. Alas, device_functions.hpp included below needs it.
static inline __device__ void __brkpt(int __c) { __brkpt(); }
#endif

// Now include *.hpp with definitions of various GPU functions.  Alas,
// a lot of thins get declared/defined with __host__ attribute which
// we don't want and we have to define it out. We also have to include
// {device,math}_functions.hpp again in order to extract the other
// branch of #if/else inside.
#define __host__
#undef __CUDABE__
#define __CUDACC__
#if CUDA_VERSION >= 9000
// Some atomic functions became compiler builtins in CUDA-9 , so we need their
// declarations.
#include "device_atomic_functions.h"
#endif
#undef __DEVICE_FUNCTIONS_HPP__
#include "device_atomic_functions.hpp"
#if CUDA_VERSION >= 9000
#include "crt/device_functions.hpp"
#include "crt/device_double_functions.hpp"
#else
#include "device_functions.hpp"
#define __CUDABE__
#include "device_double_functions.h"
#undef __CUDABE__
#endif
#include "sm_20_atomic_functions.hpp"
// Predicate functions used in `__builtin_assume` need to have no side effect.
// However, sm_20_intrinsics.hpp doesn't define them with neither pure nor
// const attribute. Rename definitions from sm_20_intrinsics.hpp and re-define
// them as pure ones.
#pragma push_macro("__isGlobal")
#pragma push_macro("__isShared")
#pragma push_macro("__isConstant")
#pragma push_macro("__isLocal")
#define __isGlobal __ignored_cuda___isGlobal
#define __isShared __ignored_cuda___isShared
#define __isConstant __ignored_cuda___isConstant
#define __isLocal __ignored_cuda___isLocal
#include "sm_20_intrinsics.hpp"
#pragma pop_macro("__isGlobal")
#pragma pop_macro("__isShared")
#pragma pop_macro("__isConstant")
#pragma pop_macro("__isLocal")
#pragma push_macro("__DEVICE__")
#define __DEVICE__ static __device__ __forceinline__ __attribute__((const))
__DEVICE__ unsigned int __isGlobal(const void *p) {
  return __nvvm_isspacep_global(p);
}
__DEVICE__ unsigned int __isShared(const void *p) {
  return __nvvm_isspacep_shared(p);
}
__DEVICE__ unsigned int __isConstant(const void *p) {
  return __nvvm_isspacep_const(p);
}
__DEVICE__ unsigned int __isLocal(const void *p) {
  return __nvvm_isspacep_local(p);
}
#pragma pop_macro("__DEVICE__")
#include "sm_32_atomic_functions.hpp"

// Don't include sm_30_intrinsics.h and sm_32_intrinsics.h.  These define the
// __shfl and __ldg intrinsics using inline (volatile) asm, but we want to
// define them using builtins so that the optimizer can reason about and across
// these instructions.  In particular, using intrinsics for ldg gets us the
// [addr+imm] addressing mode, which, although it doesn't actually exist in the
// hardware, seems to generate faster machine code because ptxas can more easily
// reason about our code.

#if CUDA_VERSION >= 8000
#pragma push_macro("__CUDA_ARCH__")
#undef __CUDA_ARCH__
#include "sm_60_atomic_functions.hpp"
#include "sm_61_intrinsics.hpp"
#pragma pop_macro("__CUDA_ARCH__")
#endif

#undef __MATH_FUNCTIONS_HPP__

// math_functions.hpp defines ::signbit as a __host__ __device__ function.  This
// conflicts with libstdc++'s constexpr ::signbit, so we have to rename
// math_function.hpp's ::signbit.  It's guarded by #undef signbit, but that's
// conditional on __GNUC__.  :)
#pragma push_macro("signbit")
#pragma push_macro("__GNUC__")
#undef __GNUC__
#define signbit __ignored_cuda_signbit

// CUDA-9 omits device-side definitions of some math functions if it sees
// include guard from math.h wrapper from libstdc++. We have to undo the header
// guard temporarily to get the definitions we need.
#pragma push_macro("_GLIBCXX_MATH_H")
#pragma push_macro("_LIBCPP_VERSION")
#if CUDA_VERSION >= 9000
#undef _GLIBCXX_MATH_H
// We also need to undo another guard that checks for libc++ 3.8+
#ifdef _LIBCPP_VERSION
#define _LIBCPP_VERSION 3700
#endif
#endif

#if CUDA_VERSION >= 9000
#include "crt/math_functions.hpp"
#else
#include "math_functions.hpp"
#endif
#pragma pop_macro("_GLIBCXX_MATH_H")
#pragma pop_macro("_LIBCPP_VERSION")
#pragma pop_macro("__GNUC__")
#pragma pop_macro("signbit")

#pragma pop_macro("__host__")

// __clang_cuda_texture_intrinsics.h must be included first in order to provide
// implementation for __nv_tex_surf_handler that CUDA's headers depend on.
// The implementation requires c++11 and only works with CUDA-9 or newer.
#if __cplusplus >= 201103L && CUDA_VERSION >= 9000
// clang-format off
#include <__clang_cuda_texture_intrinsics.h>
// clang-format on
#else
#if CUDA_VERSION >= 9000
// Provide a hint that texture support needs C++11.
template <typename T> struct __nv_tex_needs_cxx11 {
  const static bool value = false;
};
template <class T>
__host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
                                               cudaTextureObject_t obj,
                                               float x) {
  _Static_assert(__nv_tex_needs_cxx11<T>::value,
                 "Texture support requires C++11");
}
#else
// Textures in CUDA-8 and older are not supported by clang.There's no
// convenient way to intercept texture use in these versions, so we can't
// produce a meaningful error. The source code that attempts to use textures
// will continue to fail as it does now.
#endif // CUDA_VERSION
#endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
#include "texture_fetch_functions.h"
#include "texture_indirect_functions.h"

// Restore state of __CUDA_ARCH__ and __THROW we had on entry.
#pragma pop_macro("__CUDA_ARCH__")
#pragma pop_macro("__THROW")

// Set up compiler macros expected to be seen during compilation.
#undef __CUDABE__
#define __CUDACC__

extern "C" {
// Device-side CUDA system calls.
// http://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html#system-calls
// We need these declarations and wrappers for device-side
// malloc/free/printf calls to work without relying on
// -fcuda-disable-target-call-checks option.
__device__ int vprintf(const char *, const char *);
__device__ void free(void *) __attribute((nothrow));
__device__ void *malloc(size_t) __attribute((nothrow)) __attribute__((malloc));

// __assertfail() used to have a `noreturn` attribute. Unfortunately that
// contributed to triggering the longstanding bug in ptxas when assert was used
// in sufficiently convoluted code. See
// https://bugs.llvm.org/show_bug.cgi?id=27738 for the details.
__device__ void __assertfail(const char *__message, const char *__file,
                             unsigned __line, const char *__function,
                             size_t __charSize);

// In order for standard assert() macro on linux to work we need to
// provide device-side __assert_fail()
__device__ static inline void __assert_fail(const char *__message,
                                            const char *__file, unsigned __line,
                                            const char *__function) {
  __assertfail(__message, __file, __line, __function, sizeof(char));
}

// Clang will convert printf into vprintf, but we still need
// device-side declaration for it.
__device__ int printf(const char *, ...);
} // extern "C"

// We also need device-side std::malloc and std::free.
namespace std {
__device__ static inline void free(void *__ptr) { ::free(__ptr); }
__device__ static inline void *malloc(size_t __size) {
  return ::malloc(__size);
}
} // namespace std

// Out-of-line implementations from __clang_cuda_builtin_vars.h.  These need to
// come after we've pulled in the definition of uint3 and dim3.

__device__ inline __cuda_builtin_threadIdx_t::operator dim3() const {
  return dim3(x, y, z);
}

__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
  return {x, y, z};
}

__device__ inline __cuda_builtin_blockIdx_t::operator dim3() const {
  return dim3(x, y, z);
}

__device__ inline __cuda_builtin_blockIdx_t::operator uint3() const {
  return {x, y, z};
}

__device__ inline __cuda_builtin_blockDim_t::operator dim3() const {
  return dim3(x, y, z);
}

__device__ inline __cuda_builtin_blockDim_t::operator uint3() const {
  return {x, y, z};
}

__device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
  return dim3(x, y, z);
}

__device__ inline __cuda_builtin_gridDim_t::operator uint3() const {
  return {x, y, z};
}

#include <__clang_cuda_cmath.h>
#include <__clang_cuda_intrinsics.h>
#include <__clang_cuda_complex_builtins.h>

// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
// mode, giving them their "proper" types of dim3 and uint3.  This is
// incompatible with the types we give in __clang_cuda_builtin_vars.h.  As as
// hack, force-include the header (nvcc doesn't include it by default) but
// redefine dim3 and uint3 to our builtin types.  (Thankfully dim3 and uint3 are
// only used here for the redeclarations of blockDim and threadIdx.)
#pragma push_macro("dim3")
#pragma push_macro("uint3")
#define dim3 __cuda_builtin_blockDim_t
#define uint3 __cuda_builtin_threadIdx_t
#include "curand_mtgp32_kernel.h"
#pragma pop_macro("dim3")
#pragma pop_macro("uint3")
#pragma pop_macro("__USE_FAST_MATH__")
#pragma pop_macro("__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__")

// CUDA runtime uses this undocumented function to access kernel launch
// configuration. The declaration is in crt/device_functions.h but that file
// includes a lot of other stuff we don't want. Instead, we'll provide our own
// declaration for it here.
#if CUDA_VERSION >= 9020
extern "C" unsigned __cudaPushCallConfiguration(dim3 gridDim, dim3 blockDim,
                                                size_t sharedMem = 0,
                                                void *stream = 0);
#endif

#endif // __CUDA__
#endif // __CLANG_CUDA_RUNTIME_WRAPPER_H__

${this.title}

Code Editor : __clang_cuda_runtime_wrapper.h