pdhaliwal created this revision.
pdhaliwal added reviewers: ye-luo, jdoerfert, ronlieb, JonChesterfield.
Herald added subscribers: guansong, t-tye, tpr, dstuttard, yaxunl, jvesely,
kzhuravl.
pdhaliwal requested review of this revision.
Herald added subscribers: cfe-commits, sstefan1, wdng.
Herald added a project: clang.
Following nvptx approach, this patch uses complex function
definitions from complex_cmath.h. With this patch, ovo passes
23/34 complex mathematical test cases.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D109344
Files:
clang/lib/Headers/openmp_wrappers/complex
clang/test/Headers/amdgcn-openmp-device-math-complex.cpp
Index: clang/test/Headers/amdgcn-openmp-device-math-complex.cpp
===================================================================
--- /dev/null
+++ clang/test/Headers/amdgcn-openmp-device-math-complex.cpp
@@ -0,0 +1,85 @@
+// RUN: %clang_cc1 -verify -internal-isystem %S/Inputs/include -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -x c++ -triple amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -aux-triple x86_64-unknown-unknown -o - | FileCheck %s
+// expected-no-diagnostics
+
+#include <cmath>
+#include <complex>
+
+// CHECK: define weak {{.*}} @__muldc3
+// CHECK-DAG: call i32 @__ocml_isnan_f64(
+// CHECK-DAG: call i32 @__ocml_isinf_f64(
+
+// 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(
+
+// CHECK: define weak {{.*}} @__divdc3
+// CHECK-DAG: call i32 @__ocml_isnan_f64(
+// CHECK-DAG: call i32 @__ocml_isinf_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_fabs_f64(
+// CHECK-DAG: call double @__ocml_logb_f64(
+
+// CHECK: define weak {{.*}} @__divsc3
+// CHECK-DAG: call i32 @__ocml_isnan_f32(
+// CHECK-DAG: call i32 @__ocml_isinf_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_fabs_f32(
+// CHECK-DAG: call float @__ocml_logb_f32(
+
+// We actually check that there are no declarations of non-OpenMP functions.
+// That is, as long as we don't call an unkown function with a name that
+// doesn't start with '__' we are good :)
+
+// CHECK-NOT: declare.*@[^_]
+
+void test_scmplx(std::complex<float> a) {
+#pragma omp target
+ {
+ (void)(a * (a / a));
+ }
+}
+
+void test_dcmplx(std::complex<double> a) {
+#pragma omp target
+ {
+ (void)(a * (a / a));
+ }
+}
+
+template <typename T>
+std::complex<T> test_template_math_calls(std::complex<T> a) {
+ decltype(a) r = a;
+#pragma omp target
+ {
+ r = std::sin(r);
+ r = std::cos(r);
+ r = std::exp(r);
+ r = std::atan(r);
+ r = std::acos(r);
+ }
+ return r;
+}
+
+std::complex<float> test_scall(std::complex<float> a) {
+ decltype(a) r;
+#pragma omp target
+ {
+ r = std::sin(a);
+ }
+ return test_template_math_calls(r);
+}
+
+std::complex<double> test_dcall(std::complex<double> a) {
+ decltype(a) r;
+#pragma omp target
+ {
+ r = std::exp(a);
+ }
+ return test_template_math_calls(r);
+}
Index: clang/lib/Headers/openmp_wrappers/complex
===================================================================
--- clang/lib/Headers/openmp_wrappers/complex
+++ clang/lib/Headers/openmp_wrappers/complex
@@ -44,9 +44,17 @@
// handle.
#ifndef _LIBCPP_STD_VER
+#ifdef __NVPTX__
#pragma omp begin declare variant match( \
device = {arch(nvptx, nvptx64)}, \
implementation = {extension(match_any, allow_templates)})
+#endif
+
+#ifdef __AMDGCN__
+#pragma omp begin declare variant match( \
+ device = {arch(amdgcn)}, \
+ implementation = {extension(match_any, allow_templates)})
+#endif
#include <complex_cmath.h>
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
