1 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
2 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -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 %s
3 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
4 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
5 // RUN: %clang_cc1 -x c++ -O1 -disable-llvm-optzns -verify -fopenmp -internal-isystem %S/../Headers/Inputs/include -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -fexceptions -fcxx-exceptions -aux-triple powerpc64le-unknown-unknown -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
6 // expected-no-diagnostics
7 #ifndef HEADER
8 #define HEADER
9
10 #include <cmath>
11
12 // TODO: Think about teaching the OMPIRBuilder about default attributes as well so the __kmpc* declarations are annotated.
13
14 // CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}() [[attr0:#[0-9]]]
15 // CHECK: call i32 @__kmpc_target_init(
16 // CHECK: declare float @_Z3sinf(float) [[attr1:#[0-9]*]]
17 // CHECK: declare void @__kmpc_target_deinit(
18 // CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}() [[attr0]]
19 // CHECK: %call = call double @_Z3sind(double 0.000000e+00) [[attr2:#[0-9]]]
20 // CHECK: declare double @_Z3sind(double) [[attr1]]
21
22 // CHECK: attributes [[attr0]]
23 // CHECK-NOT: "llvm.assume"
24 // CHECK: attributes [[attr1]]
25 // CHECK-SAME: "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations"
26 // CHECK: attributes [[attr2]]
27 // CHECK-SAME: "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations"
28
29
30 template <typename T>
foo()31 void foo() {
32 cos(T(0));
33 }
34
35 template <typename T>
complex_reduction()36 void complex_reduction() {
37 foo<T>();
38 #pragma omp target
39 sin(T(0));
40 }
41
42 #pragma omp assumes ext_check_that_this_is_attached_to_included_functions_and_template_instantiations
43
test()44 void test() {
45 complex_reduction<float>();
46 complex_reduction<double>();
47 }
48 #endif
49