1 // Test target codegen - host bc file has to be created first.
2 // RUN: %clang_cc1 -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 -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 CHECK --check-prefix CHECK-64
4 // RUN: %clang_cc1 -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 -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 CHECK --check-prefix CHECK-32
6 // RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -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 CHECK --check-prefix CHECK-32
7 // expected-no-diagnostics
11 // Check for the data transfer medium in shared memory to transfer the reduction list to the first warp.
12 // CHECK-DAG: [[TRANSFER_STORAGE:@.+]] = common addrspace([[SHARED_ADDRSPACE:[0-9]+]]) global [32 x i64]
14 // Check that the execution mode of all 3 target regions is set to Generic Mode.
15 // CHECK-DAG: {{@__omp_offloading_.+l27}}_exec_mode = weak constant i8 1
16 // CHECK-DAG: {{@__omp_offloading_.+l33}}_exec_mode = weak constant i8 1
17 // CHECK-DAG: {{@__omp_offloading_.+l40}}_exec_mode = weak constant i8 1
28 #pragma omp teams reduction(+: e)
34 #pragma omp teams reduction(^: c) reduction(*: d)
41 #pragma omp teams reduction(|: a) reduction(max: b)
53 a += ftemplate<char>(n);
58 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l27}}_worker()
60 // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+template.+l27]](
62 // CHECK: {{call|invoke}} void [[T1]]_worker()
64 // CHECK: call void @__kmpc_kernel_init(
66 // CHECK: store double {{[0\.e\+]+}}, double* [[E:%.+]], align
67 // CHECK: [[EV:%.+]] = load double, double* [[E]], align
68 // CHECK: [[ADD:%.+]] = fadd double [[EV]], 5
69 // CHECK: store double [[ADD]], double* [[E]], align
70 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [1 x i8*]* [[RL:%.+]], i[[SZ:32|64]] 0, i{{32|64}} 0
71 // CHECK: [[E_CAST:%.+]] = bitcast double* [[E]] to i8*
72 // CHECK: store i8* [[E_CAST]], i8** [[PTR1]], align
73 // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
74 // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_teams_reduce_nowait(i32 {{.+}}, i32 1, i[[SZ]] {{4|8}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[WARP_COPY_FN:@.+]], void (i8*, i8*, i32, i32)* [[SCRATCH_COPY_FN:@.+]], void (i8*, i8*, i32, i32, i32)* [[LOAD_REDUCE_FN:@.+]])
75 // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
76 // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
79 // CHECK: [[E_INV:%.+]] = load double, double* [[E_IN:%.+]], align
80 // CHECK: [[EV:%.+]] = load double, double* [[E]], align
81 // CHECK: [[ADD:%.+]] = fadd double [[E_INV]], [[EV]]
82 // CHECK: store double [[ADD]], double* [[E_IN]], align
83 // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
84 // CHECK: br label %[[EXIT]]
87 // CHECK: call void @__kmpc_kernel_deinit(
91 // CHECK: define internal void [[REDUCTION_FUNC:@.+]](i8*, i8*)
92 // CHECK: [[VAR_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
93 // CHECK: [[VAR_RHS_VOID:%.+]] = load i8*, i8** [[VAR_RHS_REF]],
94 // CHECK: [[VAR_RHS:%.+]] = bitcast i8* [[VAR_RHS_VOID]] to double*
96 // CHECK: [[VAR_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
97 // CHECK: [[VAR_LHS_VOID:%.+]] = load i8*, i8** [[VAR_LHS_REF]],
98 // CHECK: [[VAR_LHS:%.+]] = bitcast i8* [[VAR_LHS_VOID]] to double*
100 // CHECK: [[VAR_LHS_VAL:%.+]] = load double, double* [[VAR_LHS]],
101 // CHECK: [[VAR_RHS_VAL:%.+]] = load double, double* [[VAR_RHS]],
102 // CHECK: [[RES:%.+]] = fadd double [[VAR_LHS_VAL]], [[VAR_RHS_VAL]]
103 // CHECK: store double [[RES]], double* [[VAR_LHS]],
107 // Shuffle and reduce function
108 // CHECK: define internal void [[SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
109 // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
110 // CHECK: [[REMOTE_ELT:%.+]] = alloca double
112 // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
113 // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
114 // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
116 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
117 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
118 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
119 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
121 // CHECK: [[ELT_CAST:%.+]] = bitcast double* [[ELT]] to i64*
122 // CHECK: [[REMOTE_ELT_CAST:%.+]] = bitcast double* [[REMOTE_ELT]] to i64*
123 // CHECK: [[ELT_VAL:%.+]] = load i64, i64* [[ELT_CAST]], align
124 // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
125 // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
126 // CHECK: [[REMOTE_ELT_VAL64:%.+]] = call i64 @__kmpc_shuffle_int64(i64 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
128 // CHECK: store i64 [[REMOTE_ELT_VAL64]], i64* [[REMOTE_ELT_CAST]], align
129 // CHECK: [[REMOTE_ELT_VOID:%.+]] = bitcast double* [[REMOTE_ELT]] to i8*
130 // CHECK: store i8* [[REMOTE_ELT_VOID]], i8** [[REMOTE_ELT_REF]], align
132 // Condition to reduce
133 // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
135 // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
136 // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
137 // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
139 // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
140 // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
141 // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
142 // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
143 // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
144 // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
146 // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
147 // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
148 // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
150 // CHECK: [[DO_REDUCE]]
151 // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
152 // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
153 // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
154 // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
156 // CHECK: [[REDUCE_ELSE]]
157 // CHECK: br label {{%?}}[[REDUCE_CONT]]
159 // CHECK: [[REDUCE_CONT]]
160 // Now check if we should just copy over the remote reduction list
161 // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
162 // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
163 // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
164 // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
166 // CHECK: [[DO_COPY]]
167 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
168 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
169 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
170 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
171 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to double*
172 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
173 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load double, double* [[REMOTE_ELT]], align
174 // CHECK: store double [[REMOTE_ELT_VAL]], double* [[ELT]], align
175 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
177 // CHECK: [[COPY_ELSE]]
178 // CHECK: br label {{%?}}[[COPY_CONT]]
180 // CHECK: [[COPY_CONT]]
184 // Inter warp copy function
185 // CHECK: define internal void [[WARP_COPY_FN]](i8*, i32)
186 // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
187 // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
188 // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
189 // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
190 // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
193 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
194 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
195 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
197 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
198 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to double addrspace([[SHARED_ADDRSPACE]])*
199 // CHECK: [[ELT_VAL:%.+]] = load double, double* [[ELT]], align
200 // CHECK: store double [[ELT_VAL]], double addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
201 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
203 // CHECK: [[COPY_ELSE]]
204 // CHECK: br label {{%?}}[[COPY_CONT]]
206 // Barrier after copy to shared memory storage medium.
207 // CHECK: [[COPY_CONT]]
208 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
209 // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
210 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
213 // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
214 // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
216 // CHECK: [[DO_READ]]
217 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
218 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to double addrspace([[SHARED_ADDRSPACE]])*
219 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
220 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
221 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
222 // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load double, double addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
223 // CHECK: store double [[MEDIUM_ELT_VAL]], double* [[ELT]], align
224 // CHECK: br label {{%?}}[[READ_CONT:.+]]
226 // CHECK: [[READ_ELSE]]
227 // CHECK: br label {{%?}}[[READ_CONT]]
229 // CHECK: [[READ_CONT]]
230 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
234 // Copy to scratchpad function
235 // CHECK: define internal void [[SCRATCH_COPY_FN]](i8*, i8*, i32, i32)
236 // CHECK: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
237 // CHECK: [[SCRATCHPAD_PTR:%.+]] = load i8*, i8** {{.+}}, align
238 // CHECK-64: [[TEAM32:%.+]] = load i32, i32* {{.+}}, align
239 // CHECK-64: [[TEAM:%.+]] = sext i32 [[TEAM32]] to i64
240 // CHECK-32: [[TEAM:%.+]] = load i32, i32* {{.+}}, align
241 // CHECK-64: [[NUM_TEAMS32:%.+]] = load i32, i32* {{.+}}, align
242 // CHECK-64: [[NUM_TEAMS:%.+]] = sext i32 [[NUM_TEAMS32]] to i64
243 // CHECK-32: [[NUM_TEAMS:%.+]] = load i32, i32* {{.+}}, align
244 // CHECK: [[SCRATCHPAD:%.+]] = ptrtoint i8* [[SCRATCHPAD_PTR]] to i[[SZ]]
246 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
247 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
249 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 8, [[TEAM]]
250 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[P]]
251 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
252 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
253 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to double*
254 // CHECK: [[ELT_VAL:%.+]] = load double, double* [[ELT]], align
255 // CHECK: store double [[ELT_VAL]], double* [[SCRATCHPAD_ELT_PTR]], align
260 // Load and reduce function
261 // CHECK: define internal void [[LOAD_REDUCE_FN]](i8*, i8*, i32, i32, i32)
262 // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
263 // CHECK: [[REMOTE_ELT:%.+]] = alloca double
264 // CHECK: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
265 // CHECK: [[SCRATCHPAD_PTR:%.+]] = load i8*, i8** {{.+}}, align
266 // CHECK-64: [[TEAM32:%.+]] = load i32, i32* {{.+}}, align
267 // CHECK-64: [[TEAM:%.+]] = sext i32 [[TEAM32]] to i64
268 // CHECK-32: [[TEAM:%.+]] = load i32, i32* {{.+}}, align
269 // CHECK-64: [[NUM_TEAMS32:%.+]] = load i32, i32* {{.+}}, align
270 // CHECK-64: [[NUM_TEAMS:%.+]] = sext i32 [[NUM_TEAMS32]] to i64
271 // CHECK-32: [[NUM_TEAMS:%.+]] = load i32, i32* {{.+}}, align
272 // CHECK: [[SHOULD_REDUCE:%.+]] = load i32, i32* {{.+}}, align
273 // CHECK: [[SCRATCHPAD:%.+]] = ptrtoint i8* [[SCRATCHPAD_PTR]] to i[[SZ]]
275 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 8, [[TEAM]]
276 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[P]]
277 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
279 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
280 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to double*
281 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load double, double* [[SCRATCHPAD_ELT_PTR]], align
282 // CHECK: store double [[REMOTE_ELT_VAL]], double* [[REMOTE_ELT]], align
283 // CHECK: [[REMOTE_ELT_PTR:%.+]] = bitcast double* [[REMOTE_ELT]] to i8*
284 // CHECK: store i8* [[REMOTE_ELT_PTR]], i8** [[REMOTE_ELT_REF]], align
286 // CHECK: [[REDUCE:%.+]] = icmp ne i32 [[SHOULD_REDUCE]], 0
287 // CHECK: br i1 [[REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
289 // CHECK: [[DO_REDUCE]]
290 // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
291 // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
292 // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
293 // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
295 // Copy element from remote reduce list
296 // CHECK: [[REDUCE_ELSE]]
297 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
298 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
299 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
300 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
301 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to double*
302 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
303 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load double, double* [[REMOTE_ELT]], align
304 // CHECK: store double [[REMOTE_ELT_VAL]], double* [[ELT]], align
305 // CHECK: br label {{%?}}[[REDUCE_CONT]]
307 // CHECK: [[REDUCE_CONT]]
310 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l33}}_worker()
312 // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+template.+l33]](
314 // CHECK: {{call|invoke}} void [[T2]]_worker()
316 // CHECK: call void @__kmpc_kernel_init(
318 // CHECK: store float {{1\.[0e\+]+}}, float* [[D:%.+]], align
319 // CHECK: [[C_VAL:%.+]] = load i8, i8* [[C:%.+]], align
320 // CHECK: [[CONV:%.+]] = sext i8 [[C_VAL]] to i32
321 // CHECK: [[XOR:%.+]] = xor i32 [[CONV]], 2
322 // CHECK: [[TRUNC:%.+]] = trunc i32 [[XOR]] to i8
323 // CHECK: store i8 [[TRUNC]], i8* [[C]], align
324 // CHECK: [[DV:%.+]] = load float, float* [[D]], align
325 // CHECK: [[MUL:%.+]] = fmul float [[DV]], {{[0-9e\.\+]+}}
326 // CHECK: store float [[MUL]], float* [[D]], align
327 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [2 x i8*]* [[RL:%.+]], i[[SZ]] 0, i[[SZ]] 0
328 // CHECK: store i8* [[C]], i8** [[PTR1]], align
329 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RL]], i[[SZ]] 0, i[[SZ]] 1
330 // CHECK: [[D_CAST:%.+]] = bitcast float* [[D]] to i8*
331 // CHECK: store i8* [[D_CAST]], i8** [[PTR2]], align
332 // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
333 // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_teams_reduce_nowait(i32 {{.+}}, i32 2, i[[SZ]] {{8|16}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[WARP_COPY_FN:@.+]], void (i8*, i8*, i32, i32)* [[SCRATCH_COPY_FN:@.+]], void (i8*, i8*, i32, i32, i32)* [[LOAD_REDUCE_FN:@.+]])
334 // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
335 // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
337 // CHECK: [[IFLABEL]]
338 // CHECK: [[C_INV8:%.+]] = load i8, i8* [[C_IN:%.+]], align
339 // CHECK: [[C_INV:%.+]] = sext i8 [[C_INV8]] to i32
340 // CHECK: [[CV8:%.+]] = load i8, i8* [[C]], align
341 // CHECK: [[CV:%.+]] = sext i8 [[CV8]] to i32
342 // CHECK: [[XOR:%.+]] = xor i32 [[C_INV]], [[CV]]
343 // CHECK: [[TRUNC:%.+]] = trunc i32 [[XOR]] to i8
344 // CHECK: store i8 [[TRUNC]], i8* [[C_IN]], align
345 // CHECK: [[D_INV:%.+]] = load float, float* [[D_IN:%.+]], align
346 // CHECK: [[DV:%.+]] = load float, float* [[D]], align
347 // CHECK: [[MUL:%.+]] = fmul float [[D_INV]], [[DV]]
348 // CHECK: store float [[MUL]], float* [[D_IN]], align
349 // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
350 // CHECK: br label %[[EXIT]]
353 // CHECK: call void @__kmpc_kernel_deinit(
356 // Reduction function
357 // CHECK: define internal void [[REDUCTION_FUNC:@.+]](i8*, i8*)
358 // CHECK: [[VAR1_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
359 // CHECK: [[VAR1_RHS:%.+]] = load i8*, i8** [[VAR1_RHS_REF]],
361 // CHECK: [[VAR1_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
362 // CHECK: [[VAR1_LHS:%.+]] = load i8*, i8** [[VAR1_LHS_REF]],
364 // CHECK: [[VAR2_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS]], i[[SZ]] 0, i[[SZ]] 1
365 // CHECK: [[VAR2_RHS_VOID:%.+]] = load i8*, i8** [[VAR2_RHS_REF]],
366 // CHECK: [[VAR2_RHS:%.+]] = bitcast i8* [[VAR2_RHS_VOID]] to float*
368 // CHECK: [[VAR2_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS]], i[[SZ]] 0, i[[SZ]] 1
369 // CHECK: [[VAR2_LHS_VOID:%.+]] = load i8*, i8** [[VAR2_LHS_REF]],
370 // CHECK: [[VAR2_LHS:%.+]] = bitcast i8* [[VAR2_LHS_VOID]] to float*
372 // CHECK: [[VAR1_LHS_VAL8:%.+]] = load i8, i8* [[VAR1_LHS]],
373 // CHECK: [[VAR1_LHS_VAL:%.+]] = sext i8 [[VAR1_LHS_VAL8]] to i32
374 // CHECK: [[VAR1_RHS_VAL8:%.+]] = load i8, i8* [[VAR1_RHS]],
375 // CHECK: [[VAR1_RHS_VAL:%.+]] = sext i8 [[VAR1_RHS_VAL8]] to i32
376 // CHECK: [[XOR:%.+]] = xor i32 [[VAR1_LHS_VAL]], [[VAR1_RHS_VAL]]
377 // CHECK: [[RES:%.+]] = trunc i32 [[XOR]] to i8
378 // CHECK: store i8 [[RES]], i8* [[VAR1_LHS]],
380 // CHECK: [[VAR2_LHS_VAL:%.+]] = load float, float* [[VAR2_LHS]],
381 // CHECK: [[VAR2_RHS_VAL:%.+]] = load float, float* [[VAR2_RHS]],
382 // CHECK: [[RES:%.+]] = fmul float [[VAR2_LHS_VAL]], [[VAR2_RHS_VAL]]
383 // CHECK: store float [[RES]], float* [[VAR2_LHS]],
387 // Shuffle and reduce function
388 // CHECK: define internal void [[SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
389 // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
390 // CHECK: [[REMOTE_ELT1:%.+]] = alloca i8
391 // CHECK: [[REMOTE_ELT2:%.+]] = alloca float
393 // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
394 // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
395 // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
397 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
398 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
399 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
400 // CHECK: [[ELT_VAL:%.+]] = load i8, i8* [[ELT_VOID]], align
402 // CHECK: [[ELT_CAST:%.+]] = sext i8 [[ELT_VAL]] to i32
403 // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
404 // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
405 // CHECK: [[REMOTE_ELT1_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_CAST]], i16 [[LANEOFFSET]], i16 [[WS]])
406 // CHECK: [[REMOTE_ELT1_VAL:%.+]] = trunc i32 [[REMOTE_ELT1_VAL32]] to i8
408 // CHECK: store i8 [[REMOTE_ELT1_VAL]], i8* [[REMOTE_ELT1]], align
409 // CHECK: store i8* [[REMOTE_ELT1]], i8** [[REMOTE_ELT_REF]], align
411 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
412 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
413 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
414 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
416 // CHECK: [[ELT_CAST:%.+]] = bitcast float* [[ELT]] to i32*
417 // CHECK: [[REMOTE_ELT2_CAST:%.+]] = bitcast float* [[REMOTE_ELT2]] to i32*
418 // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT_CAST]], align
419 // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
420 // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
421 // CHECK: [[REMOTE_ELT2_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
423 // CHECK: store i32 [[REMOTE_ELT2_VAL32]], i32* [[REMOTE_ELT2_CAST]], align
424 // CHECK: [[REMOTE_ELT2C:%.+]] = bitcast float* [[REMOTE_ELT2]] to i8*
425 // CHECK: store i8* [[REMOTE_ELT2C]], i8** [[REMOTE_ELT_REF]], align
427 // Condition to reduce
428 // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
430 // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
431 // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
432 // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
434 // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
435 // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
436 // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
437 // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
438 // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
439 // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
441 // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
442 // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
443 // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
445 // CHECK: [[DO_REDUCE]]
446 // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
447 // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
448 // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
449 // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
451 // CHECK: [[REDUCE_ELSE]]
452 // CHECK: br label {{%?}}[[REDUCE_CONT]]
454 // CHECK: [[REDUCE_CONT]]
455 // Now check if we should just copy over the remote reduction list
456 // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
457 // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
458 // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
459 // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
461 // CHECK: [[DO_COPY]]
462 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
463 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
464 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
465 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
466 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i8, i8* [[REMOTE_ELT_VOID]], align
467 // CHECK: store i8 [[REMOTE_ELT_VAL]], i8* [[ELT_VOID]], align
469 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
470 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
471 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
472 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
473 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to float*
474 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
475 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load float, float* [[REMOTE_ELT]], align
476 // CHECK: store float [[REMOTE_ELT_VAL]], float* [[ELT]], align
477 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
479 // CHECK: [[COPY_ELSE]]
480 // CHECK: br label {{%?}}[[COPY_CONT]]
482 // CHECK: [[COPY_CONT]]
486 // Inter warp copy function
487 // CHECK: define internal void [[WARP_COPY_FN]](i8*, i32)
488 // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
489 // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
490 // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
491 // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
492 // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
495 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
496 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
498 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
499 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i8 addrspace([[SHARED_ADDRSPACE]])*
500 // CHECK: [[ELT_VAL:%.+]] = load i8, i8* [[ELT_VOID]], align
501 // CHECK: store i8 [[ELT_VAL]], i8 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
502 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
504 // CHECK: [[COPY_ELSE]]
505 // CHECK: br label {{%?}}[[COPY_CONT]]
507 // Barrier after copy to shared memory storage medium.
508 // CHECK: [[COPY_CONT]]
509 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
510 // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
511 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
514 // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
515 // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
517 // CHECK: [[DO_READ]]
518 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
519 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i8 addrspace([[SHARED_ADDRSPACE]])*
520 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
521 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
522 // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load i8, i8 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
523 // CHECK: store i8 [[MEDIUM_ELT_VAL]], i8* [[ELT_VOID]], align
524 // CHECK: br label {{%?}}[[READ_CONT:.+]]
526 // CHECK: [[READ_ELSE]]
527 // CHECK: br label {{%?}}[[READ_CONT]]
529 // CHECK: [[READ_CONT]]
530 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
531 // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
532 // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
535 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
536 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
537 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
539 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
540 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to float addrspace([[SHARED_ADDRSPACE]])*
541 // CHECK: [[ELT_VAL:%.+]] = load float, float* [[ELT]], align
542 // CHECK: store float [[ELT_VAL]], float addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
543 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
545 // CHECK: [[COPY_ELSE]]
546 // CHECK: br label {{%?}}[[COPY_CONT]]
548 // Barrier after copy to shared memory storage medium.
549 // CHECK: [[COPY_CONT]]
550 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
551 // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
552 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
555 // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
556 // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
558 // CHECK: [[DO_READ]]
559 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
560 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to float addrspace([[SHARED_ADDRSPACE]])*
561 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
562 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
563 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
564 // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load float, float addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
565 // CHECK: store float [[MEDIUM_ELT_VAL]], float* [[ELT]], align
566 // CHECK: br label {{%?}}[[READ_CONT:.+]]
568 // CHECK: [[READ_ELSE]]
569 // CHECK: br label {{%?}}[[READ_CONT]]
571 // CHECK: [[READ_CONT]]
572 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
576 // Copy to scratchpad function
577 // CHECK: define internal void [[SCRATCH_COPY_FN]](i8*, i8*, i32, i32)
578 // CHECK: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
579 // CHECK: [[SCRATCHPAD_PTR:%.+]] = load i8*, i8** {{.+}}, align
580 // CHECK-64: [[TEAM32:%.+]] = load i32, i32* {{.+}}, align
581 // CHECK-64: [[TEAM:%.+]] = sext i32 [[TEAM32]] to i64
582 // CHECK-32: [[TEAM:%.+]] = load i32, i32* {{.+}}, align
583 // CHECK-64: [[NUM_TEAMS32:%.+]] = load i32, i32* {{.+}}, align
584 // CHECK-64: [[NUM_TEAMS:%.+]] = sext i32 [[NUM_TEAMS32]] to i64
585 // CHECK-32: [[NUM_TEAMS:%.+]] = load i32, i32* {{.+}}, align
586 // CHECK: [[SCRATCHPAD:%.+]] = ptrtoint i8* [[SCRATCHPAD_PTR]] to i[[SZ]]
588 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
589 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
591 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 1, [[TEAM]]
592 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[P]]
593 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
594 // CHECK: [[ELT_VAL:%.+]] = load i8, i8* [[ELT_VOID]], align
595 // CHECK: store i8 [[ELT_VAL]], i8* [[SCRATCHPAD_ELT_PTR]], align
597 // CHECK: [[OF:%.+]] = mul nuw i[[SZ]] [[NUM_TEAMS]], 1
598 // CHECK: [[POS1:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[OF]]
599 // CHECK: [[POS2:%.+]] = sub nuw i[[SZ]] [[POS1]], 1
600 // CHECK: [[POS3:%.+]] = udiv i[[SZ]] [[POS2]], 256
601 // CHECK: [[POS4:%.+]] = add nuw i[[SZ]] [[POS3]], 1
602 // CHECK: [[SCRATCHPAD_NEXT:%.+]] = mul nuw i[[SZ]] [[POS4]], 256
604 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
605 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
607 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 4, [[TEAM]]
608 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD_NEXT]], [[P]]
609 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
610 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
611 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to float*
612 // CHECK: [[ELT_VAL:%.+]] = load float, float* [[ELT]], align
613 // CHECK: store float [[ELT_VAL]], float* [[SCRATCHPAD_ELT_PTR]], align
618 // Load and reduce function
619 // CHECK: define internal void [[LOAD_REDUCE_FN]](i8*, i8*, i32, i32, i32)
620 // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
621 // CHECK: [[REMOTE_ELT1:%.+]] = alloca i8
622 // CHECK: [[REMOTE_ELT2:%.+]] = alloca float
623 // CHECK: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
624 // CHECK: [[SCRATCHPAD_PTR:%.+]] = load i8*, i8** {{.+}}, align
625 // CHECK-64: [[TEAM32:%.+]] = load i32, i32* {{.+}}, align
626 // CHECK-64: [[TEAM:%.+]] = sext i32 [[TEAM32]] to i64
627 // CHECK-32: [[TEAM:%.+]] = load i32, i32* {{.+}}, align
628 // CHECK-64: [[NUM_TEAMS32:%.+]] = load i32, i32* {{.+}}, align
629 // CHECK-64: [[NUM_TEAMS:%.+]] = sext i32 [[NUM_TEAMS32]] to i64
630 // CHECK-32: [[NUM_TEAMS:%.+]] = load i32, i32* {{.+}}, align
631 // CHECK: [[SHOULD_REDUCE:%.+]] = load i32, i32* {{.+}}, align
632 // CHECK: [[SCRATCHPAD:%.+]] = ptrtoint i8* [[SCRATCHPAD_PTR]] to i[[SZ]]
634 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 1, [[TEAM]]
635 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[P]]
636 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
638 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
639 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i8, i8* [[SCRATCHPAD_ELT_PTR_VOID]], align
640 // CHECK: store i8 [[REMOTE_ELT_VAL]], i8* [[REMOTE_ELT1]], align
641 // CHECK: store i8* [[REMOTE_ELT1]], i8** [[REMOTE_ELT_REF]], align
643 // CHECK: [[OF:%.+]] = mul nuw i[[SZ]] [[NUM_TEAMS]], 1
644 // CHECK: [[POS1:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[OF]]
645 // CHECK: [[POS2:%.+]] = sub nuw i[[SZ]] [[POS1]], 1
646 // CHECK: [[POS3:%.+]] = udiv i[[SZ]] [[POS2]], 256
647 // CHECK: [[POS4:%.+]] = add nuw i[[SZ]] [[POS3]], 1
648 // CHECK: [[SCRATCHPAD_NEXT:%.+]] = mul nuw i[[SZ]] [[POS4]], 256
650 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 4, [[TEAM]]
651 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD_NEXT]], [[P]]
652 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
654 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
655 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to float*
656 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load float, float* [[SCRATCHPAD_ELT_PTR]], align
657 // CHECK: store float [[REMOTE_ELT_VAL]], float* [[REMOTE_ELT2]], align
658 // CHECK: [[REMOTE_ELT_PTR:%.+]] = bitcast float* [[REMOTE_ELT2]] to i8*
659 // CHECK: store i8* [[REMOTE_ELT_PTR]], i8** [[REMOTE_ELT_REF]], align
661 // CHECK: [[REDUCE:%.+]] = icmp ne i32 [[SHOULD_REDUCE]], 0
662 // CHECK: br i1 [[REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
664 // CHECK: [[DO_REDUCE]]
665 // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
666 // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
667 // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
668 // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
670 // Copy element from remote reduce list
671 // CHECK: [[REDUCE_ELSE]]
672 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
673 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
674 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
675 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
676 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i8, i8* [[REMOTE_ELT_VOID]], align
677 // CHECK: store i8 [[REMOTE_ELT_VAL]], i8* [[ELT_VOID]], align
679 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
680 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
681 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
682 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
683 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to float*
684 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
685 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load float, float* [[REMOTE_ELT]], align
686 // CHECK: store float [[REMOTE_ELT_VAL]], float* [[ELT]], align
687 // CHECK: br label {{%?}}[[REDUCE_CONT]]
689 // CHECK: [[REDUCE_CONT]]
692 // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l40}}_worker()
694 // CHECK: define {{.*}}void [[T3:@__omp_offloading_.+template.+l40]](
696 // CHECK: {{call|invoke}} void [[T3]]_worker()
698 // CHECK: call void @__kmpc_kernel_init(
700 // CHECK: store i32 0, i32* [[A:%.+]], align
701 // CHECK: store i16 -32768, i16* [[B:%.+]], align
702 // CHECK: [[A_VAL:%.+]] = load i32, i32* [[A:%.+]], align
703 // CHECK: [[OR:%.+]] = or i32 [[A_VAL]], 1
704 // CHECK: store i32 [[OR]], i32* [[A]], align
705 // CHECK: [[BV16:%.+]] = load i16, i16* [[B]], align
706 // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
707 // CHECK: [[CMP:%.+]] = icmp sgt i32 99, [[BV]]
708 // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
711 // CHECK: br label {{%?}}[[MAX_CONT:.+]]
713 // CHECK: [[MAX_ELSE]]
714 // CHECK: [[BV:%.+]] = load i16, i16* [[B]], align
715 // CHECK: [[MAX:%.+]] = sext i16 [[BV]] to i32
716 // CHECK: br label {{%?}}[[MAX_CONT]]
718 // CHECK: [[MAX_CONT]]
719 // CHECK: [[B_LVALUE:%.+]] = phi i32 [ 99, %[[DO_MAX]] ], [ [[MAX]], %[[MAX_ELSE]] ]
720 // CHECK: [[TRUNC:%.+]] = trunc i32 [[B_LVALUE]] to i16
721 // CHECK: store i16 [[TRUNC]], i16* [[B]], align
722 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [2 x i8*]* [[RL:%.+]], i[[SZ]] 0, i[[SZ]] 0
723 // CHECK: [[A_CAST:%.+]] = bitcast i32* [[A]] to i8*
724 // CHECK: store i8* [[A_CAST]], i8** [[PTR1]], align
725 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RL]], i[[SZ]] 0, i[[SZ]] 1
726 // CHECK: [[B_CAST:%.+]] = bitcast i16* [[B]] to i8*
727 // CHECK: store i8* [[B_CAST]], i8** [[PTR2]], align
728 // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
729 // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_teams_reduce_nowait(i32 {{.+}}, i32 2, i[[SZ]] {{8|16}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[WARP_COPY_FN:@.+]], void (i8*, i8*, i32, i32)* [[SCRATCH_COPY_FN:@.+]], void (i8*, i8*, i32, i32, i32)* [[LOAD_REDUCE_FN:@.+]])
730 // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
731 // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
733 // CHECK: [[IFLABEL]]
734 // CHECK: [[A_INV:%.+]] = load i32, i32* [[A_IN:%.+]], align
735 // CHECK: [[AV:%.+]] = load i32, i32* [[A]], align
736 // CHECK: [[OR:%.+]] = or i32 [[A_INV]], [[AV]]
737 // CHECK: store i32 [[OR]], i32* [[A_IN]], align
738 // CHECK: [[B_INV16:%.+]] = load i16, i16* [[B_IN:%.+]], align
739 // CHECK: [[B_INV:%.+]] = sext i16 [[B_INV16]] to i32
740 // CHECK: [[BV16:%.+]] = load i16, i16* [[B]], align
741 // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
742 // CHECK: [[CMP:%.+]] = icmp sgt i32 [[B_INV]], [[BV]]
743 // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
746 // CHECK: [[MAX1:%.+]] = load i16, i16* [[B_IN]], align
747 // CHECK: br label {{%?}}[[MAX_CONT:.+]]
749 // CHECK: [[MAX_ELSE]]
750 // CHECK: [[MAX2:%.+]] = load i16, i16* [[B]], align
751 // CHECK: br label {{%?}}[[MAX_CONT]]
753 // CHECK: [[MAX_CONT]]
754 // CHECK: [[B_MAX:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
755 // CHECK: store i16 [[B_MAX]], i16* [[B_IN]], align
756 // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
757 // CHECK: br label %[[EXIT]]
760 // CHECK: call void @__kmpc_kernel_deinit(
763 // Reduction function
764 // CHECK: define internal void [[REDUCTION_FUNC:@.+]](i8*, i8*)
765 // CHECK: [[VAR1_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
766 // CHECK: [[VAR1_RHS_VOID:%.+]] = load i8*, i8** [[VAR1_RHS_REF]],
767 // CHECK: [[VAR1_RHS:%.+]] = bitcast i8* [[VAR1_RHS_VOID]] to i32*
769 // CHECK: [[VAR1_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
770 // CHECK: [[VAR1_LHS_VOID:%.+]] = load i8*, i8** [[VAR1_LHS_REF]],
771 // CHECK: [[VAR1_LHS:%.+]] = bitcast i8* [[VAR1_LHS_VOID]] to i32*
773 // CHECK: [[VAR2_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS]], i[[SZ]] 0, i[[SZ]] 1
774 // CHECK: [[VAR2_RHS_VOID:%.+]] = load i8*, i8** [[VAR2_RHS_REF]],
775 // CHECK: [[VAR2_RHS:%.+]] = bitcast i8* [[VAR2_RHS_VOID]] to i16*
777 // CHECK: [[VAR2_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS]], i[[SZ]] 0, i[[SZ]] 1
778 // CHECK: [[VAR2_LHS_VOID:%.+]] = load i8*, i8** [[VAR2_LHS_REF]],
779 // CHECK: [[VAR2_LHS:%.+]] = bitcast i8* [[VAR2_LHS_VOID]] to i16*
781 // CHECK: [[VAR1_LHS_VAL:%.+]] = load i32, i32* [[VAR1_LHS]],
782 // CHECK: [[VAR1_RHS_VAL:%.+]] = load i32, i32* [[VAR1_RHS]],
783 // CHECK: [[OR:%.+]] = or i32 [[VAR1_LHS_VAL]], [[VAR1_RHS_VAL]]
784 // CHECK: store i32 [[OR]], i32* [[VAR1_LHS]],
786 // CHECK: [[VAR2_LHS_VAL16:%.+]] = load i16, i16* [[VAR2_LHS]],
787 // CHECK: [[VAR2_LHS_VAL:%.+]] = sext i16 [[VAR2_LHS_VAL16]] to i32
788 // CHECK: [[VAR2_RHS_VAL16:%.+]] = load i16, i16* [[VAR2_RHS]],
789 // CHECK: [[VAR2_RHS_VAL:%.+]] = sext i16 [[VAR2_RHS_VAL16]] to i32
791 // CHECK: [[CMP:%.+]] = icmp sgt i32 [[VAR2_LHS_VAL]], [[VAR2_RHS_VAL]]
792 // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
795 // CHECK: [[MAX1:%.+]] = load i16, i16* [[VAR2_LHS]], align
796 // CHECK: br label {{%?}}[[MAX_CONT:.+]]
798 // CHECK: [[MAX_ELSE]]
799 // CHECK: [[MAX2:%.+]] = load i16, i16* [[VAR2_RHS]], align
800 // CHECK: br label {{%?}}[[MAX_CONT]]
802 // CHECK: [[MAX_CONT]]
803 // CHECK: [[MAXV:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
804 // CHECK: store i16 [[MAXV]], i16* [[VAR2_LHS]],
808 // Shuffle and reduce function
809 // CHECK: define internal void [[SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
810 // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
811 // CHECK: [[REMOTE_ELT1:%.+]] = alloca i32
812 // CHECK: [[REMOTE_ELT2:%.+]] = alloca i16
814 // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
815 // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
816 // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
818 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
819 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
820 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
821 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
822 // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
824 // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
825 // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
826 // CHECK: [[REMOTE_ELT1_VAL:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
828 // CHECK: store i32 [[REMOTE_ELT1_VAL]], i32* [[REMOTE_ELT1]], align
829 // CHECK: [[REMOTE_ELT1C:%.+]] = bitcast i32* [[REMOTE_ELT1]] to i8*
830 // CHECK: store i8* [[REMOTE_ELT1C]], i8** [[REMOTE_ELT_REF]], align
832 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
833 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
834 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
835 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
836 // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
838 // CHECK: [[ELT_CAST:%.+]] = sext i16 [[ELT_VAL]] to i32
839 // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
840 // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
841 // CHECK: [[REMOTE_ELT2_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_CAST]], i16 [[LANEOFFSET]], i16 [[WS]])
842 // CHECK: [[REMOTE_ELT2_VAL:%.+]] = trunc i32 [[REMOTE_ELT2_VAL32]] to i16
844 // CHECK: store i16 [[REMOTE_ELT2_VAL]], i16* [[REMOTE_ELT2]], align
845 // CHECK: [[REMOTE_ELT2C:%.+]] = bitcast i16* [[REMOTE_ELT2]] to i8*
846 // CHECK: store i8* [[REMOTE_ELT2C]], i8** [[REMOTE_ELT_REF]], align
848 // Condition to reduce
849 // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
851 // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
852 // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
853 // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
855 // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
856 // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
857 // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
858 // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
859 // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
860 // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
862 // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
863 // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
864 // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
866 // CHECK: [[DO_REDUCE]]
867 // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
868 // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
869 // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
870 // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
872 // CHECK: [[REDUCE_ELSE]]
873 // CHECK: br label {{%?}}[[REDUCE_CONT]]
875 // CHECK: [[REDUCE_CONT]]
876 // Now check if we should just copy over the remote reduction list
877 // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
878 // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
879 // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
880 // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
882 // CHECK: [[DO_COPY]]
883 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
884 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
885 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
886 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
887 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i32*
888 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
889 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i32, i32* [[REMOTE_ELT]], align
890 // CHECK: store i32 [[REMOTE_ELT_VAL]], i32* [[ELT]], align
892 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
893 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
894 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
895 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
896 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i16*
897 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
898 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i16, i16* [[REMOTE_ELT]], align
899 // CHECK: store i16 [[REMOTE_ELT_VAL]], i16* [[ELT]], align
900 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
902 // CHECK: [[COPY_ELSE]]
903 // CHECK: br label {{%?}}[[COPY_CONT]]
905 // CHECK: [[COPY_CONT]]
909 // Inter warp copy function
910 // CHECK: define internal void [[WARP_COPY_FN]](i8*, i32)
911 // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
912 // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
913 // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
914 // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
915 // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
918 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
919 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
920 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
922 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
923 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i32 addrspace([[SHARED_ADDRSPACE]])*
924 // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
925 // CHECK: store i32 [[ELT_VAL]], i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
926 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
928 // CHECK: [[COPY_ELSE]]
929 // CHECK: br label {{%?}}[[COPY_CONT]]
931 // Barrier after copy to shared memory storage medium.
932 // CHECK: [[COPY_CONT]]
933 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
934 // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
935 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
938 // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
939 // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
941 // CHECK: [[DO_READ]]
942 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
943 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i32 addrspace([[SHARED_ADDRSPACE]])*
944 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
945 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
946 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
947 // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load i32, i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
948 // CHECK: store i32 [[MEDIUM_ELT_VAL]], i32* [[ELT]], align
949 // CHECK: br label {{%?}}[[READ_CONT:.+]]
951 // CHECK: [[READ_ELSE]]
952 // CHECK: br label {{%?}}[[READ_CONT]]
954 // CHECK: [[READ_CONT]]
955 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
956 // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
957 // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
960 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
961 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
962 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
964 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
965 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i16 addrspace([[SHARED_ADDRSPACE]])*
966 // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
967 // CHECK: store i16 [[ELT_VAL]], i16 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
968 // CHECK: br label {{%?}}[[COPY_CONT:.+]]
970 // CHECK: [[COPY_ELSE]]
971 // CHECK: br label {{%?}}[[COPY_CONT]]
973 // Barrier after copy to shared memory storage medium.
974 // CHECK: [[COPY_CONT]]
975 // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
976 // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
977 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
980 // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
981 // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
983 // CHECK: [[DO_READ]]
984 // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
985 // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i16 addrspace([[SHARED_ADDRSPACE]])*
986 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
987 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
988 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
989 // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load i16, i16 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
990 // CHECK: store i16 [[MEDIUM_ELT_VAL]], i16* [[ELT]], align
991 // CHECK: br label {{%?}}[[READ_CONT:.+]]
993 // CHECK: [[READ_ELSE]]
994 // CHECK: br label {{%?}}[[READ_CONT]]
996 // CHECK: [[READ_CONT]]
997 // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
1001 // Copy to scratchpad function
1002 // CHECK: define internal void [[SCRATCH_COPY_FN]](i8*, i8*, i32, i32)
1003 // CHECK: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
1004 // CHECK: [[SCRATCHPAD_PTR:%.+]] = load i8*, i8** {{.+}}, align
1005 // CHECK-64: [[TEAM32:%.+]] = load i32, i32* {{.+}}, align
1006 // CHECK-64: [[TEAM:%.+]] = sext i32 [[TEAM32]] to i64
1007 // CHECK-32: [[TEAM:%.+]] = load i32, i32* {{.+}}, align
1008 // CHECK-64: [[NUM_TEAMS32:%.+]] = load i32, i32* {{.+}}, align
1009 // CHECK-64: [[NUM_TEAMS:%.+]] = sext i32 [[NUM_TEAMS32]] to i64
1010 // CHECK-32: [[NUM_TEAMS:%.+]] = load i32, i32* {{.+}}, align
1011 // CHECK: [[SCRATCHPAD:%.+]] = ptrtoint i8* [[SCRATCHPAD_PTR]] to i[[SZ]]
1013 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
1014 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
1016 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 4, [[TEAM]]
1017 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[P]]
1018 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
1019 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
1020 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to i32*
1021 // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
1022 // CHECK: store i32 [[ELT_VAL]], i32* [[SCRATCHPAD_ELT_PTR]], align
1024 // CHECK: [[OF:%.+]] = mul nuw i[[SZ]] [[NUM_TEAMS]], 4
1025 // CHECK: [[POS1:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[OF]]
1026 // CHECK: [[POS2:%.+]] = sub nuw i[[SZ]] [[POS1]], 1
1027 // CHECK: [[POS3:%.+]] = udiv i[[SZ]] [[POS2]], 256
1028 // CHECK: [[POS4:%.+]] = add nuw i[[SZ]] [[POS3]], 1
1029 // CHECK: [[SCRATCHPAD_NEXT:%.+]] = mul nuw i[[SZ]] [[POS4]], 256
1031 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
1032 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
1034 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 2, [[TEAM]]
1035 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD_NEXT]], [[P]]
1036 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
1037 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
1038 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to i16*
1039 // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
1040 // CHECK: store i16 [[ELT_VAL]], i16* [[SCRATCHPAD_ELT_PTR]], align
1045 // Load and reduce function
1046 // CHECK: define internal void [[LOAD_REDUCE_FN]](i8*, i8*, i32, i32, i32)
1047 // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
1048 // CHECK: [[REMOTE_ELT1:%.+]] = alloca i32
1049 // CHECK: [[REMOTE_ELT2:%.+]] = alloca i16
1050 // CHECK: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
1051 // CHECK: [[SCRATCHPAD_PTR:%.+]] = load i8*, i8** {{.+}}, align
1052 // CHECK-64: [[TEAM32:%.+]] = load i32, i32* {{.+}}, align
1053 // CHECK-64: [[TEAM:%.+]] = sext i32 [[TEAM32]] to i64
1054 // CHECK-32: [[TEAM:%.+]] = load i32, i32* {{.+}}, align
1055 // CHECK-64: [[NUM_TEAMS32:%.+]] = load i32, i32* {{.+}}, align
1056 // CHECK-64: [[NUM_TEAMS:%.+]] = sext i32 [[NUM_TEAMS32]] to i64
1057 // CHECK-32: [[NUM_TEAMS:%.+]] = load i32, i32* {{.+}}, align
1058 // CHECK: [[SHOULD_REDUCE:%.+]] = load i32, i32* {{.+}}, align
1059 // CHECK: [[SCRATCHPAD:%.+]] = ptrtoint i8* [[SCRATCHPAD_PTR]] to i[[SZ]]
1061 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 4, [[TEAM]]
1062 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[P]]
1063 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
1065 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
1066 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to i32*
1067 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i32, i32* [[SCRATCHPAD_ELT_PTR]], align
1068 // CHECK: store i32 [[REMOTE_ELT_VAL]], i32* [[REMOTE_ELT1]], align
1069 // CHECK: [[REMOTE_ELT1_PTR:%.+]] = bitcast i32* [[REMOTE_ELT1]] to i8*
1070 // CHECK: store i8* [[REMOTE_ELT1_PTR]], i8** [[REMOTE_ELT_REF]], align
1072 // CHECK: [[OF:%.+]] = mul nuw i[[SZ]] [[NUM_TEAMS]], 4
1073 // CHECK: [[POS1:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD]], [[OF]]
1074 // CHECK: [[POS2:%.+]] = sub nuw i[[SZ]] [[POS1]], 1
1075 // CHECK: [[POS3:%.+]] = udiv i[[SZ]] [[POS2]], 256
1076 // CHECK: [[POS4:%.+]] = add nuw i[[SZ]] [[POS3]], 1
1077 // CHECK: [[SCRATCHPAD_NEXT:%.+]] = mul nuw i[[SZ]] [[POS4]], 256
1079 // CHECK: [[P:%.+]] = mul nuw i[[SZ]] 2, [[TEAM]]
1080 // CHECK: [[SCRATCHPAD_ELT_PTR64:%.+]] = add nuw i[[SZ]] [[SCRATCHPAD_NEXT]], [[P]]
1081 // CHECK: [[SCRATCHPAD_ELT_PTR_VOID:%.+]] = inttoptr i[[SZ]] [[SCRATCHPAD_ELT_PTR64]] to i8*
1083 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
1084 // CHECK: [[SCRATCHPAD_ELT_PTR:%.+]] = bitcast i8* [[SCRATCHPAD_ELT_PTR_VOID]] to i16*
1085 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i16, i16* [[SCRATCHPAD_ELT_PTR]], align
1086 // CHECK: store i16 [[REMOTE_ELT_VAL]], i16* [[REMOTE_ELT2]], align
1087 // CHECK: [[REMOTE_ELT_PTR:%.+]] = bitcast i16* [[REMOTE_ELT2]] to i8*
1088 // CHECK: store i8* [[REMOTE_ELT_PTR]], i8** [[REMOTE_ELT_REF]], align
1090 // CHECK: [[REDUCE:%.+]] = icmp ne i32 [[SHOULD_REDUCE]], 0
1091 // CHECK: br i1 [[REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
1093 // CHECK: [[DO_REDUCE]]
1094 // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
1095 // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
1096 // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
1097 // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
1099 // Copy element from remote reduce list
1100 // CHECK: [[REDUCE_ELSE]]
1101 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
1102 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
1103 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
1104 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
1105 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i32*
1106 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
1107 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i32, i32* [[REMOTE_ELT]], align
1108 // CHECK: store i32 [[REMOTE_ELT_VAL]], i32* [[ELT]], align
1110 // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
1111 // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
1112 // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
1113 // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
1114 // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i16*
1115 // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
1116 // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i16, i16* [[REMOTE_ELT]], align
1117 // CHECK: store i16 [[REMOTE_ELT_VAL]], i16* [[ELT]], align
1118 // CHECK: br label {{%?}}[[REDUCE_CONT]]
1120 // CHECK: [[REDUCE_CONT]]