1 // Test target codegen - host bc file has to be created first.
2 // RUN: %clang_cc1 -DCK1 -verify -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 -DCK1 -verify -fopenmp -x c++ -triple nvptx64-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 CK1 --check-prefix CK1-64
4 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
5 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
6 // expected-no-diagnostics
7 #ifndef HEADER
8 #define HEADER
9
10 #ifdef CK1
11
12 template <typename T>
tmain(T argc)13 int tmain(T argc) {
14 #pragma omp target
15 #pragma omp teams
16 argc = 0;
17 return 0;
18 }
19
20
main(int argc,char ** argv)21 int main (int argc, char **argv) {
22 #pragma omp target
23 #pragma omp teams
24 {
25 argc = 0;
26 }
27 return tmain(argv);
28 }
29
30 // only nvptx side: do not outline teams region and do not call fork_teams
31 // CK1: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[ARGC:%.+]])
32 // CK1: {{.+}} = alloca i{{[0-9]+}}*,
33 // CK1: {{.+}} = alloca i{{[0-9]+}}*,
34 // CK1: [[ARGCADDR_PTR:%.+]] = alloca i{{[0-9]+}}*,
35 // CK1: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}},
36 // CK1: store {{.+}} 0, {{.+}},
37 // CK1: store i{{[0-9]+}} [[ARGC]], i{{[0-9]+}}* [[ARGCADDR]],
38 // CK1-64: [[CONV:%.+]] = bitcast i{{[0-9]+}}* [[ARGCADDR]] to i{{[0-9]+}}*
39 // CK1-64: store i{{[0-9]+}}* [[CONV]], i{{[0-9]+}}** [[ARGCADDR_PTR]],
40 // CK1-32: store i{{[0-9]+}}* [[ARGCADDR]], i{{[0-9]+}}** [[ARGCADDR_PTR]],
41 // CK1: [[ARGCADDR_PTR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[ARGCADDR_PTR]],
42 // CK1: store i{{[0-9]+}} 0, i{{[0-9]+}}* [[ARGCADDR_PTR_REF]],
43 // CK1-NOT: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(
44 // CK1: ret void
45 // CK1-NEXT: }
46
47 // target region in template
48 // CK1: define {{.*}}void @{{[^,]+}}(i{{.+}}** [[ARGC:%.+]])
49 // CK1: [[ARGCADDR_PTR:%.+]] = alloca i{{.+}}***,
50 // CK1: [[ARGCADDR:%.+]] = alloca i{{.+}}**,
51 // CK1: store i{{.+}}** [[ARGC]], i{{.+}}*** [[ARGCADDR]]
52 // CK1: store i8*** [[ARGCADDR]], i8**** [[ARGCADDR_PTR]],
53 // CK1: [[ARGCADDR_PTR_REF:%.+]] = load i{{.+}}**, i{{.+}}*** [[ARGCADDR_PTR]],
54 // CK1: store i{{[0-9]+}}** null, i{{[0-9]+}}*** [[ARGCADDR_PTR_REF]],
55 // CK1-NOT: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(
56 // CK1: ret void
57 // CK1-NEXT: }
58
59
60 #endif // CK1
61
62 // Test target codegen - host bc file has to be created first.
63 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
64 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple nvptx64-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 CK2 --check-prefix CK2-64
65 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
66 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
67 // expected-no-diagnostics
68 #ifdef CK2
69
70 template <typename T>
tmain(T argc)71 int tmain(T argc) {
72 int a = 10;
73 int b = 5;
74 #pragma omp target
75 #pragma omp teams num_teams(a) thread_limit(b)
76 {
77 argc = 0;
78 }
79 return 0;
80 }
81
main(int argc,char ** argv)82 int main (int argc, char **argv) {
83 int a = 20;
84 int b = 5;
85 #pragma omp target
86 #pragma omp teams num_teams(a) thread_limit(b)
87 {
88 argc = 0;
89 }
90 return tmain(argv);
91 }
92
93 // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[ARGC_IN:.+]])
94 // CK2: {{.}} = alloca i{{[0-9]+}}*,
95 // CK2: {{.}} = alloca i{{[0-9]+}}*,
96 // CK2: [[ARGCADDR_PTR:%.+]] = alloca i{{[0-9]+}}*,
97 // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}},
98 // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}},
99 // CK2: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}},
100 // CK2-NOT: {{%.+}} = call i32 @__kmpc_global_thread_num(
101 // CK2: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[AADDR]],
102 // CK2: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[BADDR]],
103 // CK2: store i{{[0-9]+}} [[ARGC_IN]], i{{[0-9]+}}* [[ARGCADDR]],
104 // CK2-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32*
105 // CK2-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32*
106 // CK2-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32*
107 // CK2-64: store i{{[0-9]+}}* [[CONV]], i{{[0-9]+}}** [[ARGCADDR_PTR]],
108 // CK2-32: store i{{[0-9]+}}* [[ARGCADDR]], i{{[0-9]+}}** [[ARGCADDR_PTR]],
109 // CK2: [[ARGCADDR_PTR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[ARGCADDR_PTR]],
110 // CK2: store i{{[0-9]+}} 0, i{{[0-9]+}}* [[ARGCADDR_PTR_REF]],
111 // CK2-NOT: {{.+}} = call i32 @__kmpc_push_num_teams(
112 // CK2-NOT: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(
113 // CK2: ret
114
115 // CK2: define {{.*}}void @{{[^,]+}}(i{{[0-9]+}} [[A_IN:%.+]], i{{[0-9]+}} [[BP:%.+]], i{{[0-9]+}}** [[ARGC:%.+]])
116 // CK2: [[ARGCADDR_PTR:%.+]] = alloca i{{[0-9]+}}***,
117 // CK2: [[AADDR:%.+]] = alloca i{{[0-9]+}},
118 // CK2: [[BADDR:%.+]] = alloca i{{[0-9]+}},
119 // CK2: [[ARGCADDR:%.+]] = alloca i{{[0-9]+}}**,
120 // CK2-NOT: {{%.+}} = call i32 @__kmpc_global_thread_num(
121 // CK2: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[AADDR]],
122 // CK2: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[BADDR]],
123 // CK2: store i{{[0-9]+}}** [[ARGC]], i{{[0-9]+}}*** [[ARGCADDR]],
124 // CK2: store i{{[0-9]+}}*** [[ARGCADDR]], i{{[0-9]+}}**** [[ARGCADDR_PTR]],
125 // CK2: [[ARGCADDR_PTR_REF:%.+]] = load i{{[0-9]+}}***, i{{[0-9]+}}**** [[ARGCADDR_PTR]],
126 // CK2: store i{{[0-9]+}}** null, i{{[0-9]+}}*** [[ARGCADDR_PTR_REF]],
127 // CK2-NOT: {{.+}} = call i32 @__kmpc_push_num_teams(
128 // CK2-NOT: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(
129 // CK2: ret void
130
131 #endif // CK2
132 #endif
133