From 1c40858fc28adb9b87e3f8f8bd5b4319defa9624 Mon Sep 17 00:00:00 2001 From: Gheorghe-Teodor Bercea Date: Mon, 6 May 2019 18:19:15 +0000 Subject: [OpenMP][Clang] Support for target math functions Summary: In this patch we propose a temporary solution to resolving math functions for the NVPTX toolchain, temporary until OpenMP variant is supported by Clang. We intercept the inclusion of math.h and cmath headers and if we are in the OpenMP-NVPTX case, we re-use CUDA's math function resolution mechanism. Authors: @gtbercea @jdoerfert Reviewers: hfinkel, caomhin, ABataev, tra Reviewed By: hfinkel, ABataev, tra Subscribers: mgorny, guansong, cfe-commits, jdoerfert Tags: #clang Differential Revision: https://reviews.llvm.org/D61399 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@360063 91177308-0d34-0410-b5e6-96231b3b80d8 --- test/Headers/nvptx_device_math_functions.c | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) create mode 100644 test/Headers/nvptx_device_math_functions.c (limited to 'test/Headers/nvptx_device_math_functions.c') diff --git a/test/Headers/nvptx_device_math_functions.c b/test/Headers/nvptx_device_math_functions.c new file mode 100644 index 0000000000..733ad52bd1 --- /dev/null +++ b/test/Headers/nvptx_device_math_functions.c @@ -0,0 +1,21 @@ +// Test calling of device math functions. +///==========================================================================/// + +// REQUIRES: nvptx-registered-target + +// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -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 math.h -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 + +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); + } +} -- cgit v1.2.3