[OpenMP][AMDGCN] Enable complex functions

[pdhaliwal]

Summary

This patch enables basic complex functionality using the ocml builtins.

[JonChesterfield]

Macros fine by me, haven't found time to look at what is going on with variant

clang/lib/Headers/openmp_wrappers/complex

20

I don't think this should be necessary - the declare variant stuff works for nvptx afaik

[jdoerfert]

LG

clang/lib/Headers/openmp_wrappers/complex

20

You need it because the file has include guards.

clang/lib/Headers/__clang_cuda_complex_builtins.h

 #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 and OpenMP target offloading [in C and C++ mode].)
 
 #pragma push_macro("__DEVICE__")
-#ifdef __OPENMP_NVPTX__
+#if defined(__OPENMP_NVPTX__) || defined(__OPENMP_AMDGCN__)
 #pragma omp declare target
 #define __DEVICE__ __attribute__((noinline, nothrow, cold, weak))
 #else
 #define __DEVICE__ __device__ inline
 #endif
 
 // To make the algorithms available for C and C++ in CUDA and OpenMP we select
 // different but equivalent function versions. TODO: For OpenMP we currently
 // select the native builtins as the overload support for templates is lacking.
-#if !defined(__OPENMP_NVPTX__)
+#if !defined(__OPENMP_NVPTX__) && !defined(__OPENMP_AMDGCN__)
 #define _ISNANd std::isnan
 #define _ISNANf std::isnan
 #define _ISINFd std::isinf
 #define _ISINFf std::isinf
 #define _ISFINITEd std::isfinite
 #define _ISFINITEf std::isfinite
 #define _COPYSIGNd std::copysign
 #define _COPYSIGNf std::copysign

 #define _fmaxf __nv_fmaxf
 #endif
 #endif
 
 #if defined(__cplusplus)
 extern "C" {
 #endif
 
 __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 (_ISNANd(__real__(z)) && _ISNANd(__imag__(z))) {
     int __recalc = 0;
     if (_ISINFd(__a) || _ISINFd(__b)) {
       __a = _COPYSIGNd(_ISINFd(__a) ? 1 : 0, __a);
       __b = _COPYSIGNd(_ISINFd(__b) ? 1 : 0, __b);
       if (_ISNANd(__c))
         __c = _COPYSIGNd(0, __c);
       if (_ISNANd(__d))
         __d = _COPYSIGNd(0, __d);
       __recalc = 1;
     }
     if (_ISINFd(__c) || _ISINFd(__d)) {
       __c = _COPYSIGNd(_ISINFd(__c) ? 1 : 0, __c);
       __d = _COPYSIGNd(_ISINFd(__d) ? 1 : 0, __d);
       if (_ISNANd(__a))
         __a = _COPYSIGNd(0, __a);
       if (_ISNANd(__b))
         __b = _COPYSIGNd(0, __b);
       __recalc = 1;
     }
     if (!__recalc &&
         (_ISINFd(__ac) || _ISINFd(__bd) || _ISINFd(__ad) || _ISINFd(__bc))) {

 #undef _SCALBNf
 #undef _ABSd
 #undef _ABSf
 #undef _LOGBd
 #undef _LOGBf
 #undef _fmaxd
 #undef _fmaxf
 
-#ifdef __OPENMP_NVPTX__
+#if defined(__OPENMP_NVPTX__) || defined(__OPENMP_AMDGCN__)
 #pragma omp end declare target
 #endif
 
 #pragma pop_macro("__DEVICE__")
 
 #endif // __CLANG_CUDA_COMPLEX_BUILTINS

clang/lib/Headers/openmp_wrappers/complex

 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 // We require std::math functions in the complex builtins below.
 #include <cmath>
 
+#ifdef __NVPTX__
 #define __OPENMP_NVPTX__
 #include <__clang_cuda_complex_builtins.h>
 #undef __OPENMP_NVPTX__
+#endif // __NVPTX__
+
+#ifdef __AMDGCN__
+#define __OPENMP_AMDGCN__
+#include <__clang_cuda_complex_builtins.h>
+#undef __OPENMP_AMDGCN__
+#endif // __AMDGCN__
+
 #endif
 
 // Grab the host header too.
 #include_next <complex>
 
 // If we are compiling against libc++, the macro _LIBCPP_STD_VER should be set
 // after including <cmath> above. Since the complex header we use is a
 // simplified version of the libc++, we don't need it in this case. If we
 // compile against libstdc++, or any other standard library, we will overload
 // the (hopefully template) functions in the <complex> header with the ones we
 // got from libc++ which decomposes math functions, like `std::sin`, into
 // arithmetic and calls to non-complex functions, all of which we can then
 // handle.
 #ifndef _LIBCPP_STD_VER
 
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)},                                           \
     implementation = {extension(match_any, allow_templates)})
 
 #include <complex_cmath.h>
 
 #pragma omp end declare variant
 
-#endif
+#endif // _LIBCPP_STD_VER

clang/lib/Headers/openmp_wrappers/complex.h

 /*===-- complex --- OpenMP complex wrapper for target regions --------- c++ -===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_OPENMP_COMPLEX_H__
 #define __CLANG_OPENMP_COMPLEX_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 // We require math functions in the complex builtins below.
 #include <math.h>
 
+#ifdef __NVPTX__
 #define __OPENMP_NVPTX__
 #include <__clang_cuda_complex_builtins.h>
 #undef __OPENMP_NVPTX__
 #endif
 
+#ifdef __AMDGCN__
+#define __OPENMP_AMDGCN__
+#include <__clang_cuda_complex_builtins.h>
+#undef __OPENMP_AMDGCN__
+#endif
+
+#endif
+
 // Grab the host header too.
 #include_next <complex.h>

clang/test/Headers/amdgcn-openmp-device-math-complex.c

+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK
+
+#include <complex.h>
+
+void test_complex_f64(double _Complex a) {
+// CHECK-LABEL: define {{.*}}test_complex_f64
+#pragma omp target
+  {
+    // CHECK: call { double, double } @__divdc3
+    // CHECK: call { double, double } @__muldc3
+    (void)(a * (a / a));
+  }
+}
+
+// CHECK: define weak {{.*}} @__divdc3
+// CHECK-DAG: call double @__ocml_fabs_f64(
+// CHECK-DAG: call i32 @__ocml_isnan_f64(
+// CHECK-DAG: call i32 @__ocml_isfinite_f64(
+// CHECK-DAG: call double @__ocml_copysign_f64(
+// CHECK-DAG: call double @__ocml_scalbn_f64(
+// CHECK-DAG: call double @__ocml_logb_f64(
+
+// CHECK: define weak {{.*}} @__muldc3
+// CHECK-DAG: call i32 @__ocml_isnan_f64(
+// CHECK-DAG: call i32 @__ocml_isinf_f64(
+// CHECK-DAG: call double @__ocml_copysign_f64(
+
+void test_complex_f32(float _Complex a) {
+// CHECK-LABEL: define {{.*}}test_complex_f32
+#pragma omp target
+  {
+    // CHECK: call [2 x i32] @__divsc3
+    // CHECK: call [2 x i32] @__mulsc3
+    (void)(a * (a / a));
+  }
+}
+
+// CHECK: define weak {{.*}} @__divsc3
+// CHECK-DAG: call float @__ocml_fabs_f32(
+// CHECK-DAG: call i32 @__ocml_isnan_f32(
+// CHECK-DAG: call i32 @__ocml_isfinite_f32(
+// CHECK-DAG: call float @__ocml_copysign_f32(
+// CHECK-DAG: call float @__ocml_scalbn_f32(
+// CHECK-DAG: call float @__ocml_logb_f32(
+
+// CHECK: define weak {{.*}} @__mulsc3
+// CHECK-DAG: call i32 @__ocml_isnan_f32(
+// CHECK-DAG: call i32 @__ocml_isinf_f32(
+// CHECK-DAG: call float @__ocml_copysign_f32(