gtbercea updated this revision to Diff 160598.
gtbercea added a comment.
Herald added a subscriber: jholewinski.
Add __NO_MATH_INLINES macro for the NVPTX toolchain to prevent any host
assembly from seeping onto the device.
Repository:
rC Clang
https://reviews.llvm.org/D47849
Files:
include/clang/Driver/ToolChain.h
lib/Basic/Targets/NVPTX.cpp
lib/Driver/ToolChains/Clang.cpp
lib/Driver/ToolChains/Cuda.cpp
lib/Driver/ToolChains/Cuda.h
lib/Headers/CMakeLists.txt
lib/Headers/__clang_cuda_device_functions.h
lib/Headers/__clang_cuda_libdevice_declares.h
test/CodeGen/nvptx_device_math_functions.c
test/Driver/openmp-offload-gpu.c
Index: test/Driver/openmp-offload-gpu.c
===================================================================
--- test/Driver/openmp-offload-gpu.c
+++ test/Driver/openmp-offload-gpu.c
@@ -76,9 +76,9 @@
// RUN: -no-canonical-prefixes -save-temps %t.o -fopenmp-use-target-bundling 2>&1 \
// RUN: | FileCheck -check-prefix=CHK-CUBIN-UNBUNDLING-NVLINK %s
-/// Use DAG to ensure that cubin file has been unbundled.
+/// Use DAG to ensure that object file has not been unbundled.
// CHK-CUBIN-UNBUNDLING-NVLINK-DAG: nvlink{{.*}}" {{.*}}"[[CUBIN:.*\.cubin]]"
-// CHK-CUBIN-UNBUNDLING-NVLINK-DAG: clang-offload-bundler{{.*}}" "-type=o" {{.*}}"-outputs={{.*}}[[CUBIN]]
+// CHK-CUBIN-UNBUNDLING-NVLINK-DAG: clang-offload-bundler{{.*}}" "-type=o" {{.*}}[[CUBIN]]
// CHK-CUBIN-UNBUNDLING-NVLINK-DAG-SAME: "-unbundle"
/// ###########################################################################
Index: test/CodeGen/nvptx_device_math_functions.c
===================================================================
--- /dev/null
+++ test/CodeGen/nvptx_device_math_functions.c
@@ -0,0 +1,20 @@
+// Test calling of device math functions.
+///==========================================================================///
+
+// RUN: %clang -fmath-errno -S -emit-llvm -o - %s -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda | FileCheck -check-prefix CHECK-YES %s
+
+void test_sqrt(double a1) {
+ #pragma omp target
+ {
+ // CHECK-YES: call double @llvm.nvvm.sqrt.rn.d(double
+ double l1 = sqrt(a1);
+ }
+}
+
+void test_pow(float a0, double a1, long double a2) {
+ #pragma omp target
+ {
+ // CHECK-YES: call double @__internal_accurate_pow(double
+ double l1 = pow(a1, a1);
+ }
+}
Index: lib/Headers/__clang_cuda_libdevice_declares.h
===================================================================
--- lib/Headers/__clang_cuda_libdevice_declares.h
+++ lib/Headers/__clang_cuda_libdevice_declares.h
@@ -24,443 +24,455 @@
#ifndef __CLANG_CUDA_LIBDEVICE_DECLARES_H__
#define __CLANG_CUDA_LIBDEVICE_DECLARES_H__
+#if defined(_OPENMP)
+#define __DEVICE__
+#elif defined(__CUDA__)
+#define __DEVICE__ __device__
+#endif
+
+#if defined(__cplusplus)
extern "C" {
+#endif
-__device__ int __nv_abs(int __a);
-__device__ double __nv_acos(double __a);
-__device__ float __nv_acosf(float __a);
-__device__ double __nv_acosh(double __a);
-__device__ float __nv_acoshf(float __a);
-__device__ double __nv_asin(double __a);
-__device__ float __nv_asinf(float __a);
-__device__ double __nv_asinh(double __a);
-__device__ float __nv_asinhf(float __a);
-__device__ double __nv_atan2(double __a, double __b);
-__device__ float __nv_atan2f(float __a, float __b);
-__device__ double __nv_atan(double __a);
-__device__ float __nv_atanf(float __a);
-__device__ double __nv_atanh(double __a);
-__device__ float __nv_atanhf(float __a);
-__device__ int __nv_brev(int __a);
-__device__ long long __nv_brevll(long long __a);
-__device__ int __nv_byte_perm(int __a, int __b, int __c);
-__device__ double __nv_cbrt(double __a);
-__device__ float __nv_cbrtf(float __a);
-__device__ double __nv_ceil(double __a);
-__device__ float __nv_ceilf(float __a);
-__device__ int __nv_clz(int __a);
-__device__ int __nv_clzll(long long __a);
-__device__ double __nv_copysign(double __a, double __b);
-__device__ float __nv_copysignf(float __a, float __b);
-__device__ double __nv_cos(double __a);
-__device__ float __nv_cosf(float __a);
-__device__ double __nv_cosh(double __a);
-__device__ float __nv_coshf(float __a);
-__device__ double __nv_cospi(double __a);
-__device__ float __nv_cospif(float __a);
-__device__ double __nv_cyl_bessel_i0(double __a);
-__device__ float __nv_cyl_bessel_i0f(float __a);
-__device__ double __nv_cyl_bessel_i1(double __a);
-__device__ float __nv_cyl_bessel_i1f(float __a);
-__device__ double __nv_dadd_rd(double __a, double __b);
-__device__ double __nv_dadd_rn(double __a, double __b);
-__device__ double __nv_dadd_ru(double __a, double __b);
-__device__ double __nv_dadd_rz(double __a, double __b);
-__device__ double __nv_ddiv_rd(double __a, double __b);
-__device__ double __nv_ddiv_rn(double __a, double __b);
-__device__ double __nv_ddiv_ru(double __a, double __b);
-__device__ double __nv_ddiv_rz(double __a, double __b);
-__device__ double __nv_dmul_rd(double __a, double __b);
-__device__ double __nv_dmul_rn(double __a, double __b);
-__device__ double __nv_dmul_ru(double __a, double __b);
-__device__ double __nv_dmul_rz(double __a, double __b);
-__device__ float __nv_double2float_rd(double __a);
-__device__ float __nv_double2float_rn(double __a);
-__device__ float __nv_double2float_ru(double __a);
-__device__ float __nv_double2float_rz(double __a);
-__device__ int __nv_double2hiint(double __a);
-__device__ int __nv_double2int_rd(double __a);
-__device__ int __nv_double2int_rn(double __a);
-__device__ int __nv_double2int_ru(double __a);
-__device__ int __nv_double2int_rz(double __a);
-__device__ long long __nv_double2ll_rd(double __a);
-__device__ long long __nv_double2ll_rn(double __a);
-__device__ long long __nv_double2ll_ru(double __a);
-__device__ long long __nv_double2ll_rz(double __a);
-__device__ int __nv_double2loint(double __a);
-__device__ unsigned int __nv_double2uint_rd(double __a);
-__device__ unsigned int __nv_double2uint_rn(double __a);
-__device__ unsigned int __nv_double2uint_ru(double __a);
-__device__ unsigned int __nv_double2uint_rz(double __a);
-__device__ unsigned long long __nv_double2ull_rd(double __a);
-__device__ unsigned long long __nv_double2ull_rn(double __a);
-__device__ unsigned long long __nv_double2ull_ru(double __a);
-__device__ unsigned long long __nv_double2ull_rz(double __a);
-__device__ unsigned long long __nv_double_as_longlong(double __a);
-__device__ double __nv_drcp_rd(double __a);
-__device__ double __nv_drcp_rn(double __a);
-__device__ double __nv_drcp_ru(double __a);
-__device__ double __nv_drcp_rz(double __a);
-__device__ double __nv_dsqrt_rd(double __a);
-__device__ double __nv_dsqrt_rn(double __a);
-__device__ double __nv_dsqrt_ru(double __a);
-__device__ double __nv_dsqrt_rz(double __a);
-__device__ double __nv_dsub_rd(double __a, double __b);
-__device__ double __nv_dsub_rn(double __a, double __b);
-__device__ double __nv_dsub_ru(double __a, double __b);
-__device__ double __nv_dsub_rz(double __a, double __b);
-__device__ double __nv_erfc(double __a);
-__device__ float __nv_erfcf(float __a);
-__device__ double __nv_erfcinv(double __a);
-__device__ float __nv_erfcinvf(float __a);
-__device__ double __nv_erfcx(double __a);
-__device__ float __nv_erfcxf(float __a);
-__device__ double __nv_erf(double __a);
-__device__ float __nv_erff(float __a);
-__device__ double __nv_erfinv(double __a);
-__device__ float __nv_erfinvf(float __a);
-__device__ double __nv_exp10(double __a);
-__device__ float __nv_exp10f(float __a);
-__device__ double __nv_exp2(double __a);
-__device__ float __nv_exp2f(float __a);
-__device__ double __nv_exp(double __a);
-__device__ float __nv_expf(float __a);
-__device__ double __nv_expm1(double __a);
-__device__ float __nv_expm1f(float __a);
-__device__ double __nv_fabs(double __a);
-__device__ float __nv_fabsf(float __a);
-__device__ float __nv_fadd_rd(float __a, float __b);
-__device__ float __nv_fadd_rn(float __a, float __b);
-__device__ float __nv_fadd_ru(float __a, float __b);
-__device__ float __nv_fadd_rz(float __a, float __b);
-__device__ float __nv_fast_cosf(float __a);
-__device__ float __nv_fast_exp10f(float __a);
-__device__ float __nv_fast_expf(float __a);
-__device__ float __nv_fast_fdividef(float __a, float __b);
-__device__ float __nv_fast_log10f(float __a);
-__device__ float __nv_fast_log2f(float __a);
-__device__ float __nv_fast_logf(float __a);
-__device__ float __nv_fast_powf(float __a, float __b);
-__device__ void __nv_fast_sincosf(float __a, float *__sptr, float *__cptr);
-__device__ float __nv_fast_sinf(float __a);
-__device__ float __nv_fast_tanf(float __a);
-__device__ double __nv_fdim(double __a, double __b);
-__device__ float __nv_fdimf(float __a, float __b);
-__device__ float __nv_fdiv_rd(float __a, float __b);
-__device__ float __nv_fdiv_rn(float __a, float __b);
-__device__ float __nv_fdiv_ru(float __a, float __b);
-__device__ float __nv_fdiv_rz(float __a, float __b);
-__device__ int __nv_ffs(int __a);
-__device__ int __nv_ffsll(long long __a);
-__device__ int __nv_finitef(float __a);
-__device__ unsigned short __nv_float2half_rn(float __a);
-__device__ int __nv_float2int_rd(float __a);
-__device__ int __nv_float2int_rn(float __a);
-__device__ int __nv_float2int_ru(float __a);
-__device__ int __nv_float2int_rz(float __a);
-__device__ long long __nv_float2ll_rd(float __a);
-__device__ long long __nv_float2ll_rn(float __a);
-__device__ long long __nv_float2ll_ru(float __a);
-__device__ long long __nv_float2ll_rz(float __a);
-__device__ unsigned int __nv_float2uint_rd(float __a);
-__device__ unsigned int __nv_float2uint_rn(float __a);
-__device__ unsigned int __nv_float2uint_ru(float __a);
-__device__ unsigned int __nv_float2uint_rz(float __a);
-__device__ unsigned long long __nv_float2ull_rd(float __a);
-__device__ unsigned long long __nv_float2ull_rn(float __a);
-__device__ unsigned long long __nv_float2ull_ru(float __a);
-__device__ unsigned long long __nv_float2ull_rz(float __a);
-__device__ int __nv_float_as_int(float __a);
-__device__ unsigned int __nv_float_as_uint(float __a);
-__device__ double __nv_floor(double __a);
-__device__ float __nv_floorf(float __a);
-__device__ double __nv_fma(double __a, double __b, double __c);
-__device__ float __nv_fmaf(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_rd(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_rn(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_ru(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ieee_rz(float __a, float __b, float __c);
-__device__ float __nv_fmaf_rd(float __a, float __b, float __c);
-__device__ float __nv_fmaf_rn(float __a, float __b, float __c);
-__device__ float __nv_fmaf_ru(float __a, float __b, float __c);
-__device__ float __nv_fmaf_rz(float __a, float __b, float __c);
-__device__ double __nv_fma_rd(double __a, double __b, double __c);
-__device__ double __nv_fma_rn(double __a, double __b, double __c);
-__device__ double __nv_fma_ru(double __a, double __b, double __c);
-__device__ double __nv_fma_rz(double __a, double __b, double __c);
-__device__ double __nv_fmax(double __a, double __b);
-__device__ float __nv_fmaxf(float __a, float __b);
-__device__ double __nv_fmin(double __a, double __b);
-__device__ float __nv_fminf(float __a, float __b);
-__device__ double __nv_fmod(double __a, double __b);
-__device__ float __nv_fmodf(float __a, float __b);
-__device__ float __nv_fmul_rd(float __a, float __b);
-__device__ float __nv_fmul_rn(float __a, float __b);
-__device__ float __nv_fmul_ru(float __a, float __b);
-__device__ float __nv_fmul_rz(float __a, float __b);
-__device__ float __nv_frcp_rd(float __a);
-__device__ float __nv_frcp_rn(float __a);
-__device__ float __nv_frcp_ru(float __a);
-__device__ float __nv_frcp_rz(float __a);
-__device__ double __nv_frexp(double __a, int *__b);
-__device__ float __nv_frexpf(float __a, int *__b);
-__device__ float __nv_frsqrt_rn(float __a);
-__device__ float __nv_fsqrt_rd(float __a);
-__device__ float __nv_fsqrt_rn(float __a);
-__device__ float __nv_fsqrt_ru(float __a);
-__device__ float __nv_fsqrt_rz(float __a);
-__device__ float __nv_fsub_rd(float __a, float __b);
-__device__ float __nv_fsub_rn(float __a, float __b);
-__device__ float __nv_fsub_ru(float __a, float __b);
-__device__ float __nv_fsub_rz(float __a, float __b);
-__device__ int __nv_hadd(int __a, int __b);
-__device__ float __nv_half2float(unsigned short __h);
-__device__ double __nv_hiloint2double(int __a, int __b);
-__device__ double __nv_hypot(double __a, double __b);
-__device__ float __nv_hypotf(float __a, float __b);
-__device__ int __nv_ilogb(double __a);
-__device__ int __nv_ilogbf(float __a);
-__device__ double __nv_int2double_rn(int __a);
-__device__ float __nv_int2float_rd(int __a);
-__device__ float __nv_int2float_rn(int __a);
-__device__ float __nv_int2float_ru(int __a);
-__device__ float __nv_int2float_rz(int __a);
-__device__ float __nv_int_as_float(int __a);
-__device__ int __nv_isfinited(double __a);
-__device__ int __nv_isinfd(double __a);
-__device__ int __nv_isinff(float __a);
-__device__ int __nv_isnand(double __a);
-__device__ int __nv_isnanf(float __a);
-__device__ double __nv_j0(double __a);
-__device__ float __nv_j0f(float __a);
-__device__ double __nv_j1(double __a);
-__device__ float __nv_j1f(float __a);
-__device__ float __nv_jnf(int __a, float __b);
-__device__ double __nv_jn(int __a, double __b);
-__device__ double __nv_ldexp(double __a, int __b);
-__device__ float __nv_ldexpf(float __a, int __b);
-__device__ double __nv_lgamma(double __a);
-__device__ float __nv_lgammaf(float __a);
-__device__ double __nv_ll2double_rd(long long __a);
-__device__ double __nv_ll2double_rn(long long __a);
-__device__ double __nv_ll2double_ru(long long __a);
-__device__ double __nv_ll2double_rz(long long __a);
-__device__ float __nv_ll2float_rd(long long __a);
-__device__ float __nv_ll2float_rn(long long __a);
-__device__ float __nv_ll2float_ru(long long __a);
-__device__ float __nv_ll2float_rz(long long __a);
-__device__ long long __nv_llabs(long long __a);
-__device__ long long __nv_llmax(long long __a, long long __b);
-__device__ long long __nv_llmin(long long __a, long long __b);
-__device__ long long __nv_llrint(double __a);
-__device__ long long __nv_llrintf(float __a);
-__device__ long long __nv_llround(double __a);
-__device__ long long __nv_llroundf(float __a);
-__device__ double __nv_log10(double __a);
-__device__ float __nv_log10f(float __a);
-__device__ double __nv_log1p(double __a);
-__device__ float __nv_log1pf(float __a);
-__device__ double __nv_log2(double __a);
-__device__ float __nv_log2f(float __a);
-__device__ double __nv_logb(double __a);
-__device__ float __nv_logbf(float __a);
-__device__ double __nv_log(double __a);
-__device__ float __nv_logf(float __a);
-__device__ double __nv_longlong_as_double(long long __a);
-__device__ int __nv_max(int __a, int __b);
-__device__ int __nv_min(int __a, int __b);
-__device__ double __nv_modf(double __a, double *__b);
-__device__ float __nv_modff(float __a, float *__b);
-__device__ int __nv_mul24(int __a, int __b);
-__device__ long long __nv_mul64hi(long long __a, long long __b);
-__device__ int __nv_mulhi(int __a, int __b);
-__device__ double __nv_nan(const signed char *__a);
-__device__ float __nv_nanf(const signed char *__a);
-__device__ double __nv_nearbyint(double __a);
-__device__ float __nv_nearbyintf(float __a);
-__device__ double __nv_nextafter(double __a, double __b);
-__device__ float __nv_nextafterf(float __a, float __b);
-__device__ double __nv_norm3d(double __a, double __b, double __c);
-__device__ float __nv_norm3df(float __a, float __b, float __c);
-__device__ double __nv_norm4d(double __a, double __b, double __c, double __d);
-__device__ float __nv_norm4df(float __a, float __b, float __c, float __d);
-__device__ double __nv_normcdf(double __a);
-__device__ float __nv_normcdff(float __a);
-__device__ double __nv_normcdfinv(double __a);
-__device__ float __nv_normcdfinvf(float __a);
-__device__ float __nv_normf(int __a, const float *__b);
-__device__ double __nv_norm(int __a, const double *__b);
-__device__ int __nv_popc(int __a);
-__device__ int __nv_popcll(long long __a);
-__device__ double __nv_pow(double __a, double __b);
-__device__ float __nv_powf(float __a, float __b);
-__device__ double __nv_powi(double __a, int __b);
-__device__ float __nv_powif(float __a, int __b);
-__device__ double __nv_rcbrt(double __a);
-__device__ float __nv_rcbrtf(float __a);
-__device__ double __nv_rcp64h(double __a);
-__device__ double __nv_remainder(double __a, double __b);
-__device__ float __nv_remainderf(float __a, float __b);
-__device__ double __nv_remquo(double __a, double __b, int *__c);
-__device__ float __nv_remquof(float __a, float __b, int *__c);
-__device__ int __nv_rhadd(int __a, int __b);
-__device__ double __nv_rhypot(double __a, double __b);
-__device__ float __nv_rhypotf(float __a, float __b);
-__device__ double __nv_rint(double __a);
-__device__ float __nv_rintf(float __a);
-__device__ double __nv_rnorm3d(double __a, double __b, double __c);
-__device__ float __nv_rnorm3df(float __a, float __b, float __c);
-__device__ double __nv_rnorm4d(double __a, double __b, double __c, double __d);
-__device__ float __nv_rnorm4df(float __a, float __b, float __c, float __d);
-__device__ float __nv_rnormf(int __a, const float *__b);
-__device__ double __nv_rnorm(int __a, const double *__b);
-__device__ double __nv_round(double __a);
-__device__ float __nv_roundf(float __a);
-__device__ double __nv_rsqrt(double __a);
-__device__ float __nv_rsqrtf(float __a);
-__device__ int __nv_sad(int __a, int __b, int __c);
-__device__ float __nv_saturatef(float __a);
-__device__ double __nv_scalbn(double __a, int __b);
-__device__ float __nv_scalbnf(float __a, int __b);
-__device__ int __nv_signbitd(double __a);
-__device__ int __nv_signbitf(float __a);
-__device__ void __nv_sincos(double __a, double *__b, double *__c);
-__device__ void __nv_sincosf(float __a, float *__b, float *__c);
-__device__ void __nv_sincospi(double __a, double *__b, double *__c);
-__device__ void __nv_sincospif(float __a, float *__b, float *__c);
-__device__ double __nv_sin(double __a);
-__device__ float __nv_sinf(float __a);
-__device__ double __nv_sinh(double __a);
-__device__ float __nv_sinhf(float __a);
-__device__ double __nv_sinpi(double __a);
-__device__ float __nv_sinpif(float __a);
-__device__ double __nv_sqrt(double __a);
-__device__ float __nv_sqrtf(float __a);
-__device__ double __nv_tan(double __a);
-__device__ float __nv_tanf(float __a);
-__device__ double __nv_tanh(double __a);
-__device__ float __nv_tanhf(float __a);
-__device__ double __nv_tgamma(double __a);
-__device__ float __nv_tgammaf(float __a);
-__device__ double __nv_trunc(double __a);
-__device__ float __nv_truncf(float __a);
-__device__ int __nv_uhadd(unsigned int __a, unsigned int __b);
-__device__ double __nv_uint2double_rn(unsigned int __i);
-__device__ float __nv_uint2float_rd(unsigned int __a);
-__device__ float __nv_uint2float_rn(unsigned int __a);
-__device__ float __nv_uint2float_ru(unsigned int __a);
-__device__ float __nv_uint2float_rz(unsigned int __a);
-__device__ float __nv_uint_as_float(unsigned int __a);
-__device__ double __nv_ull2double_rd(unsigned long long __a);
-__device__ double __nv_ull2double_rn(unsigned long long __a);
-__device__ double __nv_ull2double_ru(unsigned long long __a);
-__device__ double __nv_ull2double_rz(unsigned long long __a);
-__device__ float __nv_ull2float_rd(unsigned long long __a);
-__device__ float __nv_ull2float_rn(unsigned long long __a);
-__device__ float __nv_ull2float_ru(unsigned long long __a);
-__device__ float __nv_ull2float_rz(unsigned long long __a);
-__device__ unsigned long long __nv_ullmax(unsigned long long __a,
+__DEVICE__ int __nv_abs(int __a);
+__DEVICE__ double __nv_acos(double __a);
+__DEVICE__ float __nv_acosf(float __a);
+__DEVICE__ double __nv_acosh(double __a);
+__DEVICE__ float __nv_acoshf(float __a);
+__DEVICE__ double __nv_asin(double __a);
+__DEVICE__ float __nv_asinf(float __a);
+__DEVICE__ double __nv_asinh(double __a);
+__DEVICE__ float __nv_asinhf(float __a);
+__DEVICE__ double __nv_atan2(double __a, double __b);
+__DEVICE__ float __nv_atan2f(float __a, float __b);
+__DEVICE__ double __nv_atan(double __a);
+__DEVICE__ float __nv_atanf(float __a);
+__DEVICE__ double __nv_atanh(double __a);
+__DEVICE__ float __nv_atanhf(float __a);
+__DEVICE__ int __nv_brev(int __a);
+__DEVICE__ long long __nv_brevll(long long __a);
+__DEVICE__ int __nv_byte_perm(int __a, int __b, int __c);
+__DEVICE__ double __nv_cbrt(double __a);
+__DEVICE__ float __nv_cbrtf(float __a);
+__DEVICE__ double __nv_ceil(double __a);
+__DEVICE__ float __nv_ceilf(float __a);
+__DEVICE__ int __nv_clz(int __a);
+__DEVICE__ int __nv_clzll(long long __a);
+__DEVICE__ double __nv_copysign(double __a, double __b);
+__DEVICE__ float __nv_copysignf(float __a, float __b);
+__DEVICE__ double __nv_cos(double __a);
+__DEVICE__ float __nv_cosf(float __a);
+__DEVICE__ double __nv_cosh(double __a);
+__DEVICE__ float __nv_coshf(float __a);
+__DEVICE__ double __nv_cospi(double __a);
+__DEVICE__ float __nv_cospif(float __a);
+__DEVICE__ double __nv_cyl_bessel_i0(double __a);
+__DEVICE__ float __nv_cyl_bessel_i0f(float __a);
+__DEVICE__ double __nv_cyl_bessel_i1(double __a);
+__DEVICE__ float __nv_cyl_bessel_i1f(float __a);
+__DEVICE__ double __nv_dadd_rd(double __a, double __b);
+__DEVICE__ double __nv_dadd_rn(double __a, double __b);
+__DEVICE__ double __nv_dadd_ru(double __a, double __b);
+__DEVICE__ double __nv_dadd_rz(double __a, double __b);
+__DEVICE__ double __nv_ddiv_rd(double __a, double __b);
+__DEVICE__ double __nv_ddiv_rn(double __a, double __b);
+__DEVICE__ double __nv_ddiv_ru(double __a, double __b);
+__DEVICE__ double __nv_ddiv_rz(double __a, double __b);
+__DEVICE__ double __nv_dmul_rd(double __a, double __b);
+__DEVICE__ double __nv_dmul_rn(double __a, double __b);
+__DEVICE__ double __nv_dmul_ru(double __a, double __b);
+__DEVICE__ double __nv_dmul_rz(double __a, double __b);
+__DEVICE__ float __nv_double2float_rd(double __a);
+__DEVICE__ float __nv_double2float_rn(double __a);
+__DEVICE__ float __nv_double2float_ru(double __a);
+__DEVICE__ float __nv_double2float_rz(double __a);
+__DEVICE__ int __nv_double2hiint(double __a);
+__DEVICE__ int __nv_double2int_rd(double __a);
+__DEVICE__ int __nv_double2int_rn(double __a);
+__DEVICE__ int __nv_double2int_ru(double __a);
+__DEVICE__ int __nv_double2int_rz(double __a);
+__DEVICE__ long long __nv_double2ll_rd(double __a);
+__DEVICE__ long long __nv_double2ll_rn(double __a);
+__DEVICE__ long long __nv_double2ll_ru(double __a);
+__DEVICE__ long long __nv_double2ll_rz(double __a);
+__DEVICE__ int __nv_double2loint(double __a);
+__DEVICE__ unsigned int __nv_double2uint_rd(double __a);
+__DEVICE__ unsigned int __nv_double2uint_rn(double __a);
+__DEVICE__ unsigned int __nv_double2uint_ru(double __a);
+__DEVICE__ unsigned int __nv_double2uint_rz(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_rd(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_rn(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_ru(double __a);
+__DEVICE__ unsigned long long __nv_double2ull_rz(double __a);
+__DEVICE__ unsigned long long __nv_double_as_longlong(double __a);
+__DEVICE__ double __nv_drcp_rd(double __a);
+__DEVICE__ double __nv_drcp_rn(double __a);
+__DEVICE__ double __nv_drcp_ru(double __a);
+__DEVICE__ double __nv_drcp_rz(double __a);
+__DEVICE__ double __nv_dsqrt_rd(double __a);
+__DEVICE__ double __nv_dsqrt_rn(double __a);
+__DEVICE__ double __nv_dsqrt_ru(double __a);
+__DEVICE__ double __nv_dsqrt_rz(double __a);
+__DEVICE__ double __nv_dsub_rd(double __a, double __b);
+__DEVICE__ double __nv_dsub_rn(double __a, double __b);
+__DEVICE__ double __nv_dsub_ru(double __a, double __b);
+__DEVICE__ double __nv_dsub_rz(double __a, double __b);
+__DEVICE__ double __nv_erfc(double __a);
+__DEVICE__ float __nv_erfcf(float __a);
+__DEVICE__ double __nv_erfcinv(double __a);
+__DEVICE__ float __nv_erfcinvf(float __a);
+__DEVICE__ double __nv_erfcx(double __a);
+__DEVICE__ float __nv_erfcxf(float __a);
+__DEVICE__ double __nv_erf(double __a);
+__DEVICE__ float __nv_erff(float __a);
+__DEVICE__ double __nv_erfinv(double __a);
+__DEVICE__ float __nv_erfinvf(float __a);
+__DEVICE__ double __nv_exp10(double __a);
+__DEVICE__ float __nv_exp10f(float __a);
+__DEVICE__ double __nv_exp2(double __a);
+__DEVICE__ float __nv_exp2f(float __a);
+__DEVICE__ double __nv_exp(double __a);
+__DEVICE__ float __nv_expf(float __a);
+__DEVICE__ double __nv_expm1(double __a);
+__DEVICE__ float __nv_expm1f(float __a);
+__DEVICE__ double __nv_fabs(double __a);
+__DEVICE__ float __nv_fabsf(float __a);
+__DEVICE__ float __nv_fadd_rd(float __a, float __b);
+__DEVICE__ float __nv_fadd_rn(float __a, float __b);
+__DEVICE__ float __nv_fadd_ru(float __a, float __b);
+__DEVICE__ float __nv_fadd_rz(float __a, float __b);
+__DEVICE__ float __nv_fast_cosf(float __a);
+__DEVICE__ float __nv_fast_exp10f(float __a);
+__DEVICE__ float __nv_fast_expf(float __a);
+__DEVICE__ float __nv_fast_fdividef(float __a, float __b);
+__DEVICE__ float __nv_fast_log10f(float __a);
+__DEVICE__ float __nv_fast_log2f(float __a);
+__DEVICE__ float __nv_fast_logf(float __a);
+__DEVICE__ float __nv_fast_powf(float __a, float __b);
+__DEVICE__ void __nv_fast_sincosf(float __a, float *__sptr, float *__cptr);
+__DEVICE__ float __nv_fast_sinf(float __a);
+__DEVICE__ float __nv_fast_tanf(float __a);
+__DEVICE__ double __nv_fdim(double __a, double __b);
+__DEVICE__ float __nv_fdimf(float __a, float __b);
+__DEVICE__ float __nv_fdiv_rd(float __a, float __b);
+__DEVICE__ float __nv_fdiv_rn(float __a, float __b);
+__DEVICE__ float __nv_fdiv_ru(float __a, float __b);
+__DEVICE__ float __nv_fdiv_rz(float __a, float __b);
+__DEVICE__ int __nv_ffs(int __a);
+__DEVICE__ int __nv_ffsll(long long __a);
+__DEVICE__ int __nv_finitef(float __a);
+__DEVICE__ unsigned short __nv_float2half_rn(float __a);
+__DEVICE__ int __nv_float2int_rd(float __a);
+__DEVICE__ int __nv_float2int_rn(float __a);
+__DEVICE__ int __nv_float2int_ru(float __a);
+__DEVICE__ int __nv_float2int_rz(float __a);
+__DEVICE__ long long __nv_float2ll_rd(float __a);
+__DEVICE__ long long __nv_float2ll_rn(float __a);
+__DEVICE__ long long __nv_float2ll_ru(float __a);
+__DEVICE__ long long __nv_float2ll_rz(float __a);
+__DEVICE__ unsigned int __nv_float2uint_rd(float __a);
+__DEVICE__ unsigned int __nv_float2uint_rn(float __a);
+__DEVICE__ unsigned int __nv_float2uint_ru(float __a);
+__DEVICE__ unsigned int __nv_float2uint_rz(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_rd(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_rn(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_ru(float __a);
+__DEVICE__ unsigned long long __nv_float2ull_rz(float __a);
+__DEVICE__ int __nv_float_as_int(float __a);
+__DEVICE__ unsigned int __nv_float_as_uint(float __a);
+__DEVICE__ double __nv_floor(double __a);
+__DEVICE__ float __nv_floorf(float __a);
+__DEVICE__ double __nv_fma(double __a, double __b, double __c);
+__DEVICE__ float __nv_fmaf(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_rd(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_rn(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_ru(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ieee_rz(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_rd(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_rn(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_ru(float __a, float __b, float __c);
+__DEVICE__ float __nv_fmaf_rz(float __a, float __b, float __c);
+__DEVICE__ double __nv_fma_rd(double __a, double __b, double __c);
+__DEVICE__ double __nv_fma_rn(double __a, double __b, double __c);
+__DEVICE__ double __nv_fma_ru(double __a, double __b, double __c);
+__DEVICE__ double __nv_fma_rz(double __a, double __b, double __c);
+__DEVICE__ double __nv_fmax(double __a, double __b);
+__DEVICE__ float __nv_fmaxf(float __a, float __b);
+__DEVICE__ double __nv_fmin(double __a, double __b);
+__DEVICE__ float __nv_fminf(float __a, float __b);
+__DEVICE__ double __nv_fmod(double __a, double __b);
+__DEVICE__ float __nv_fmodf(float __a, float __b);
+__DEVICE__ float __nv_fmul_rd(float __a, float __b);
+__DEVICE__ float __nv_fmul_rn(float __a, float __b);
+__DEVICE__ float __nv_fmul_ru(float __a, float __b);
+__DEVICE__ float __nv_fmul_rz(float __a, float __b);
+__DEVICE__ float __nv_frcp_rd(float __a);
+__DEVICE__ float __nv_frcp_rn(float __a);
+__DEVICE__ float __nv_frcp_ru(float __a);
+__DEVICE__ float __nv_frcp_rz(float __a);
+__DEVICE__ double __nv_frexp(double __a, int *__b);
+__DEVICE__ float __nv_frexpf(float __a, int *__b);
+__DEVICE__ float __nv_frsqrt_rn(float __a);
+__DEVICE__ float __nv_fsqrt_rd(float __a);
+__DEVICE__ float __nv_fsqrt_rn(float __a);
+__DEVICE__ float __nv_fsqrt_ru(float __a);
+__DEVICE__ float __nv_fsqrt_rz(float __a);
+__DEVICE__ float __nv_fsub_rd(float __a, float __b);
+__DEVICE__ float __nv_fsub_rn(float __a, float __b);
+__DEVICE__ float __nv_fsub_ru(float __a, float __b);
+__DEVICE__ float __nv_fsub_rz(float __a, float __b);
+__DEVICE__ int __nv_hadd(int __a, int __b);
+__DEVICE__ float __nv_half2float(unsigned short __h);
+__DEVICE__ double __nv_hiloint2double(int __a, int __b);
+__DEVICE__ double __nv_hypot(double __a, double __b);
+__DEVICE__ float __nv_hypotf(float __a, float __b);
+__DEVICE__ int __nv_ilogb(double __a);
+__DEVICE__ int __nv_ilogbf(float __a);
+__DEVICE__ double __nv_int2double_rn(int __a);
+__DEVICE__ float __nv_int2float_rd(int __a);
+__DEVICE__ float __nv_int2float_rn(int __a);
+__DEVICE__ float __nv_int2float_ru(int __a);
+__DEVICE__ float __nv_int2float_rz(int __a);
+__DEVICE__ float __nv_int_as_float(int __a);
+__DEVICE__ int __nv_isfinited(double __a);
+__DEVICE__ int __nv_isinfd(double __a);
+__DEVICE__ int __nv_isinff(float __a);
+__DEVICE__ int __nv_isnand(double __a);
+__DEVICE__ int __nv_isnanf(float __a);
+__DEVICE__ double __nv_j0(double __a);
+__DEVICE__ float __nv_j0f(float __a);
+__DEVICE__ double __nv_j1(double __a);
+__DEVICE__ float __nv_j1f(float __a);
+__DEVICE__ float __nv_jnf(int __a, float __b);
+__DEVICE__ double __nv_jn(int __a, double __b);
+__DEVICE__ double __nv_ldexp(double __a, int __b);
+__DEVICE__ float __nv_ldexpf(float __a, int __b);
+__DEVICE__ double __nv_lgamma(double __a);
+__DEVICE__ float __nv_lgammaf(float __a);
+__DEVICE__ double __nv_ll2double_rd(long long __a);
+__DEVICE__ double __nv_ll2double_rn(long long __a);
+__DEVICE__ double __nv_ll2double_ru(long long __a);
+__DEVICE__ double __nv_ll2double_rz(long long __a);
+__DEVICE__ float __nv_ll2float_rd(long long __a);
+__DEVICE__ float __nv_ll2float_rn(long long __a);
+__DEVICE__ float __nv_ll2float_ru(long long __a);
+__DEVICE__ float __nv_ll2float_rz(long long __a);
+__DEVICE__ long long __nv_llabs(long long __a);
+__DEVICE__ long long __nv_llmax(long long __a, long long __b);
+__DEVICE__ long long __nv_llmin(long long __a, long long __b);
+__DEVICE__ long long __nv_llrint(double __a);
+__DEVICE__ long long __nv_llrintf(float __a);
+__DEVICE__ long long __nv_llround(double __a);
+__DEVICE__ long long __nv_llroundf(float __a);
+__DEVICE__ double __nv_log10(double __a);
+__DEVICE__ float __nv_log10f(float __a);
+__DEVICE__ double __nv_log1p(double __a);
+__DEVICE__ float __nv_log1pf(float __a);
+__DEVICE__ double __nv_log2(double __a);
+__DEVICE__ float __nv_log2f(float __a);
+__DEVICE__ double __nv_logb(double __a);
+__DEVICE__ float __nv_logbf(float __a);
+__DEVICE__ double __nv_log(double __a);
+__DEVICE__ float __nv_logf(float __a);
+__DEVICE__ double __nv_longlong_as_double(long long __a);
+__DEVICE__ int __nv_max(int __a, int __b);
+__DEVICE__ int __nv_min(int __a, int __b);
+__DEVICE__ double __nv_modf(double __a, double *__b);
+__DEVICE__ float __nv_modff(float __a, float *__b);
+__DEVICE__ int __nv_mul24(int __a, int __b);
+__DEVICE__ long long __nv_mul64hi(long long __a, long long __b);
+__DEVICE__ int __nv_mulhi(int __a, int __b);
+__DEVICE__ double __nv_nan(const signed char *__a);
+__DEVICE__ float __nv_nanf(const signed char *__a);
+__DEVICE__ double __nv_nearbyint(double __a);
+__DEVICE__ float __nv_nearbyintf(float __a);
+__DEVICE__ double __nv_nextafter(double __a, double __b);
+__DEVICE__ float __nv_nextafterf(float __a, float __b);
+__DEVICE__ double __nv_norm3d(double __a, double __b, double __c);
+__DEVICE__ float __nv_norm3df(float __a, float __b, float __c);
+__DEVICE__ double __nv_norm4d(double __a, double __b, double __c, double __d);
+__DEVICE__ float __nv_norm4df(float __a, float __b, float __c, float __d);
+__DEVICE__ double __nv_normcdf(double __a);
+__DEVICE__ float __nv_normcdff(float __a);
+__DEVICE__ double __nv_normcdfinv(double __a);
+__DEVICE__ float __nv_normcdfinvf(float __a);
+__DEVICE__ float __nv_normf(int __a, const float *__b);
+__DEVICE__ double __nv_norm(int __a, const double *__b);
+__DEVICE__ int __nv_popc(int __a);
+__DEVICE__ int __nv_popcll(long long __a);
+__DEVICE__ double __nv_pow(double __a, double __b);
+__DEVICE__ float __nv_powf(float __a, float __b);
+__DEVICE__ double __nv_powi(double __a, int __b);
+__DEVICE__ float __nv_powif(float __a, int __b);
+__DEVICE__ double __nv_rcbrt(double __a);
+__DEVICE__ float __nv_rcbrtf(float __a);
+__DEVICE__ double __nv_rcp64h(double __a);
+__DEVICE__ double __nv_remainder(double __a, double __b);
+__DEVICE__ float __nv_remainderf(float __a, float __b);
+__DEVICE__ double __nv_remquo(double __a, double __b, int *__c);
+__DEVICE__ float __nv_remquof(float __a, float __b, int *__c);
+__DEVICE__ int __nv_rhadd(int __a, int __b);
+__DEVICE__ double __nv_rhypot(double __a, double __b);
+__DEVICE__ float __nv_rhypotf(float __a, float __b);
+__DEVICE__ double __nv_rint(double __a);
+__DEVICE__ float __nv_rintf(float __a);
+__DEVICE__ double __nv_rnorm3d(double __a, double __b, double __c);
+__DEVICE__ float __nv_rnorm3df(float __a, float __b, float __c);
+__DEVICE__ double __nv_rnorm4d(double __a, double __b, double __c, double __d);
+__DEVICE__ float __nv_rnorm4df(float __a, float __b, float __c, float __d);
+__DEVICE__ float __nv_rnormf(int __a, const float *__b);
+__DEVICE__ double __nv_rnorm(int __a, const double *__b);
+__DEVICE__ double __nv_round(double __a);
+__DEVICE__ float __nv_roundf(float __a);
+__DEVICE__ double __nv_rsqrt(double __a);
+__DEVICE__ float __nv_rsqrtf(float __a);
+__DEVICE__ int __nv_sad(int __a, int __b, int __c);
+__DEVICE__ float __nv_saturatef(float __a);
+__DEVICE__ double __nv_scalbn(double __a, int __b);
+__DEVICE__ float __nv_scalbnf(float __a, int __b);
+__DEVICE__ int __nv_signbitd(double __a);
+__DEVICE__ int __nv_signbitf(float __a);
+__DEVICE__ void __nv_sincos(double __a, double *__b, double *__c);
+__DEVICE__ void __nv_sincosf(float __a, float *__b, float *__c);
+__DEVICE__ void __nv_sincospi(double __a, double *__b, double *__c);
+__DEVICE__ void __nv_sincospif(float __a, float *__b, float *__c);
+__DEVICE__ double __nv_sin(double __a);
+__DEVICE__ float __nv_sinf(float __a);
+__DEVICE__ double __nv_sinh(double __a);
+__DEVICE__ float __nv_sinhf(float __a);
+__DEVICE__ double __nv_sinpi(double __a);
+__DEVICE__ float __nv_sinpif(float __a);
+__DEVICE__ double __nv_sqrt(double __a);
+__DEVICE__ float __nv_sqrtf(float __a);
+__DEVICE__ double __nv_tan(double __a);
+__DEVICE__ float __nv_tanf(float __a);
+__DEVICE__ double __nv_tanh(double __a);
+__DEVICE__ float __nv_tanhf(float __a);
+__DEVICE__ double __nv_tgamma(double __a);
+__DEVICE__ float __nv_tgammaf(float __a);
+__DEVICE__ double __nv_trunc(double __a);
+__DEVICE__ float __nv_truncf(float __a);
+__DEVICE__ int __nv_uhadd(unsigned int __a, unsigned int __b);
+__DEVICE__ double __nv_uint2double_rn(unsigned int __i);
+__DEVICE__ float __nv_uint2float_rd(unsigned int __a);
+__DEVICE__ float __nv_uint2float_rn(unsigned int __a);
+__DEVICE__ float __nv_uint2float_ru(unsigned int __a);
+__DEVICE__ float __nv_uint2float_rz(unsigned int __a);
+__DEVICE__ float __nv_uint_as_float(unsigned int __a);
+__DEVICE__ double __nv_ull2double_rd(unsigned long long __a);
+__DEVICE__ double __nv_ull2double_rn(unsigned long long __a);
+__DEVICE__ double __nv_ull2double_ru(unsigned long long __a);
+__DEVICE__ double __nv_ull2double_rz(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_rd(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_rn(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_ru(unsigned long long __a);
+__DEVICE__ float __nv_ull2float_rz(unsigned long long __a);
+__DEVICE__ unsigned long long __nv_ullmax(unsigned long long __a,
unsigned long long __b);
-__device__ unsigned long long __nv_ullmin(unsigned long long __a,
+__DEVICE__ unsigned long long __nv_ullmin(unsigned long long __a,
unsigned long long __b);
-__device__ unsigned int __nv_umax(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_umin(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_umul24(unsigned int __a, unsigned int __b);
-__device__ unsigned long long __nv_umul64hi(unsigned long long __a,
+__DEVICE__ unsigned int __nv_umax(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_umin(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_umul24(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned long long __nv_umul64hi(unsigned long long __a,
unsigned long long __b);
-__device__ unsigned int __nv_umulhi(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_urhadd(unsigned int __a, unsigned int __b);
-__device__ unsigned int __nv_usad(unsigned int __a, unsigned int __b,
+__DEVICE__ unsigned int __nv_umulhi(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_urhadd(unsigned int __a, unsigned int __b);
+__DEVICE__ unsigned int __nv_usad(unsigned int __a, unsigned int __b,
unsigned int __c);
+
#if CUDA_VERSION >= 9000 && CUDA_VERSION < 9020
-__device__ int __nv_vabs2(int __a);
-__device__ int __nv_vabs4(int __a);
-__device__ int __nv_vabsdiffs2(int __a, int __b);
-__device__ int __nv_vabsdiffs4(int __a, int __b);
-__device__ int __nv_vabsdiffu2(int __a, int __b);
-__device__ int __nv_vabsdiffu4(int __a, int __b);
-__device__ int __nv_vabsss2(int __a);
-__device__ int __nv_vabsss4(int __a);
-__device__ int __nv_vadd2(int __a, int __b);
-__device__ int __nv_vadd4(int __a, int __b);
-__device__ int __nv_vaddss2(int __a, int __b);
-__device__ int __nv_vaddss4(int __a, int __b);
-__device__ int __nv_vaddus2(int __a, int __b);
-__device__ int __nv_vaddus4(int __a, int __b);
-__device__ int __nv_vavgs2(int __a, int __b);
-__device__ int __nv_vavgs4(int __a, int __b);
-__device__ int __nv_vavgu2(int __a, int __b);
-__device__ int __nv_vavgu4(int __a, int __b);
-__device__ int __nv_vcmpeq2(int __a, int __b);
-__device__ int __nv_vcmpeq4(int __a, int __b);
-__device__ int __nv_vcmpges2(int __a, int __b);
-__device__ int __nv_vcmpges4(int __a, int __b);
-__device__ int __nv_vcmpgeu2(int __a, int __b);
-__device__ int __nv_vcmpgeu4(int __a, int __b);
-__device__ int __nv_vcmpgts2(int __a, int __b);
-__device__ int __nv_vcmpgts4(int __a, int __b);
-__device__ int __nv_vcmpgtu2(int __a, int __b);
-__device__ int __nv_vcmpgtu4(int __a, int __b);
-__device__ int __nv_vcmples2(int __a, int __b);
-__device__ int __nv_vcmples4(int __a, int __b);
-__device__ int __nv_vcmpleu2(int __a, int __b);
-__device__ int __nv_vcmpleu4(int __a, int __b);
-__device__ int __nv_vcmplts2(int __a, int __b);
-__device__ int __nv_vcmplts4(int __a, int __b);
-__device__ int __nv_vcmpltu2(int __a, int __b);
-__device__ int __nv_vcmpltu4(int __a, int __b);
-__device__ int __nv_vcmpne2(int __a, int __b);
-__device__ int __nv_vcmpne4(int __a, int __b);
-__device__ int __nv_vhaddu2(int __a, int __b);
-__device__ int __nv_vhaddu4(int __a, int __b);
-__device__ int __nv_vmaxs2(int __a, int __b);
-__device__ int __nv_vmaxs4(int __a, int __b);
-__device__ int __nv_vmaxu2(int __a, int __b);
-__device__ int __nv_vmaxu4(int __a, int __b);
-__device__ int __nv_vmins2(int __a, int __b);
-__device__ int __nv_vmins4(int __a, int __b);
-__device__ int __nv_vminu2(int __a, int __b);
-__device__ int __nv_vminu4(int __a, int __b);
-__device__ int __nv_vneg2(int __a);
-__device__ int __nv_vneg4(int __a);
-__device__ int __nv_vnegss2(int __a);
-__device__ int __nv_vnegss4(int __a);
-__device__ int __nv_vsads2(int __a, int __b);
-__device__ int __nv_vsads4(int __a, int __b);
-__device__ int __nv_vsadu2(int __a, int __b);
-__device__ int __nv_vsadu4(int __a, int __b);
-__device__ int __nv_vseteq2(int __a, int __b);
-__device__ int __nv_vseteq4(int __a, int __b);
-__device__ int __nv_vsetges2(int __a, int __b);
-__device__ int __nv_vsetges4(int __a, int __b);
-__device__ int __nv_vsetgeu2(int __a, int __b);
-__device__ int __nv_vsetgeu4(int __a, int __b);
-__device__ int __nv_vsetgts2(int __a, int __b);
-__device__ int __nv_vsetgts4(int __a, int __b);
-__device__ int __nv_vsetgtu2(int __a, int __b);
-__device__ int __nv_vsetgtu4(int __a, int __b);
-__device__ int __nv_vsetles2(int __a, int __b);
-__device__ int __nv_vsetles4(int __a, int __b);
-__device__ int __nv_vsetleu2(int __a, int __b);
-__device__ int __nv_vsetleu4(int __a, int __b);
-__device__ int __nv_vsetlts2(int __a, int __b);
-__device__ int __nv_vsetlts4(int __a, int __b);
-__device__ int __nv_vsetltu2(int __a, int __b);
-__device__ int __nv_vsetltu4(int __a, int __b);
-__device__ int __nv_vsetne2(int __a, int __b);
-__device__ int __nv_vsetne4(int __a, int __b);
-__device__ int __nv_vsub2(int __a, int __b);
-__device__ int __nv_vsub4(int __a, int __b);
-__device__ int __nv_vsubss2(int __a, int __b);
-__device__ int __nv_vsubss4(int __a, int __b);
-__device__ int __nv_vsubus2(int __a, int __b);
-__device__ int __nv_vsubus4(int __a, int __b);
+__DEVICE__ int __nv_vabs2(int __a);
+__DEVICE__ int __nv_vabs4(int __a);
+__DEVICE__ int __nv_vabsdiffs2(int __a, int __b);
+__DEVICE__ int __nv_vabsdiffs4(int __a, int __b);
+__DEVICE__ int __nv_vabsdiffu2(int __a, int __b);
+__DEVICE__ int __nv_vabsdiffu4(int __a, int __b);
+__DEVICE__ int __nv_vabsss2(int __a);
+__DEVICE__ int __nv_vabsss4(int __a);
+__DEVICE__ int __nv_vadd2(int __a, int __b);
+__DEVICE__ int __nv_vadd4(int __a, int __b);
+__DEVICE__ int __nv_vaddss2(int __a, int __b);
+__DEVICE__ int __nv_vaddss4(int __a, int __b);
+__DEVICE__ int __nv_vaddus2(int __a, int __b);
+__DEVICE__ int __nv_vaddus4(int __a, int __b);
+__DEVICE__ int __nv_vavgs2(int __a, int __b);
+__DEVICE__ int __nv_vavgs4(int __a, int __b);
+__DEVICE__ int __nv_vavgu2(int __a, int __b);
+__DEVICE__ int __nv_vavgu4(int __a, int __b);
+__DEVICE__ int __nv_vcmpeq2(int __a, int __b);
+__DEVICE__ int __nv_vcmpeq4(int __a, int __b);
+__DEVICE__ int __nv_vcmpges2(int __a, int __b);
+__DEVICE__ int __nv_vcmpges4(int __a, int __b);
+__DEVICE__ int __nv_vcmpgeu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpgeu4(int __a, int __b);
+__DEVICE__ int __nv_vcmpgts2(int __a, int __b);
+__DEVICE__ int __nv_vcmpgts4(int __a, int __b);
+__DEVICE__ int __nv_vcmpgtu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpgtu4(int __a, int __b);
+__DEVICE__ int __nv_vcmples2(int __a, int __b);
+__DEVICE__ int __nv_vcmples4(int __a, int __b);
+__DEVICE__ int __nv_vcmpleu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpleu4(int __a, int __b);
+__DEVICE__ int __nv_vcmplts2(int __a, int __b);
+__DEVICE__ int __nv_vcmplts4(int __a, int __b);
+__DEVICE__ int __nv_vcmpltu2(int __a, int __b);
+__DEVICE__ int __nv_vcmpltu4(int __a, int __b);
+__DEVICE__ int __nv_vcmpne2(int __a, int __b);
+__DEVICE__ int __nv_vcmpne4(int __a, int __b);
+__DEVICE__ int __nv_vhaddu2(int __a, int __b);
+__DEVICE__ int __nv_vhaddu4(int __a, int __b);
+__DEVICE__ int __nv_vmaxs2(int __a, int __b);
+__DEVICE__ int __nv_vmaxs4(int __a, int __b);
+__DEVICE__ int __nv_vmaxu2(int __a, int __b);
+__DEVICE__ int __nv_vmaxu4(int __a, int __b);
+__DEVICE__ int __nv_vmins2(int __a, int __b);
+__DEVICE__ int __nv_vmins4(int __a, int __b);
+__DEVICE__ int __nv_vminu2(int __a, int __b);
+__DEVICE__ int __nv_vminu4(int __a, int __b);
+__DEVICE__ int __nv_vneg2(int __a);
+__DEVICE__ int __nv_vneg4(int __a);
+__DEVICE__ int __nv_vnegss2(int __a);
+__DEVICE__ int __nv_vnegss4(int __a);
+__DEVICE__ int __nv_vsads2(int __a, int __b);
+__DEVICE__ int __nv_vsads4(int __a, int __b);
+__DEVICE__ int __nv_vsadu2(int __a, int __b);
+__DEVICE__ int __nv_vsadu4(int __a, int __b);
+__DEVICE__ int __nv_vseteq2(int __a, int __b);
+__DEVICE__ int __nv_vseteq4(int __a, int __b);
+__DEVICE__ int __nv_vsetges2(int __a, int __b);
+__DEVICE__ int __nv_vsetges4(int __a, int __b);
+__DEVICE__ int __nv_vsetgeu2(int __a, int __b);
+__DEVICE__ int __nv_vsetgeu4(int __a, int __b);
+__DEVICE__ int __nv_vsetgts2(int __a, int __b);
+__DEVICE__ int __nv_vsetgts4(int __a, int __b);
+__DEVICE__ int __nv_vsetgtu2(int __a, int __b);
+__DEVICE__ int __nv_vsetgtu4(int __a, int __b);
+__DEVICE__ int __nv_vsetles2(int __a, int __b);
+__DEVICE__ int __nv_vsetles4(int __a, int __b);
+__DEVICE__ int __nv_vsetleu2(int __a, int __b);
+__DEVICE__ int __nv_vsetleu4(int __a, int __b);
+__DEVICE__ int __nv_vsetlts2(int __a, int __b);
+__DEVICE__ int __nv_vsetlts4(int __a, int __b);
+__DEVICE__ int __nv_vsetltu2(int __a, int __b);
+__DEVICE__ int __nv_vsetltu4(int __a, int __b);
+__DEVICE__ int __nv_vsetne2(int __a, int __b);
+__DEVICE__ int __nv_vsetne4(int __a, int __b);
+__DEVICE__ int __nv_vsub2(int __a, int __b);
+__DEVICE__ int __nv_vsub4(int __a, int __b);
+__DEVICE__ int __nv_vsubss2(int __a, int __b);
+__DEVICE__ int __nv_vsubss4(int __a, int __b);
+__DEVICE__ int __nv_vsubus2(int __a, int __b);
+__DEVICE__ int __nv_vsubus4(int __a, int __b);
#endif // CUDA_VERSION
-__device__ double __nv_y0(double __a);
-__device__ float __nv_y0f(float __a);
-__device__ double __nv_y1(double __a);
-__device__ float __nv_y1f(float __a);
-__device__ float __nv_ynf(int __a, float __b);
-__device__ double __nv_yn(int __a, double __b);
+__DEVICE__ double __nv_y0(double __a);
+__DEVICE__ float __nv_y0f(float __a);
+__DEVICE__ double __nv_y1(double __a);
+__DEVICE__ float __nv_y1f(float __a);
+__DEVICE__ float __nv_ynf(int __a, float __b);
+__DEVICE__ double __nv_yn(int __a, double __b);
+
+#if defined(__cplusplus)
} // extern "C"
+#endif
#endif // __CLANG_CUDA_LIBDEVICE_DECLARES_H__
Index: lib/Headers/__clang_cuda_device_functions.h
===================================================================
--- lib/Headers/__clang_cuda_device_functions.h
+++ lib/Headers/__clang_cuda_device_functions.h
@@ -24,15 +24,25 @@
#ifndef __CLANG_CUDA_DEVICE_FUNCTIONS_H__
#define __CLANG_CUDA_DEVICE_FUNCTIONS_H__
-#if CUDA_VERSION < 9000
+#if defined(_OPENMP)
+#include <__clang_cuda_libdevice_declares.h>
+#include <stddef.h>
+#include <limits.h>
+#endif
+
+#if defined(__CUDA__) && CUDA_VERSION < 9000
#error This file is intended to be used with CUDA-9+ only.
#endif
// __DEVICE__ is a helper macro with common set of attributes for the wrappers
// we implement in this file. We need static in order to avoid emitting unused
// functions and __forceinline__ helps inlining these wrappers at -O1.
#pragma push_macro("__DEVICE__")
+#if defined(__CUDA__)
#define __DEVICE__ static __device__ __forceinline__
+#elif defined(_OPENMP)
+#define __DEVICE__ static __attribute__((always_inline))
+#endif
// libdevice provides fast low precision and slow full-recision implementations
// for some functions. Which one gets selected depends on
@@ -53,7 +63,9 @@
return __nv_brevll(__a);
}
__DEVICE__ void __brkpt() { asm volatile("brkpt;"); }
+#if defined(__cplusplus)
__DEVICE__ void __brkpt(int __a) { __brkpt(); }
+#endif
__DEVICE__ unsigned int __byte_perm(unsigned int __a, unsigned int __b,
unsigned int __c) {
return __nv_byte_perm(__a, __b, __c);
@@ -519,7 +531,9 @@
}
__DEVICE__ float __saturatef(float __a) { return __nv_saturatef(__a); }
__DEVICE__ int __signbitd(double __a) { return __nv_signbitd(__a); }
+#if defined(__CUDA__)
__DEVICE__ int __signbitf(float __a) { return __nv_signbitf(__a); }
+#endif
__DEVICE__ void __sincosf(float __a, float *__sptr, float *__cptr) {
return __nv_fast_sincosf(__a, __sptr, __cptr);
}
@@ -1487,7 +1501,9 @@
__DEVICE__ float cbrtf(float __a) { return __nv_cbrtf(__a); }
__DEVICE__ double ceil(double __a) { return __nv_ceil(__a); }
__DEVICE__ float ceilf(float __a) { return __nv_ceilf(__a); }
+#if defined(__CUDA__)
__DEVICE__ int clock() { return __nvvm_read_ptx_sreg_clock(); }
+#endif
__DEVICE__ long long clock64() { return __nvvm_read_ptx_sreg_clock64(); }
__DEVICE__ double copysign(double __a, double __b) {
return __nv_copysign(__a, __b);
@@ -1527,8 +1543,10 @@
__DEVICE__ float expm1f(float __a) { return __nv_expm1f(__a); }
__DEVICE__ double fabs(double __a) { return __nv_fabs(__a); }
__DEVICE__ float fabsf(float __a) { return __nv_fabsf(__a); }
+#if defined(__CUDA__)
__DEVICE__ double fdim(double __a, double __b) { return __nv_fdim(__a, __b); }
__DEVICE__ float fdimf(float __a, float __b) { return __nv_fdimf(__a, __b); }
+#endif
__DEVICE__ double fdivide(double __a, double __b) { return __a / __b; }
__DEVICE__ float fdividef(float __a, float __b) {
#if __FAST_MATH__ && !__CUDA_PREC_DIV
@@ -1564,7 +1582,7 @@
__DEVICE__ double jn(int __n, double __a) { return __nv_jn(__n, __a); }
__DEVICE__ float jnf(int __n, float __a) { return __nv_jnf(__n, __a); }
#if defined(__LP64__)
-__DEVICE__ long labs(long __a) { return llabs(__a); };
+__DEVICE__ long labs(long __a) { return __nv_llabs(__a); };
#else
__DEVICE__ long labs(long __a) { return __nv_abs(__a); };
#endif
@@ -1698,6 +1716,7 @@
__DEVICE__ float rsqrtf(float __a) { return __nv_rsqrtf(__a); }
__DEVICE__ double scalbn(double __a, int __b) { return __nv_scalbn(__a, __b); }
__DEVICE__ float scalbnf(float __a, int __b) { return __nv_scalbnf(__a, __b); }
+#if defined(__CUDA__)
__DEVICE__ double scalbln(double __a, long __b) {
if (__b > INT_MAX)
return __a > 0 ? HUGE_VAL : -HUGE_VAL;
@@ -1712,6 +1731,7 @@
return __a > 0 ? 0.f : -0.f;
return scalbnf(__a, (int)__b);
}
+#endif
__DEVICE__ double sin(double __a) { return __nv_sin(__a); }
__DEVICE__ void sincos(double __a, double *__sptr, double *__cptr) {
return __nv_sincos(__a, __sptr, __cptr);
Index: lib/Headers/CMakeLists.txt
===================================================================
--- lib/Headers/CMakeLists.txt
+++ lib/Headers/CMakeLists.txt
@@ -144,7 +144,7 @@
list(APPEND out_files ${dst})
endforeach( f )
-add_custom_command(OUTPUT ${output_dir}/arm_neon.h
+add_custom_command(OUTPUT ${output_dir}/arm_neon.h
DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/arm_neon.h
COMMAND ${CMAKE_COMMAND} -E copy_if_different ${CMAKE_CURRENT_BINARY_DIR}/arm_neon.h ${output_dir}/arm_neon.h
COMMENT "Copying clang's arm_neon.h...")
Index: lib/Driver/ToolChains/Cuda.h
===================================================================
--- lib/Driver/ToolChains/Cuda.h
+++ lib/Driver/ToolChains/Cuda.h
@@ -46,6 +46,9 @@
CudaInstallationDetector(const Driver &D, const llvm::Triple &HostTriple,
const llvm::opt::ArgList &Args);
+ void AddMathDeviceFunctions(const llvm::opt::ArgList &DriverArgs,
+ llvm::opt::ArgStringList &CC1Args) const;
+
void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
llvm::opt::ArgStringList &CC1Args) const;
@@ -164,6 +167,9 @@
void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
llvm::opt::ArgStringList &CC1Args) const override;
+ void AddMathDeviceFunctions(const llvm::opt::ArgList &DriverArgs,
+ llvm::opt::ArgStringList &CC1Args) const override;
+
void addClangWarningOptions(llvm::opt::ArgStringList &CC1Args) const override;
CXXStdlibType GetCXXStdlibType(const llvm::opt::ArgList &Args) const override;
void
Index: lib/Driver/ToolChains/Cuda.cpp
===================================================================
--- lib/Driver/ToolChains/Cuda.cpp
+++ lib/Driver/ToolChains/Cuda.cpp
@@ -226,6 +226,16 @@
}
}
+void CudaInstallationDetector::AddMathDeviceFunctions(
+ const ArgList &DriverArgs, ArgStringList &CC1Args) const {
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(getIncludePath()));
+ CC1Args.push_back("-include");
+ CC1Args.push_back("__clang_cuda_device_functions.h");
+ CC1Args.push_back("-I");
+ CC1Args.push_back(DriverArgs.MakeArgString(getIncludePath()));
+}
+
void CudaInstallationDetector::AddCudaIncludeArgs(
const ArgList &DriverArgs, ArgStringList &CC1Args) const {
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
@@ -853,6 +863,11 @@
CudaInstallation.AddCudaIncludeArgs(DriverArgs, CC1Args);
}
+void CudaToolChain::AddMathDeviceFunctions(
+ const ArgList &DriverArgs, ArgStringList &CC1Args) const {
+ CudaInstallation.AddMathDeviceFunctions(DriverArgs, CC1Args);
+}
+
llvm::opt::DerivedArgList *
CudaToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
StringRef BoundArch,
Index: lib/Driver/ToolChains/Clang.cpp
===================================================================
--- lib/Driver/ToolChains/Clang.cpp
+++ lib/Driver/ToolChains/Clang.cpp
@@ -1083,6 +1083,14 @@
if (JA.isOffloading(Action::OFK_Cuda))
getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
+ // If we are offloading to a target via OpenMP and this target happens
+ // to be an NVIDIA GPU then we need to include the CUDA runtime wrapper
+ // to ensure the correct math functions are called in the offloaded
+ // code.
+ if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
+ getToolChain().getTriple().isNVPTX())
+ getToolChain().AddMathDeviceFunctions(Args, CmdArgs);
+
// Add -i* options, and automatically translate to
// -include-pch/-include-pth for transparent PCH support. It's
// wonky, but we include looking for .gch so we can support seamless
@@ -4745,6 +4753,9 @@
CmdArgs.push_back("-fopenmp-host-ir-file-path");
CmdArgs.push_back(Args.MakeArgString(Inputs.back().getFilename()));
}
+ // Prevent usage of math.h builtins for device
+ // toolchain.
+ CmdArgs.push_back("-fno-math-builtin");
}
// For all the host OpenMP offloading compile jobs we need to pass the targets
Index: lib/Basic/Targets/NVPTX.cpp
===================================================================
--- lib/Basic/Targets/NVPTX.cpp
+++ lib/Basic/Targets/NVPTX.cpp
@@ -226,6 +226,10 @@
}();
Builder.defineMacro("__CUDA_ARCH__", CUDAArchCode);
}
+
+ // Prevent any host specific assembly instructions from
+ // getting inlined on the device.
+ Builder.defineMacro("__NO_MATH_INLINES");
}
ArrayRef<Builtin::Info> NVPTXTargetInfo::getTargetBuiltins() const {
Index: include/clang/Driver/ToolChain.h
===================================================================
--- include/clang/Driver/ToolChain.h
+++ include/clang/Driver/ToolChain.h
@@ -526,6 +526,10 @@
virtual void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs,
llvm::opt::ArgStringList &CC1Args) const;
+ /// Add arguments to use system-specific CUDA includes.
+ virtual void AddMathDeviceFunctions(const llvm::opt::ArgList &DriverArgs,
+ llvm::opt::ArgStringList &CC1Args) const {};
+
/// Add arguments to use MCU GCC toolchain includes.
virtual void AddIAMCUIncludeArgs(const llvm::opt::ArgList &DriverArgs,
llvm::opt::ArgStringList &CC1Args) const;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
