Author: Pushpinder Singh Date: 2021-09-09T16:20:23-07:00 New Revision: 718280c8a2d1beec1ec4c130f96648ab3e41f38c
URL: https://github.com/llvm/llvm-project/commit/718280c8a2d1beec1ec4c130f96648ab3e41f38c DIFF: https://github.com/llvm/llvm-project/commit/718280c8a2d1beec1ec4c130f96648ab3e41f38c.diff LOG: [AMDGPU][OpenMP] Use complex definitions from complex_cmath.h Following nvptx approach, this patch uses complex function definitions from complex_cmath.h. With this patch, ovo passes 23/34 complex mathematical test cases. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D109344 (cherry picked from commit 12dcbf913c49db839b3669db0dcacd5de25facde) Added: clang/test/Headers/amdgcn-openmp-device-math-complex.cpp Modified: clang/lib/Headers/openmp_wrappers/complex Removed: ################################################################################ diff --git a/clang/lib/Headers/openmp_wrappers/complex b/clang/lib/Headers/openmp_wrappers/complex index dfd6193c97cbd..eb1ead207d582 100644 --- a/clang/lib/Headers/openmp_wrappers/complex +++ b/clang/lib/Headers/openmp_wrappers/complex @@ -36,7 +36,7 @@ #ifndef _LIBCPP_STD_VER #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, \ + device = {arch(amdgcn, nvptx, nvptx64)}, \ implementation = {extension(match_any, allow_templates)}) #include <complex_cmath.h> diff --git a/clang/test/Headers/amdgcn-openmp-device-math-complex.cpp b/clang/test/Headers/amdgcn-openmp-device-math-complex.cpp new file mode 100644 index 0000000000000..d1a2cf31fabae --- /dev/null +++ b/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); +} _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
