64 lines · cpp
1// REQUIRES: nvptx-registered-target2// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc3// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s4// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -DCMATH -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s5 6#ifdef CMATH7#include <cmath>8#else9#include <math.h>10#endif11 12// 4 calls to sincos(f), all translated to __nv_sincos calls:13 14// CHECK-NOT: _Z.sincos15// CHECK: call void @__nv_sincos(double16// CHECK-NOT: _Z.sincos17// CHECK: call void @__nv_sincosf(float18// CHECK-NOT: _Z.sincos19// CHECK: call void @__nv_sincos(double20// CHECK-NOT: _Z.sincos21// CHECK: call void @__nv_sincosf(float22// CHECK-NOT: _Z.sincos23 24// single precision wrapper25inline void sincos(float x, float* __restrict__ sin, float* __restrict__ cos)26{27 sincosf(x, sin, cos);28}29 30template<typename T>31void test_sincos(T x)32{33 T res_sin, res_cos;34 35 #pragma omp target map(from: res_sin, res_cos)36 {37 sincos(x, &res_sin, &res_cos);38 }39 40}41 42int main(int argc, char **argv)43{44 45#if !defined(C_ONLY)46 test_sincos<double>(0.0);47 test_sincos<float>(0.0);48#endif49 50 #pragma omp target51 {52 double s, c;53 sincos(0, &s, &c);54 }55 56 #pragma omp target57 {58 float s, c;59 sincosf(0.f, &s, &c);60 }61 62 return 0;63}64