[OpenMP][Bugfix] Move double and float versions of abs under c++ macro

[gtbercea]

Summary

This is a fix for the reported bug:

41861

abs functions need to be moved under the c++ macro to avoid conflicts with included headers.

lib/Headers/__clang_cuda_cmath.h

 #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); }
-#endif
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
+#endif
 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
 __DEVICE__ float exp(float __x) { return ::expf(__x); }

lib/Headers/__clang_cuda_math_forward_declares.h

 #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) __NOEXCEPT;
 __DEVICE__ double abs(double);
 __DEVICE__ float abs(float);
+#endif
+__DEVICE__ int abs(int) __NOEXCEPT;
 __DEVICE__ double acos(double);
 __DEVICE__ float acos(float);
 __DEVICE__ double acosh(double);
 __DEVICE__ float acosh(float);
 __DEVICE__ double asin(double);
 __DEVICE__ float asin(float);
 __DEVICE__ double asinh(double);
 __DEVICE__ float asinh(float);

test/Headers/Inputs/include/cstdlib

 #pragma once
 
 #if __cplusplus >= 201703L
 extern int abs (int __x) throw()  __attribute__ ((__const__)) ;
 extern long int labs (long int __x) throw() __attribute__ ((__const__)) ;
+extern float fabs (float __x) throw() __attribute__ ((__const__)) ;
 #else
 extern int abs (int __x) __attribute__ ((__const__)) ;
 extern long int labs (long int __x) __attribute__ ((__const__)) ;
+extern float fabs (float __x) __attribute__ ((__const__)) ;
 #endif
 
 namespace std
 {
 
 using ::abs;
 
 inline long
 abs(long __i) { return __builtin_labs(__i); }
 
 inline long long
 abs(long long __x) { return __builtin_llabs (__x); }
 }

test/Headers/nvptx_device_cmath_functions.c

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);
+    // CHECK-YES: call double @__nv_fabs(double
+    double l4 = fabs(a1);
+    // CHECK-YES: call i32 @__nv_abs(i32
+    double l5 = abs((int)a1);
   }
 }

test/Headers/nvptx_device_cmath_functions.cpp

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);
+    // CHECK-YES: call double @__nv_fabs(double
+    double l4 = fabs(a1);
+    // CHECK-YES: call i32 @__nv_abs(i32
+    double l5 = abs((int)a1);
   }
 }

test/Headers/nvptx_device_cmath_functions_cxx17.cpp

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);
+    // CHECK-YES: call double @__nv_fabs(double
+    double l4 = fabs(a1);
+    // CHECK-YES: call i32 @__nv_abs(i32
+    double l5 = abs((int)a1);
   }
 }

test/Headers/nvptx_device_math_functions.c

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);
+    // CHECK-YES: call double @__nv_fabs(double
+    double l4 = fabs(a1);
+    // CHECK-YES: call i32 @__nv_abs(i32
+    double l5 = abs((int)a1);
   }
 }

test/Headers/nvptx_device_math_functions.cpp

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);
+    // CHECK-YES: call double @__nv_fabs(double
+    double l4 = fabs(a1);
+    // CHECK-YES: call i32 @__nv_abs(i32
+    double l5 = abs((int)a1);
   }
 }

test/Headers/nvptx_device_math_functions_cxx17.cpp

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);
+    // CHECK-YES: call double @__nv_fabs(double
+    double l4 = fabs(a1);
+    // CHECK-YES: call i32 @__nv_abs(i32
+    double l5 = abs((int)a1);
   }
 }