[OPENMP][NVPTX]Fix PR45003: add support for complex functions.

[ABataev]

Summary

Inherited support for complex math functions from CUDA implementation.

[ABataev]

Ping

[jdoerfert]

I am unsure we need this with the proper math support. Sema patch for that is going for review today. I'll try this out soon.

[ABataev]

In D75209#1909976, @jdoerfert wrote:

I am unsure we need this with the proper math support. Sema patch for that is going for review today. I'll try this out soon.

It has nothing to do with the math functions support. These functions are required for the definition of muldc3/divdc3/__mulsc2/__divsc3 functions emitted for the complex types. CUDA does absolutely the same.

[jdoerfert]

In D75209#1910047, @ABataev wrote:

In D75209#1909976, @jdoerfert wrote:

I am unsure we need this with the proper math support. Sema patch for that is going for review today. I'll try this out soon.

It has nothing to do with the math functions support. These functions are required for the definition of muldc3/divdc3/__mulsc2/__divsc3 functions emitted for the complex types. CUDA does absolutely the same.

The fact that CUDA does the same is the important part here. Why should we copy it then? See https://reviews.llvm.org/D75788#change-0Ax9LSfw7OAk for working support of complex math functions in target regions without copying what cuda does but by just using it.

clang/lib/Headers/__clang_cuda_complex_builtins.h

 #ifndef __CLANG_CUDA_COMPLEX_BUILTINS
 #define __CLANG_CUDA_COMPLEX_BUILTINS
 
 // This header defines __muldc3, __mulsc3, __divdc3, and __divsc3.  These are
 // libgcc functions that clang assumes are available when compiling c99 complex
 // operations.  (These implementations come from libc++, and have been modified
 // to work with CUDA.)
 
-extern "C" inline __device__ double _Complex __muldc3(double __a, double __b,
-                                                      double __c, double __d) {
+#ifdef _OPENMP
+#ifdef __cplusplus
+#define __DEVICE__ extern "C" inline
+#else
+#define __DEVICE__ __attribute__((always_inline))
+#endif // __cplusplus
+#else
+#define __DEVICE__ extern "C" inline __device__
+#endif // _OPENMP
+
+#ifdef _OPENMP
+#define MAX(x, y) fmax((x), (y))
+#define MAXF(x, y) fmaxf((x), (y))
+#define IS_NAN(x) __isnan(x)
+#define IS_NANF(x) __isnanf(x)
+#define IS_INF(x) __isinf(x)
+#define IS_INFF(x) __isinff(x)
+#define IS_FINITE(x) __finite(x)
+#define IS_FINITEF(x) __finitef(x)
+#define COPYSIGN(x, y) copysign((x), (y))
+#define COPYSIGNF(x, y) copysignf((x), (y))
+#define LOGB(x) logb(x)
+#define LOGBF(x) logbf(x)
+#define ABS(x) fabs(x)
+#define ABSF(x) fabsf(x)
+#define SCALBN(x, y) scalbn((x), (y))
+#define SCALBNF(x, y) scalbnf((x), (y))
+#else
+// Can't use std::max, because that's defined in <algorithm>, and we don't
+// want to pull that in for every compile.  The CUDA headers define
+// ::max(float, float) and ::max(double, double), which is sufficient for us.
+#define MAX(x, y) max((x), (y))
+#define MAXF(x, y) max((x), (y))
+#define IS_NAN(x) std::isnan(x)
+#define IS_NANF(x) std::isnan(x)
+#define IS_INF(x) std::isinf(x)
+#define IS_INFF(x) std::isinf(x)
+#define IS_FINITE(x) std::isfinite(x)
+#define IS_FINITEF(x) std::isfinite(x)
+#define COPYSIGN(x, y) std::copysign((x), (y))
+#define COPYSIGNF(x, y) std::copysign((x), (y))
+#define LOGB(x) std::logb(x)
+#define LOGBF(x) std::logb(x)
+#define ABS(x) std::abs(x)
+#define ABSF(x) std::abs(x)
+#define SCALBN(x, y) std::scalbn((x), (y))
+#define SCALBNF(x, y) std::scalbn((x), (y))
+#endif // _OPENMP
+
+#ifdef _OPENMP
+#pragma omp declare target
+#endif // _OPENMP
+
+__DEVICE__ double _Complex __muldc3(double __a, double __b, double __c,
+                                    double __d) {
   double __ac = __a * __c;
   double __bd = __b * __d;
   double __ad = __a * __d;
   double __bc = __b * __c;
   double _Complex z;
   __real__(z) = __ac - __bd;
   __imag__(z) = __ad + __bc;
-  if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
+  if (IS_NAN(__real__(z)) && IS_NAN(__imag__(z))) {
     int __recalc = 0;
-    if (std::isinf(__a) || std::isinf(__b)) {
-      __a = std::copysign(std::isinf(__a) ? 1 : 0, __a);
-      __b = std::copysign(std::isinf(__b) ? 1 : 0, __b);
-      if (std::isnan(__c))
-        __c = std::copysign(0, __c);
-      if (std::isnan(__d))
-        __d = std::copysign(0, __d);
-      __recalc = 1;
-    }
-    if (std::isinf(__c) || std::isinf(__d)) {
-      __c = std::copysign(std::isinf(__c) ? 1 : 0, __c);
-      __d = std::copysign(std::isinf(__d) ? 1 : 0, __d);
-      if (std::isnan(__a))
-        __a = std::copysign(0, __a);
-      if (std::isnan(__b))
-        __b = std::copysign(0, __b);
-      __recalc = 1;
-    }
-    if (!__recalc && (std::isinf(__ac) || std::isinf(__bd) ||
-                      std::isinf(__ad) || std::isinf(__bc))) {
-      if (std::isnan(__a))
-        __a = std::copysign(0, __a);
-      if (std::isnan(__b))
-        __b = std::copysign(0, __b);
-      if (std::isnan(__c))
-        __c = std::copysign(0, __c);
-      if (std::isnan(__d))
-        __d = std::copysign(0, __d);
+    if (IS_INF(__a) || IS_INF(__b)) {
+      __a = COPYSIGN(IS_INF(__a) ? 1 : 0, __a);
+      __b = COPYSIGN(IS_INF(__b) ? 1 : 0, __b);
+      if (IS_NAN(__c))
+        __c = COPYSIGN(0, __c);
+      if (IS_NAN(__d))
+        __d = COPYSIGN(0, __d);
+      __recalc = 1;
+    }
+    if (IS_INF(__c) || IS_INF(__d)) {
+      __c = COPYSIGN(IS_INF(__c) ? 1 : 0, __c);
+      __d = COPYSIGN(IS_INF(__d) ? 1 : 0, __d);
+      if (IS_NAN(__a))
+        __a = COPYSIGN(0, __a);
+      if (IS_NAN(__b))
+        __b = COPYSIGN(0, __b);
+      __recalc = 1;
+    }
+    if (!__recalc && (IS_INF(__ac) || IS_INF(__bd) ||
+                      IS_INF(__ad) || IS_INF(__bc))) {
+      if (IS_NAN(__a))
+        __a = COPYSIGN(0, __a);
+      if (IS_NAN(__b))
+        __b = COPYSIGN(0, __b);
+      if (IS_NAN(__c))
+        __c = COPYSIGN(0, __c);
+      if (IS_NAN(__d))
+        __d = COPYSIGN(0, __d);
       __recalc = 1;
     }
     if (__recalc) {
       // Can't use std::numeric_limits<double>::infinity() -- that doesn't have
       // a device overload (and isn't constexpr before C++11, naturally).
-      __real__(z) = __builtin_huge_valf() * (__a * __c - __b * __d);
-      __imag__(z) = __builtin_huge_valf() * (__a * __d + __b * __c);
+      __real__(z) = __builtin_huge_val() * (__a * __c - __b * __d);
+      __imag__(z) = __builtin_huge_val() * (__a * __d + __b * __c);
     }
   }
   return z;
 }
 
-extern "C" inline __device__ float _Complex __mulsc3(float __a, float __b,
-                                                     float __c, float __d) {
+__DEVICE__ float _Complex __mulsc3(float __a, float __b, float __c, float __d) {
   float __ac = __a * __c;
   float __bd = __b * __d;
   float __ad = __a * __d;
   float __bc = __b * __c;
   float _Complex z;
   __real__(z) = __ac - __bd;
   __imag__(z) = __ad + __bc;
-  if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
+  if (IS_NANF(__real__(z)) && IS_NANF(__imag__(z))) {
     int __recalc = 0;
-    if (std::isinf(__a) || std::isinf(__b)) {
-      __a = std::copysign(std::isinf(__a) ? 1 : 0, __a);
-      __b = std::copysign(std::isinf(__b) ? 1 : 0, __b);
-      if (std::isnan(__c))
-        __c = std::copysign(0, __c);
-      if (std::isnan(__d))
-        __d = std::copysign(0, __d);
-      __recalc = 1;
-    }
-    if (std::isinf(__c) || std::isinf(__d)) {
-      __c = std::copysign(std::isinf(__c) ? 1 : 0, __c);
-      __d = std::copysign(std::isinf(__d) ? 1 : 0, __d);
-      if (std::isnan(__a))
-        __a = std::copysign(0, __a);
-      if (std::isnan(__b))
-        __b = std::copysign(0, __b);
-      __recalc = 1;
-    }
-    if (!__recalc && (std::isinf(__ac) || std::isinf(__bd) ||
-                      std::isinf(__ad) || std::isinf(__bc))) {
-      if (std::isnan(__a))
-        __a = std::copysign(0, __a);
-      if (std::isnan(__b))
-        __b = std::copysign(0, __b);
-      if (std::isnan(__c))
-        __c = std::copysign(0, __c);
-      if (std::isnan(__d))
-        __d = std::copysign(0, __d);
+    if (IS_INFF(__a) || IS_INFF(__b)) {
+      __a = COPYSIGNF(IS_INFF(__a) ? 1 : 0, __a);
+      __b = COPYSIGNF(IS_INFF(__b) ? 1 : 0, __b);
+      if (IS_NANF(__c))
+        __c = COPYSIGNF(0, __c);
+      if (IS_NANF(__d))
+        __d = COPYSIGNF(0, __d);
+      __recalc = 1;
+    }
+    if (IS_INFF(__c) || IS_INFF(__d)) {
+      __c = COPYSIGNF(IS_INFF(__c) ? 1 : 0, __c);
+      __d = COPYSIGNF(IS_INFF(__d) ? 1 : 0, __d);
+      if (IS_NANF(__a))
+        __a = COPYSIGNF(0, __a);
+      if (IS_NANF(__b))
+        __b = COPYSIGNF(0, __b);
+      __recalc = 1;
+    }
+    if (!__recalc && (IS_INFF(__ac) || IS_INFF(__bd) ||
+                      IS_INFF(__ad) || IS_INFF(__bc))) {
+      if (IS_NANF(__a))
+        __a = COPYSIGNF(0, __a);
+      if (IS_NANF(__b))
+        __b = COPYSIGNF(0, __b);
+      if (IS_NANF(__c))
+        __c = COPYSIGNF(0, __c);
+      if (IS_NANF(__d))
+        __d = COPYSIGNF(0, __d);
       __recalc = 1;
     }
     if (__recalc) {
       __real__(z) = __builtin_huge_valf() * (__a * __c - __b * __d);
       __imag__(z) = __builtin_huge_valf() * (__a * __d + __b * __c);
     }
   }
   return z;
 }
 
-extern "C" inline __device__ double _Complex __divdc3(double __a, double __b,
-                                                      double __c, double __d) {
+__DEVICE__ double _Complex __divdc3(double __a, double __b, double __c,
+                                    double __d) {
   int __ilogbw = 0;
-  // Can't use std::max, because that's defined in <algorithm>, and we don't
-  // want to pull that in for every compile.  The CUDA headers define
-  // ::max(float, float) and ::max(double, double), which is sufficient for us.
-  double __logbw = std::logb(max(std::abs(__c), std::abs(__d)));
-  if (std::isfinite(__logbw)) {
+  double __logbw = LOGB(MAX(ABS(__c), ABS(__d)));
+  if (IS_FINITE(__logbw)) {
     __ilogbw = (int)__logbw;
-    __c = std::scalbn(__c, -__ilogbw);
-    __d = std::scalbn(__d, -__ilogbw);
+    __c = SCALBN(__c, -__ilogbw);
+    __d = SCALBN(__d, -__ilogbw);
   }
   double __denom = __c * __c + __d * __d;
   double _Complex z;
-  __real__(z) = std::scalbn((__a * __c + __b * __d) / __denom, -__ilogbw);
-  __imag__(z) = std::scalbn((__b * __c - __a * __d) / __denom, -__ilogbw);
-  if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
-    if ((__denom == 0.0) && (!std::isnan(__a) || !std::isnan(__b))) {
-      __real__(z) = std::copysign(__builtin_huge_valf(), __c) * __a;
-      __imag__(z) = std::copysign(__builtin_huge_valf(), __c) * __b;
-    } else if ((std::isinf(__a) || std::isinf(__b)) && std::isfinite(__c) &&
-               std::isfinite(__d)) {
-      __a = std::copysign(std::isinf(__a) ? 1.0 : 0.0, __a);
-      __b = std::copysign(std::isinf(__b) ? 1.0 : 0.0, __b);
-      __real__(z) = __builtin_huge_valf() * (__a * __c + __b * __d);
-      __imag__(z) = __builtin_huge_valf() * (__b * __c - __a * __d);
-    } else if (std::isinf(__logbw) && __logbw > 0.0 && std::isfinite(__a) &&
-               std::isfinite(__b)) {
-      __c = std::copysign(std::isinf(__c) ? 1.0 : 0.0, __c);
-      __d = std::copysign(std::isinf(__d) ? 1.0 : 0.0, __d);
+  __real__(z) = SCALBN((__a * __c + __b * __d) / __denom, -__ilogbw);
+  __imag__(z) = SCALBN((__b * __c - __a * __d) / __denom, -__ilogbw);
+  if (IS_NAN(__real__(z)) && IS_NAN(__imag__(z))) {
+    if ((__denom == 0.0) && (!IS_NAN(__a) || !IS_NAN(__b))) {
+      __real__(z) = COPYSIGN(__builtin_huge_val(), __c) * __a;
+      __imag__(z) = COPYSIGN(__builtin_huge_val(), __c) * __b;
+    } else if ((IS_INF(__a) || IS_INF(__b)) && IS_FINITE(__c) &&
+               IS_FINITE(__d)) {
+      __a = COPYSIGN(IS_INF(__a) ? 1.0 : 0.0, __a);
+      __b = COPYSIGN(IS_INF(__b) ? 1.0 : 0.0, __b);
+      __real__(z) = __builtin_huge_val() * (__a * __c + __b * __d);
+      __imag__(z) = __builtin_huge_val() * (__b * __c - __a * __d);
+    } else if (IS_INF(__logbw) && __logbw > 0.0 && IS_FINITE(__a) &&
+               IS_FINITE(__b)) {
+      __c = COPYSIGN(IS_INF(__c) ? 1.0 : 0.0, __c);
+      __d = COPYSIGN(IS_INF(__d) ? 1.0 : 0.0, __d);
       __real__(z) = 0.0 * (__a * __c + __b * __d);
       __imag__(z) = 0.0 * (__b * __c - __a * __d);
     }
   }
   return z;
 }
 
-extern "C" inline __device__ float _Complex __divsc3(float __a, float __b,
-                                                     float __c, float __d) {
+__DEVICE__ float _Complex __divsc3(float __a, float __b, float __c, float __d) {
   int __ilogbw = 0;
-  float __logbw = std::logb(max(std::abs(__c), std::abs(__d)));
-  if (std::isfinite(__logbw)) {
+  float __logbw = LOGBF(MAXF(ABSF(__c), ABSF(__d)));
+  if (IS_FINITEF(__logbw)) {
     __ilogbw = (int)__logbw;
-    __c = std::scalbn(__c, -__ilogbw);
-    __d = std::scalbn(__d, -__ilogbw);
+    __c = SCALBNF(__c, -__ilogbw);
+    __d = SCALBNF(__d, -__ilogbw);
   }
   float __denom = __c * __c + __d * __d;
   float _Complex z;
-  __real__(z) = std::scalbn((__a * __c + __b * __d) / __denom, -__ilogbw);
-  __imag__(z) = std::scalbn((__b * __c - __a * __d) / __denom, -__ilogbw);
-  if (std::isnan(__real__(z)) && std::isnan(__imag__(z))) {
-    if ((__denom == 0) && (!std::isnan(__a) || !std::isnan(__b))) {
-      __real__(z) = std::copysign(__builtin_huge_valf(), __c) * __a;
-      __imag__(z) = std::copysign(__builtin_huge_valf(), __c) * __b;
-    } else if ((std::isinf(__a) || std::isinf(__b)) && std::isfinite(__c) &&
-               std::isfinite(__d)) {
-      __a = std::copysign(std::isinf(__a) ? 1 : 0, __a);
-      __b = std::copysign(std::isinf(__b) ? 1 : 0, __b);
+  __real__(z) = SCALBNF((__a * __c + __b * __d) / __denom, -__ilogbw);
+  __imag__(z) = SCALBNF((__b * __c - __a * __d) / __denom, -__ilogbw);
+  if (IS_NANF(__real__(z)) && IS_NANF(__imag__(z))) {
+    if ((__denom == 0) && (!IS_NANF(__a) || !IS_NANF(__b))) {
+      __real__(z) = COPYSIGNF(__builtin_huge_valf(), __c) * __a;
+      __imag__(z) = COPYSIGNF(__builtin_huge_valf(), __c) * __b;
+    } else if ((IS_INFF(__a) || IS_INFF(__b)) && IS_FINITEF(__c) &&
+               IS_FINITEF(__d)) {
+      __a = COPYSIGNF(IS_INFF(__a) ? 1 : 0, __a);
+      __b = COPYSIGNF(IS_INFF(__b) ? 1 : 0, __b);
       __real__(z) = __builtin_huge_valf() * (__a * __c + __b * __d);
       __imag__(z) = __builtin_huge_valf() * (__b * __c - __a * __d);
-    } else if (std::isinf(__logbw) && __logbw > 0 && std::isfinite(__a) &&
-               std::isfinite(__b)) {
-      __c = std::copysign(std::isinf(__c) ? 1 : 0, __c);
-      __d = std::copysign(std::isinf(__d) ? 1 : 0, __d);
+    } else if (IS_INFF(__logbw) && __logbw > 0 && IS_FINITEF(__a) &&
+               IS_FINITEF(__b)) {
+      __c = COPYSIGNF(IS_INFF(__c) ? 1 : 0, __c);
+      __d = COPYSIGNF(IS_INFF(__d) ? 1 : 0, __d);
       __real__(z) = 0 * (__a * __c + __b * __d);
       __imag__(z) = 0 * (__b * __c - __a * __d);
     }
   }
   return z;
 }
 
+#ifdef _OPENMP
+#pragma omp end declare target
+#endif // _OPENMP
+
+#undef IS_NAN(x)
+#undef IS_NANF(x)
+#undef IS_INF(x)
+#undef IS_INFF(x)
+#undef IS_FINITE(x)
+#undef IS_FINITEF(x)
+#undef COPYSIGN(x, y)
+#undef COPYSIGNF(x, y)
+#undef LOGB(x)
+#undef LOGBF(x)
+#undef ABS(x)
+#undef ABSF(x)
+#undef SCALBN(x, y)
+#undef SCALBNF(x, y)
+
 #endif // __CLANG_CUDA_COMPLEX_BUILTINS

clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

 #if defined(__cplusplus)
   #include <__clang_cuda_math_forward_declares.h>
 #endif
 
 /// Include declarations for libdevice functions.
 #include <__clang_cuda_libdevice_declares.h>
 /// Provide definitions for these functions.
 #include <__clang_cuda_device_functions.h>
+/// Provide definitions for complex math functions.
+#include <__clang_cuda_complex_builtins.h>
 
 #undef __CUDA__
 
 #endif
 #endif

clang/test/Headers/nvptx_device_cmath_functions.c

 // Test calling of device math functions.
 ///==========================================================================///
 
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // 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 -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 -o - | FileCheck -check-prefix CHECK-YES %s
 
 #include <cmath>
 
+// CHECK-YES-DAG: define {{.*}}hidden { double, double } @__muldc3(
+// CHECK-YES-DAG: define {{.*}}hidden { float, float } @__mulsc3(
+// CHECK-YES-DAG: define {{.*}}hidden { double, double } @__divdc3(
+// CHECK-YES-DAG: define {{.*}}hidden { float, float } @__divsc3(
+
 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);
   }
 }
+
+// CHECK-YES-NOT: @{{__mulsc3|__divsc3}}(
+void test_cmplx(float _Complex a) {
+#pragma omp target
+  {
+    (void)(a * (a / a));
+  }
+}

clang/test/Headers/nvptx_device_cmath_functions.cpp

#pragma omp target
     // 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);
   }
 }
+
+// CHECK-YES-DAG: call { float, float } @__mulsc3(
+// CHECK-YES-DAG: call { float, float } @__divsc3(
+void test_cmplx(float _Complex a) {
+#pragma omp target
+  {
+    (void)(a * (a / a));
+  }
+}
+// CHECK-YES-DAG: define {{.*}}hidden { float, float } @__mulsc3(
+// CHECK-YES-DAG: define {{.*}}hidden { float, float } @__divsc3(
+

clang/test/Headers/nvptx_device_cmath_functions_cxx17.cpp

#pragma omp target
     // 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);
   }
 }
+
+// CHECK-YES-DAG: call { double, double } @__muldc3(
+// CHECK-YES-DAG: call { double, double } @__divdc3(
+void test_cmplx(double _Complex a) {
+#pragma omp target
+  {
+    (void)(a * (a / a));
+  }
+}
+// CHECK-YES-DAG: define {{.*}}hidden { double, double } @__muldc3(
+// CHECK-YES-DAG: define {{.*}}hidden { double, double } @__divdc3(
+