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-device -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 -emit-llvm %s -fopenmp-is-device -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-device -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-device -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 -emit-llvm %s -fopenmp-is-device -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-device -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:@.+]] ={{ hidden | }}global %struct.S zeroinitializer,
20 // HOST-DAG: @[[C_ADDR:.+]] = internal global i32 0,
21 // HOST-DAG: @[[CD_ADDR:.+]] ={{( hidden | dso_local)?}} global %struct.S zeroinitializer,
22
23 #pragma omp declare target
foo()24 int foo() { return 0; }
25 #pragma omp end declare target
bar()26 int bar() { return 0; }
27 #pragma omp declare target (bar)
baz()28 int baz() { return 0; }
29
30 #pragma omp declare target
doo()31 int doo() { return 0; }
32 #pragma omp end declare target
car()33 int car() { return 0; }
34 #pragma omp declare target (bar)
caz()35 int caz() { return 0; }
36
37 // DEVICE-DAG: define{{ hidden | }}i32 [[FOO:@.*foo.*]]()
38 // DEVICE-DAG: define{{ hidden | }}i32 [[BAR:@.*bar.*]]()
39 // DEVICE-DAG: define{{ hidden | }}i32 [[BAZ:@.*baz.*]]()
40 // DEVICE-DAG: define{{ hidden | }}i32 [[DOO:@.*doo.*]]()
41 // DEVICE-DAG: define{{ hidden | }}i32 [[CAR:@.*car.*]]()
42 // DEVICE-DAG: define{{ hidden | }}i32 [[CAZ:@.*caz.*]]()
43
44 static int c = foo() + bar() + baz();
45 #pragma omp declare target (c)
46 // HOST-DAG: @[[C_CTOR:__omp_offloading__.+_c_l44_ctor]] = private constant i8 0
47 // DEVICE-DAG: define internal void [[C_CTOR:@__omp_offloading__.+_c_l44_ctor]]()
48 // DEVICE-DAG: call i32 [[FOO]]()
49 // DEVICE-DAG: call i32 [[BAR]]()
50 // DEVICE-DAG: call i32 [[BAZ]]()
51 // DEVICE-DAG: ret void
52
53 struct S {
54 int a;
55 S() = default;
SS56 S(int a) : a(a) {}
~SS57 ~S() { a = 0; }
58 };
59
60 #pragma omp declare target
61 S cd = doo() + car() + caz() + baz();
62 #pragma omp end declare target
63 // HOST-DAG: @[[CD_CTOR:__omp_offloading__.+_cd_l61_ctor]] = private constant i8 0
64 // DEVICE-DAG: define internal void [[CD_CTOR:@__omp_offloading__.+_cd_l61_ctor]]()
65 // DEVICE-DAG: call i32 [[DOO]]()
66 // DEVICE-DAG: call i32 [[CAR]]()
67 // DEVICE-DAG: call i32 [[CAZ]]()
68 // DEVICE-DAG: ret void
69
70 // HOST-DAG: @[[CD_DTOR:__omp_offloading__.+_cd_l61_dtor]] = private constant i8 0
71 // DEVICE-DAG: define internal void [[CD_DTOR:@__omp_offloading__.+_cd_l61_dtor]]()
72 // DEVICE-DAG: call void
73 // DEVICE-DAG: ret void
74
75 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_ADDR]]\00"
76 // HOST-DAG: @.omp_offloading.entry.[[C_ADDR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (i32* @[[C_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section "omp_offloading_entries", align 1
77 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_ADDR]]\00"
78 // HOST-DAG: @.omp_offloading.entry.[[CD_ADDR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* bitcast (%struct.S* @[[CD_ADDR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 4, i32 0, i32 0 }, section "omp_offloading_entries", align 1
79 // HOST-DAG: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_CTOR]]\00"
80 // HOST-DAG: @.omp_offloading.entry.[[C_CTOR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* @[[C_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section "omp_offloading_entries", align 1
81 // HOST-DAG: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_CTOR]]\00"
82 // HOST-DAG: @.omp_offloading.entry.[[CD_CTOR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* @[[CD_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section "omp_offloading_entries", align 1
83 // HOST-DAG: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_DTOR]]\00"
84 // HOST-DAG: @.omp_offloading.entry.[[CD_DTOR]] = weak{{.*}} constant %struct.__tgt_offload_entry { i8* @[[CD_DTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 4, i32 0 }, section "omp_offloading_entries", align 1
maini1()85 int maini1() {
86 int a;
87 #pragma omp target map(tofrom : a)
88 {
89 a = c;
90 }
91 return 0;
92 }
93
94 // DEVICE-DAG: define weak{{.*}} void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* nonnull align {{[0-9]+}} dereferenceable{{[^,]*}}
95 // DEVICE-DAG: [[C:%.+]] = load i32, i32* [[C_ADDR]],
96 // DEVICE-DAG: store i32 [[C]], i32* %
97
98 // HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-11]](i32* nonnull align {{[0-9]+}} dereferenceable{{.*}})
99 // HOST: [[C:%.*]] = load i32, i32* @[[C_ADDR]],
100 // HOST: store i32 [[C]], i32* %
101
102 // HOST-DAG: !{i32 1, !"[[CD_ADDR]]", i32 0, i32 {{[0-9]+}}}
103 // HOST-DAG: !{i32 1, !"[[C_ADDR]]", i32 0, i32 {{[0-9]+}}}
104
105 // DEVICE: !nvvm.annotations
106 // DEVICE-DAG: !{void ()* [[C_CTOR]], !"kernel", i32 1}
107 // DEVICE-DAG: !{void ()* [[CD_CTOR]], !"kernel", i32 1}
108 // DEVICE-DAG: !{void ()* [[CD_DTOR]], !"kernel", i32 1}
109
110 #endif // HEADER
111
112