gtbercea created this revision.
gtbercea added reviewers: ABataev, tra, jdoerfert, hfinkel, caomhin.
Herald added subscribers: cfe-commits, guansong.
Herald added a project: clang.
In OpenMP device offloading we must ensure that unde C++ 17, the inclusion of
cstdlib will works correctly.
Repository:
rC Clang
https://reviews.llvm.org/D61949
Files:
lib/Headers/__clang_cuda_device_functions.h
lib/Headers/__clang_cuda_math_forward_declares.h
lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h
test/Headers/Inputs/include/cstdlib
test/Headers/nvptx_device_cmath_functions_cxx17.cpp
test/Headers/nvptx_device_math_functions_cxx17.cpp
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 <cstdlib>
+#include <math.h>
+
+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_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 <cmath>
+#include <cstdlib>
+
+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/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: 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 <cstdlib>
+ #include <stdlib.h>
#endif
/// Include declarations for libdevice functions.
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/__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);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
