Author: gbercea Date: Wed May 15 13:18:21 2019 New Revision: 360804 URL: http://llvm.org/viewvc/llvm-project?rev=360804&view=rev Log: [OpenMP][bugfix] Fix issues with C++ 17 compilation when handling math functions
Summary: In OpenMP device offloading we must ensure that unde C++ 17, the inclusion of cstdlib will works correctly. Reviewers: ABataev, tra, jdoerfert, hfinkel, caomhin Reviewed By: jdoerfert Subscribers: Hahnfeld, guansong, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D61949 Added: cfe/trunk/test/Headers/nvptx_device_cmath_functions_cxx17.cpp cfe/trunk/test/Headers/nvptx_device_math_functions_cxx17.cpp Modified: cfe/trunk/lib/Headers/__clang_cuda_cmath.h cfe/trunk/lib/Headers/__clang_cuda_device_functions.h cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h cfe/trunk/test/Headers/Inputs/include/cstdlib Modified: cfe/trunk/lib/Headers/__clang_cuda_cmath.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_cmath.h?rev=360804&r1=360803&r2=360804&view=diff ============================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_cmath.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_cmath.h Wed May 15 13:18:21 2019 @@ -36,6 +36,15 @@ #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline)) #endif +// For C++ 17 we need to include noexcept attribute to be compatible +// with the header-defined version. This may be removed once +// variant is supported. +#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L +#define __NOEXCEPT noexcept +#else +#define __NOEXCEPT +#endif + #if !(defined(_OPENMP) && defined(__cplusplus)) __DEVICE__ long long abs(long long __n) { return ::llabs(__n); } __DEVICE__ long abs(long __n) { return ::labs(__n); } @@ -50,7 +59,7 @@ __DEVICE__ float ceil(float __x) { retur __DEVICE__ float cos(float __x) { return ::cosf(__x); } __DEVICE__ float cosh(float __x) { return ::coshf(__x); } __DEVICE__ float exp(float __x) { return ::expf(__x); } -__DEVICE__ float fabs(float __x) { return ::fabsf(__x); } +__DEVICE__ float fabs(float __x) __NOEXCEPT { return ::fabsf(__x); } __DEVICE__ float floor(float __x) { return ::floorf(__x); } __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); } // TODO: remove when variant is supported @@ -465,6 +474,7 @@ _GLIBCXX_END_NAMESPACE_VERSION } // namespace std #endif +#undef __NOEXCEPT #undef __DEVICE__ #endif Modified: cfe/trunk/lib/Headers/__clang_cuda_device_functions.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_device_functions.h?rev=360804&r1=360803&r2=360804&view=diff ============================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_device_functions.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_device_functions.h Wed May 15 13:18:21 2019 @@ -37,6 +37,15 @@ #define __FAST_OR_SLOW(fast, slow) slow #endif +// For C++ 17 we need to include noexcept attribute to be compatible +// with the header-defined version. This may be removed once +// variant is supported. +#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L +#define __NOEXCEPT noexcept +#else +#define __NOEXCEPT +#endif + __DEVICE__ int __all(int __a) { return __nvvm_vote_all(__a); } __DEVICE__ int __any(int __a) { return __nvvm_vote_any(__a); } __DEVICE__ unsigned int __ballot(int __a) { return __nvvm_vote_ballot(__a); } @@ -1474,7 +1483,8 @@ __DEVICE__ unsigned int __vsubus4(unsign return r; } #endif // CUDA_VERSION >= 9020 -__DEVICE__ int abs(int __a) { return __nv_abs(__a); } +__DEVICE__ int abs(int __a) __NOEXCEPT { return __nv_abs(__a); } +__DEVICE__ double fabs(double __a) __NOEXCEPT { return __nv_fabs(__a); } __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); } @@ -1533,7 +1543,6 @@ __DEVICE__ float exp2f(float __a) { retu __DEVICE__ float expf(float __a) { return __nv_expf(__a); } __DEVICE__ double expm1(double __a) { return __nv_expm1(__a); } __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); } __DEVICE__ double fdim(double __a, double __b) { return __nv_fdim(__a, __b); } __DEVICE__ float fdimf(float __a, float __b) { return __nv_fdimf(__a, __b); } @@ -1572,15 +1581,15 @@ __DEVICE__ float j1f(float __a) { return __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) -__DEVICE__ long labs(long __a) { return __nv_llabs(__a); }; +__DEVICE__ long labs(long __a) __NOEXCEPT { return __nv_llabs(__a); }; #else -__DEVICE__ long labs(long __a) { return __nv_abs(__a); }; +__DEVICE__ long labs(long __a) __NOEXCEPT { return __nv_abs(__a); }; #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); } -__DEVICE__ long long llabs(long long __a) { return __nv_llabs(__a); } +__DEVICE__ long long llabs(long long __a) __NOEXCEPT { return __nv_llabs(__a); } __DEVICE__ long long llmax(long long __a, long long __b) { return __nv_llmax(__a, __b); } @@ -1778,6 +1787,7 @@ __DEVICE__ float y1f(float __a) { return __DEVICE__ double yn(int __a, double __b) { return __nv_yn(__a, __b); } __DEVICE__ float ynf(int __a, float __b) { return __nv_ynf(__a, __b); } +#undef __NOEXCEPT #pragma pop_macro("__DEVICE__") #pragma pop_macro("__FAST_OR_SLOW") #endif // __CLANG_CUDA_DEVICE_FUNCTIONS_H__ Modified: cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h?rev=360804&r1=360803&r2=360804&view=diff ============================================================================== --- cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h (original) +++ cfe/trunk/lib/Headers/__clang_cuda_math_forward_declares.h Wed May 15 13:18:21 2019 @@ -27,11 +27,20 @@ static __inline__ __attribute__((always_inline)) __attribute__((device)) #endif +// For C++ 17 we need to include noexcept attribute to be compatible +// with the header-defined version. This may be removed once +// variant is supported. +#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L +#define __NOEXCEPT noexcept +#else +#define __NOEXCEPT +#endif + #if !(defined(_OPENMP) && defined(__cplusplus)) __DEVICE__ long abs(long); __DEVICE__ long long abs(long long); #endif -__DEVICE__ int abs(int); +__DEVICE__ int abs(int) __NOEXCEPT; __DEVICE__ double abs(double); __DEVICE__ float abs(float); __DEVICE__ double acos(double); @@ -68,8 +77,8 @@ __DEVICE__ double exp(double); __DEVICE__ float exp(float); __DEVICE__ double expm1(double); __DEVICE__ float expm1(float); -__DEVICE__ double fabs(double); -__DEVICE__ float fabs(float); +__DEVICE__ double fabs(double) __NOEXCEPT; +__DEVICE__ float fabs(float) __NOEXCEPT; __DEVICE__ double fdim(double, double); __DEVICE__ float fdim(float, float); __DEVICE__ double floor(double); @@ -119,12 +128,12 @@ __DEVICE__ bool isnormal(double); __DEVICE__ bool isnormal(float); __DEVICE__ bool isunordered(double, double); __DEVICE__ bool isunordered(float, float); -__DEVICE__ long labs(long); +__DEVICE__ long labs(long) __NOEXCEPT; __DEVICE__ double ldexp(double, int); __DEVICE__ float ldexp(float, int); __DEVICE__ double lgamma(double); __DEVICE__ float lgamma(float); -__DEVICE__ long long llabs(long long); +__DEVICE__ long long llabs(long long) __NOEXCEPT; __DEVICE__ long long llrint(double); __DEVICE__ long long llrint(float); __DEVICE__ double log10(double); @@ -282,6 +291,7 @@ _GLIBCXX_END_NAMESPACE_VERSION } // namespace std #endif +#undef __NOEXCEPT #pragma pop_macro("__DEVICE__") #endif Modified: cfe/trunk/test/Headers/Inputs/include/cstdlib URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Headers/Inputs/include/cstdlib?rev=360804&r1=360803&r2=360804&view=diff ============================================================================== --- cfe/trunk/test/Headers/Inputs/include/cstdlib (original) +++ cfe/trunk/test/Headers/Inputs/include/cstdlib Wed May 15 13:18:21 2019 @@ -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 { Added: cfe/trunk/test/Headers/nvptx_device_cmath_functions_cxx17.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Headers/nvptx_device_cmath_functions_cxx17.cpp?rev=360804&view=auto ============================================================================== --- cfe/trunk/test/Headers/nvptx_device_cmath_functions_cxx17.cpp (added) +++ cfe/trunk/test/Headers/nvptx_device_cmath_functions_cxx17.cpp Wed May 15 13:18:21 2019 @@ -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); + } +} Added: cfe/trunk/test/Headers/nvptx_device_math_functions_cxx17.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Headers/nvptx_device_math_functions_cxx17.cpp?rev=360804&view=auto ============================================================================== --- cfe/trunk/test/Headers/nvptx_device_math_functions_cxx17.cpp (added) +++ cfe/trunk/test/Headers/nvptx_device_math_functions_cxx17.cpp Wed May 15 13:18:21 2019 @@ -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); + } +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
