Index: lib/Headers/__clang_cuda_device_functions.h =================================================================== --- lib/Headers/__clang_cuda_device_functions.h +++ lib/Headers/__clang_cuda_device_functions.h @@ -1474,7 +1474,11 @@ return r; } #endif // CUDA_VERSION >= 9020 +#if __cplusplus >= 201703L +__DEVICE__ int abs(int __a) noexcept { return __nv_abs(__a); } +#else __DEVICE__ int abs(int __a) { return __nv_abs(__a); } +#endif __DEVICE__ double acos(double __a) { return __nv_acos(__a); } __DEVICE__ float acosf(float __a) { return __nv_acosf(__a); } __DEVICE__ double acosh(double __a) { return __nv_acosh(__a); } @@ -1572,15 +1576,27 @@ __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__) || defined(_WIN64) +#if __cplusplus >= 201703L +__DEVICE__ long labs(long __a) noexcept { return __nv_llabs(__a); }; +#else __DEVICE__ long labs(long __a) { return __nv_llabs(__a); }; +#endif +#else +#if __cplusplus >= 201703L +__DEVICE__ long labs(long __a) noexcept { return __nv_abs(__a); }; #else __DEVICE__ long labs(long __a) { return __nv_abs(__a); }; #endif +#endif __DEVICE__ double ldexp(double __a, int __b) { return __nv_ldexp(__a, __b); } __DEVICE__ float ldexpf(float __a, int __b) { return __nv_ldexpf(__a, __b); } __DEVICE__ double lgamma(double __a) { return __nv_lgamma(__a); } __DEVICE__ float lgammaf(float __a) { return __nv_lgammaf(__a); } +#if __cplusplus >= 201703L +__DEVICE__ long long llabs(long long __a) noexcept { return __nv_llabs(__a); } +#else __DEVICE__ long long llabs(long long __a) { return __nv_llabs(__a); } +#endif __DEVICE__ long long llmax(long long __a, long long __b) { return __nv_llmax(__a, __b); } Index: lib/Headers/__clang_cuda_math_forward_declares.h =================================================================== --- lib/Headers/__clang_cuda_math_forward_declares.h +++ lib/Headers/__clang_cuda_math_forward_declares.h @@ -31,7 +31,11 @@ __DEVICE__ long abs(long); __DEVICE__ long long abs(long long); #endif +#if __cplusplus >= 201703L +__DEVICE__ int abs(int) noexcept; +#else __DEVICE__ int abs(int); +#endif __DEVICE__ double abs(double); __DEVICE__ float abs(float); __DEVICE__ double acos(double); @@ -119,12 +123,20 @@ __DEVICE__ bool isnormal(float); __DEVICE__ bool isunordered(double, double); __DEVICE__ bool isunordered(float, float); +#if __cplusplus >= 201703L +__DEVICE__ long labs(long) noexcept; +#else __DEVICE__ long labs(long); +#endif __DEVICE__ double ldexp(double, int); __DEVICE__ float ldexp(float, int); __DEVICE__ double lgamma(double); __DEVICE__ float lgamma(float); +#if __cplusplus >= 201703L +__DEVICE__ long long llabs(long long) noexcept; +#else __DEVICE__ long long llabs(long long); +#endif __DEVICE__ long long llrint(double); __DEVICE__ long long llrint(float); __DEVICE__ double log10(double); Index: lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h =================================================================== --- lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h +++ lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h @@ -20,6 +20,8 @@ #if defined(__cplusplus) #include <__clang_cuda_math_forward_declares.h> + #include + #include #endif /// Include declarations for libdevice functions. Index: test/Headers/Inputs/include/cstdlib =================================================================== --- test/Headers/Inputs/include/cstdlib +++ test/Headers/Inputs/include/cstdlib @@ -1,7 +1,12 @@ #pragma once +#if __cplusplus >= 201703L +extern int abs (int __x) throw() __attribute__ ((__const__)) ; +extern long int labs (long int __x) throw() __attribute__ ((__const__)) ; +#else extern int abs (int __x) __attribute__ ((__const__)) ; extern long int labs (long int __x) __attribute__ ((__const__)) ; +#endif namespace std { Index: test/Headers/nvptx_device_cmath_functions_cxx17.cpp =================================================================== --- /dev/null +++ test/Headers/nvptx_device_cmath_functions_cxx17.cpp @@ -0,0 +1,22 @@ +// Test calling of device math functions. +///==========================================================================/// + +// REQUIRES: nvptx-registered-target + +// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -std=c++17 +// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include cmath -internal-isystem %S/Inputs/include -include stdlib.h -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -std=c++17 -o - | FileCheck -check-prefix CHECK-YES %s + +#include +#include + +void test_sqrt(double a1) { + #pragma omp target + { + // CHECK-YES: call double @__nv_sqrt(double + double l1 = sqrt(a1); + // CHECK-YES: call double @__nv_pow(double + double l2 = pow(a1, a1); + // CHECK-YES: call double @__nv_modf(double + double l3 = modf(a1 + 3.5, &a1); + } +} Index: test/Headers/nvptx_device_math_functions_cxx17.cpp =================================================================== --- /dev/null +++ test/Headers/nvptx_device_math_functions_cxx17.cpp @@ -0,0 +1,22 @@ +// Test calling of device math functions. +///==========================================================================/// + +// REQUIRES: nvptx-registered-target + +// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -std=c++17 +// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_math_declares.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -internal-isystem %S/Inputs/include -include stdlib.h -include limits -include cstdlib -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -std=c++17 -o - | FileCheck -check-prefix CHECK-YES %s + +#include +#include + +void test_sqrt(double a1) { + #pragma omp target + { + // CHECK-YES: call double @__nv_sqrt(double + double l1 = sqrt(a1); + // CHECK-YES: call double @__nv_pow(double + double l2 = pow(a1, a1); + // CHECK-YES: call double @__nv_modf(double + double l3 = modf(a1 + 3.5, &a1); + } +}