1 | // Test host codegen. |
2 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
3 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
4 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 |
5 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
6 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
7 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 |
8 | |
9 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
10 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s |
11 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
12 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
13 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
14 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s |
15 | // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} |
16 | |
17 | // Test target codegen - host bc file has to be created first. |
18 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
19 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
20 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
21 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
22 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
23 | // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
24 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
25 | // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
26 | |
27 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
28 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s |
29 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
30 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s |
31 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
32 | // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s |
33 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
34 | // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s |
35 | // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} |
36 | |
37 | // expected-no-diagnostics |
38 | #ifndef HEADER |
39 | #define HEADER |
40 | |
41 | // CHECK-DAG: [[TT:%.+]] = type { i64, i8 } |
42 | // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } |
43 | // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } |
44 | // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } |
45 | |
46 | // TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 } |
47 | |
48 | // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat |
49 | |
50 | // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4] |
51 | // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800] |
52 | // CHECK-DAG: @{{.*}} = weak constant i8 0 |
53 | |
54 | // TCHECK: @{{.+}} = weak constant [[ENTTY]] |
55 | // TCHECK: @{{.+}} = {{.*}}constant [[ENTTY]] |
56 | // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] |
57 | |
58 | // Check if offloading descriptor is created. |
59 | // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]] |
60 | // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]] |
61 | // CHECK: [[DEVBEGIN:@.+]] = extern_weak constant i8 |
62 | // CHECK: [[DEVEND:@.+]] = extern_weak constant i8 |
63 | // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }], comdat($[[REGFN]]) |
64 | // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }, comdat($[[REGFN]]) |
65 | |
66 | // Check target registration is registered as a Ctor. |
67 | // CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* @[[REGFN]], i8* bitcast (void ()* @[[REGFN]] to i8*) }] |
68 | |
69 | |
70 | template<typename tx, typename ty> |
71 | struct TT{ |
72 | tx X; |
73 | ty Y; |
74 | }; |
75 | |
76 | int global; |
77 | extern int global; |
78 | |
79 | // CHECK: define {{.*}}[[FOO:@.+]]( |
80 | int foo(int n) { |
81 | int a = 0; |
82 | short aa = 0; |
83 | float b[10]; |
84 | float bn[n]; |
85 | double c[5][10]; |
86 | double cn[5][n]; |
87 | TT<long long, char> d; |
88 | static long *plocal; |
89 | |
90 | // CHECK: [[ADD:%.+]] = add nsw i32 |
91 | // CHECK: store i32 [[ADD]], i32* [[DEVICE_CAP:%.+]], |
92 | // CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 0 |
93 | // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]], |
94 | // CHECK: store i32 [[DEV]], i32* [[GEP]], |
95 | // CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @0, i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*)) |
96 | // CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]* |
97 | // CHECK: getelementptr inbounds [4 x %struct.kmp_depend_info], [4 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
98 | // CHECK: getelementptr inbounds [4 x %struct.kmp_depend_info], [4 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 1 |
99 | // CHECK: getelementptr inbounds [4 x %struct.kmp_depend_info], [4 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 2 |
100 | // CHECK: getelementptr inbounds [4 x %struct.kmp_depend_info], [4 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 3 |
101 | // CHECK: [[DEP_START:%.+]] = getelementptr inbounds [4 x %struct.kmp_depend_info], [4 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
102 | // CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8* |
103 | // CHECK: call void @__kmpc_omp_wait_deps(%struct.ident_t* @0, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null) |
104 | // CHECK: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @0, i32 [[GTID]], i8* [[TASK]]) |
105 | // CHECK: call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]]) |
106 | // CHECK: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @0, i32 [[GTID]], i8* [[TASK]]) |
107 | #pragma omp target parallel device(global + a) depend(in: global) depend(out: a, b, cn[4]) |
108 | { |
109 | } |
110 | |
111 | // CHECK: [[ADD:%.+]] = add nsw i32 |
112 | // CHECK: store i32 [[ADD]], i32* [[DEVICE_CAP:%.+]], |
113 | |
114 | // CHECK: [[BOOL:%.+]] = icmp ne i32 %{{.+}}, 0 |
115 | // CHECK: br i1 [[BOOL]], label %[[THEN:.+]], label %[[ELSE:.+]] |
116 | // CHECK: [[THEN]]: |
117 | // CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], i32 0, i32 0 |
118 | // CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i32 0, i32 0 |
119 | // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to i[[SZ]]** |
120 | // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to i[[SZ]]** |
121 | // CHECK-DAG: store i[[SZ]]* [[BP0:%[^,]+]], i[[SZ]]** [[CBPADDR0]] |
122 | // CHECK-DAG: store i[[SZ]]* [[BP0]], i[[SZ]]** [[CPADDR0]] |
123 | |
124 | // CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 1 |
125 | // CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 1 |
126 | // CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]* |
127 | // CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* |
128 | // CHECK-DAG: store i[[SZ]] [[BP1:%[^,]+]], i[[SZ]]* [[CBPADDR1]] |
129 | // CHECK-DAG: store i[[SZ]] [[BP1]], i[[SZ]]* [[CPADDR1]] |
130 | // CHECK-DAG: getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 0 |
131 | // CHECK-DAG: getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 0 |
132 | // CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 |
133 | // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]], |
134 | // CHECK: store i32 [[DEV]], i32* [[GEP]], |
135 | |
136 | // CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @0, i32 [[GTID]], i32 1, i[[SZ]] {{104|52}}, i[[SZ]] {{16|12}}, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY1_:@.+]] to i32 (i32, i8*)*)) |
137 | // CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1_:%.+]]* |
138 | // CHECK: getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
139 | // CHECK: getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 1 |
140 | // CHECK: getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 2 |
141 | // CHECK: [[DEP_START:%.+]] = getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
142 | // CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8* |
143 | // CHECK: call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @0, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null) |
144 | // CHECK: br label %[[EXIT:.+]] |
145 | |
146 | // CHECK: [[ELSE]]: |
147 | // CHECK-NOT: getelementptr inbounds [2 x i8*], [2 x i8*]* |
148 | // CHECK: [[GEP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 |
149 | // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]], |
150 | // CHECK: store i32 [[DEV]], i32* [[GEP]], |
151 | |
152 | // CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @0, i32 [[GTID]], i32 1, i[[SZ]] {{56|28}}, i[[SZ]] {{16|12}}, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY1__:@.+]] to i32 (i32, i8*)*)) |
153 | // CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY1__:%.+]]* |
154 | // CHECK: getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
155 | // CHECK: getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 1 |
156 | // CHECK: getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 2 |
157 | // CHECK: [[DEP_START:%.+]] = getelementptr inbounds [3 x %struct.kmp_depend_info], [3 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
158 | // CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8* |
159 | // CHECK: call i32 @__kmpc_omp_task_with_deps(%struct.ident_t* @0, i32 [[GTID]], i8* [[TASK]], i32 3, i8* [[DEP]], i32 0, i8* null) |
160 | // CHECK: br label %[[EXIT:.+]] |
161 | // CHECK: [[EXIT]]: |
162 | |
163 | #pragma omp target parallel device(global + a) nowait depend(inout: global, a, bn) if(target:a) |
164 | { |
165 | static int local1; |
166 | *plocal = global; |
167 | local1 = global; |
168 | } |
169 | |
170 | // CHECK: [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @0, i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*)) |
171 | // CHECK: [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]* |
172 | // CHECK: getelementptr inbounds [1 x %struct.kmp_depend_info], [1 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
173 | // CHECK: [[DEP_START:%.+]] = getelementptr inbounds [1 x %struct.kmp_depend_info], [1 x %struct.kmp_depend_info]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
174 | // CHECK: [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8* |
175 | // CHECK: call void @__kmpc_omp_wait_deps(%struct.ident_t* @0, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null) |
176 | // CHECK: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @0, i32 [[GTID]], i8* [[TASK]]) |
177 | // CHECK: call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]]) |
178 | // CHECK: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @0, i32 [[GTID]], i8* [[TASK]]) |
179 | #pragma omp target parallel if(0) firstprivate(global) depend(out:global) |
180 | { |
181 | global += 1; |
182 | } |
183 | |
184 | return a; |
185 | } |
186 | |
187 | // Check that the offloading functions are emitted and that the arguments are |
188 | // correct and loaded correctly for the target regions in foo(). |
189 | |
190 | // CHECK: define internal void [[HVT0:@.+]]() |
191 | |
192 | // CHECK: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, [[TASK_TY0]]* noalias) |
193 | // CHECK: store void (i8*, ...)* null, void (i8*, ...)** % |
194 | // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 0 |
195 | // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]], |
196 | // CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64 |
197 | // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* null, i64* null, i32 1, i32 0) |
198 | // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
199 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] |
200 | // CHECK: [[FAIL]] |
201 | // CHECK: call void [[HVT0]]() |
202 | // CHECK-NEXT: br label %[[END]] |
203 | // CHECK: [[END]] |
204 | // CHECK: ret i32 0 |
205 | |
206 | // CHECK: define internal void [[HVT1:@.+]](i[[SZ]]* %{{.+}}, i[[SZ]] %{{.+}}) |
207 | |
208 | // CHECK: define internal{{.*}} i32 [[TASK_ENTRY1_]](i32{{.*}}, [[TASK_TY1_]]* noalias) |
209 | // CHECK: call void (i8*, ...) % |
210 | // CHECK: [[SZT:%.+]] = getelementptr inbounds [2 x i[[SZ]]], [2 x i[[SZ]]]* %{{.+}}, i[[SZ]] 0, i[[SZ]] 0 |
211 | // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 |
212 | // CHECK: [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]], |
213 | // CHECK: [[DEVICE:%.+]] = sext i32 [[DEV]] to i64 |
214 | // CHECK: [[RET:%.+]] = call i32 @__tgt_target_teams_nowait(i64 [[DEVICE]], i8* @{{[^,]+}}, i32 2, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* [[SZT]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPT]], i32 0, i32 0), i32 1, i32 0) |
215 | |
216 | // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 |
217 | // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] |
218 | // CHECK: [[FAIL]] |
219 | // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % |
220 | // CHECK: [[BP1_I32:%.+]] = load i32, i32* % |
221 | // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* |
222 | // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], |
223 | // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], |
224 | // CHECK: [[BP1:%.+]] = load i[[SZ]], i[[SZ]]* [[BP1_PTR]], |
225 | // CHECK: call void [[HVT1]](i[[SZ]]* [[BP0]], i[[SZ]] [[BP1]]) |
226 | // CHECK-NEXT: br label %[[END]] |
227 | // CHECK: [[END]] |
228 | // CHECK: ret i32 0 |
229 | |
230 | // CHECK: define internal{{.*}} i32 [[TASK_ENTRY1__]](i32{{.*}}, [[TASK_TY1__]]* noalias) |
231 | // CHECK: call void (i8*, ...) % |
232 | // CHECK: [[DEVICE_CAP:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* %{{.+}}, i32 0, i32 2 |
233 | // CHECK: [[BP0:%.+]] = load i[[SZ]]*, i[[SZ]]** % |
234 | // CHECK: [[BP1_I32:%.+]] = load i32, i32* % |
235 | // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* |
236 | // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], |
237 | // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], |
238 | // CHECK: [[BP1:%.+]] = load i[[SZ]], i[[SZ]]* [[BP1_PTR]], |
239 | // CHECK: call void [[HVT1]](i[[SZ]]* [[BP0]], i[[SZ]] [[BP1]]) |
240 | // CHECK: ret i32 0 |
241 | |
242 | // CHECK: define internal void [[HVT2:@.+]](i[[SZ]] %{{.+}}) |
243 | // Create stack storage and store argument in there. |
244 | // CHECK: [[AA_ADDR:%.+]] = alloca i[[SZ]], align |
245 | // CHECK: store i[[SZ]] %{{.+}}, i[[SZ]]* [[AA_ADDR]], align |
246 | // CHECK-64: [[AA_CADDR:%.+]] = bitcast i[[SZ]]* [[AA_ADDR]] to i32* |
247 | // CHECK-64: load i32, i32* [[AA_CADDR]], align |
248 | // CHECK-32: load i32, i32* [[AA_ADDR]], align |
249 | |
250 | // CHECK: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, [[TASK_TY2]]* noalias) |
251 | // CHECK: call void (i8*, ...) % |
252 | // CHECK: [[BP1_I32:%.+]] = load i32, i32* % |
253 | // CHECK-64: [[BP1_CAST:%.+]] = bitcast i[[SZ]]* [[BP1_PTR:%.+]] to i32* |
254 | // CHECK-64: store i32 [[BP1_I32]], i32* [[BP1_CAST]], |
255 | // CHECK-32: store i32 [[BP1_I32]], i32* [[BP1_PTR:%.+]], |
256 | // CHECK: [[BP1:%.+]] = load i[[SZ]], i[[SZ]]* [[BP1_PTR]], |
257 | // CHECK: call void [[HVT2]](i[[SZ]] [[BP1]]) |
258 | // CHECK: ret i32 0 |
259 | |
260 | |
261 | #endif |
262 | |