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 internal void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}_worker() [[attr0:#[0-9]*]]
15 // CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIfEvv_{{.*}}() [[attr0]]
16 // CHECK: %call = call float @_Z3sinf(float 0.000000e+00) [[attr5:#[0-9]*]]
17 // CHECK-DAG: declare i32 @llvm.nvvm.read.ptx.sreg.warpsize() [[attr1:#[0-9]*]]
18 // CHECK-DAG: declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() [[attr1]]
19 // CHECK-DAG: declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() [[attr1]]
20 // CHECK: declare void @__kmpc_kernel_init(i32, i16)
21 // CHECK-NOT: #
22 // CHECK: declare void @__kmpc_data_sharing_init_stack()
23 // CHECK-NOT: #
24 // CHECK: declare float @_Z3sinf(float) [[attr2:#[0-9]*]]
25 // CHECK: declare void @__kmpc_kernel_deinit(i16)
26 // CHECK-NOT: #
27 // CHECK: declare void @__kmpc_barrier_simple_spmd(%struct.ident_t*, i32) [[attr3:#[0-9]*]]
28 // CHECK: declare i1 @__kmpc_kernel_parallel(i8**)
29 // CHECK-NOT: #
30 // CHECK: declare i32 @__kmpc_global_thread_num(%struct.ident_t*) [[attr4:#[0-9]*]]
31 // CHECK: declare void @__kmpc_kernel_end_parallel()
32 // CHECK-NOT: #
33 // CHECK: define internal void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}_worker() [[attr0]]
34 // CHECK: define weak void @__omp_offloading_{{.*}}__Z17complex_reductionIdEvv_{{.*}}() [[attr0]]
35 // CHECK: %call = call double @_Z3sind(double 0.000000e+00) [[attr5]]
36 // CHECK: declare double @_Z3sind(double) [[attr2]]
37
38 // CHECK: attributes [[attr0]]
39 // CHECK-NOT: "llvm.assume"
40 // CHECK: attributes [[attr1]]
41 // CHECK-NOT: "llvm.assume"
42 // CHECK: attributes [[attr2]]
43 // CHECK-SAME: "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations"
44 // CHECK: attributes [[attr3]]
45 // CHECK-NOT: "llvm.assume"
46 // CHECK: attributes [[attr4]]
47 // CHECK-NOT: "llvm.assume"
48 // CHECK: attributes [[attr5]]
49 // CHECK-SAME: "llvm.assume"="check_that_this_is_attached_to_included_functions_and_template_instantiations"
50
51
52 template <typename T>
foo()53 void foo() {
54 cos(T(0));
55 }
56
57 template <typename T>
complex_reduction()58 void complex_reduction() {
59 foo<T>();
60 #pragma omp target
61 sin(T(0));
62 }
63
64 #pragma omp assumes ext_check_that_this_is_attached_to_included_functions_and_template_instantiations
65
test()66 void test() {
67 complex_reduction<float>();
68 complex_reduction<double>();
69 }
70 #endif
71