1 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix HOST --check-prefix CHECK
2 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
3 // RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix DEVICE --check-prefix CHECK
4 // RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
5 // RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix DEVICE --check-prefix CHECK
6
7 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o - | FileCheck %s --check-prefix SIMD-ONLY
8 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
9 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY
10 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
11 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix SIMD-ONLY
12
13 #ifndef HEADER
14 #define HEADER
15
16 // SIMD-ONLY-NOT: {{__kmpc|__tgt}}
17
18 // DEVICE-DAG: [[C_ADDR:.+]] = internal global i32 0,
19 // DEVICE-DAG: [[CD_ADDR:@.+]] ={{ protected | }}global %struct.S zeroinitializer,
20 // HOST-DAG: @[[C_ADDR:.+]] = internal global i32 0,
21 // HOST-DAG: @[[CD_ADDR:.+]] ={{( protected | dso_local)?}} global %struct.S zeroinitializer,
22
23 // DEVICE-DAG: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @[[CTOR:.+]], ptr null }]
24 // DEVICE-DAG: @llvm.global_dtors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 65535, ptr @[[DTOR:.+]], ptr null }]
25
26 #pragma omp declare target
foo()27 int foo() { return 0; }
28 #pragma omp end declare target
bar()29 int bar() { return 0; }
30 #pragma omp declare target (bar)
baz()31 int baz() { return 0; }
32
33 #pragma omp declare target
doo()34 int doo() { return 0; }
35 #pragma omp end declare target
car()36 int car() { return 0; }
37 #pragma omp declare target (bar)
caz()38 int caz() { return 0; }
39
40 // DEVICE-DAG: define hidden noundef i32 [[FOO:@.*foo.*]]()
41 // DEVICE-DAG: define hidden noundef i32 [[BAR:@.*bar.*]]()
42 // DEVICE-DAG: define hidden noundef i32 [[BAZ:@.*baz.*]]()
43 // DEVICE-DAG: define hidden noundef i32 [[DOO:@.*doo.*]]()
44 // DEVICE-DAG: define hidden noundef i32 [[CAR:@.*car.*]]()
45 // DEVICE-DAG: define hidden noundef i32 [[CAZ:@.*caz.*]]()
46
47 static int c = foo() + bar() + baz();
48 #pragma omp declare target (c)
49
50 struct S {
51 int a;
52 S() = default;
SS53 S(int a) : a(a) {}
~SS54 ~S() { a = 0; }
55 };
56
57 #pragma omp declare target
58 S cd = doo() + car() + caz() + baz();
59 #pragma omp end declare target
60
maini1()61 int maini1() {
62 int a;
63 #pragma omp target map(tofrom : a)
64 {
65 a = c;
66 }
67 return 0;
68 }
69
70 // DEVICE-DAG: define weak{{.*}} void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](ptr {{[^,]*}}, ptr noundef nonnull align {{[0-9]+}} dereferenceable{{[^,]*}}
71 // DEVICE-DAG: [[C:%.+]] = load i32, ptr [[C_ADDR]],
72 // DEVICE-DAG: store i32 [[C]], ptr %
73
74 // HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-11]](ptr noundef nonnull align {{[0-9]+}} dereferenceable{{.*}})
75 // HOST: [[C:%.*]] = load i32, ptr @[[C_ADDR]],
76 // HOST: store i32 [[C]], ptr %
77
78 // HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}}
79 // HOST-DAG: !{i32 1, !"[[C_ADDR]]", i32 0, i32 {{[0-9]+}}}
80
81 #endif // HEADER
82
83