1 // Test target codegen - host bc file has to be created first.
2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -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 -fopenmp-cuda-mode -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 -fopenmp-cuda-mode -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 -fopenmp-cuda-mode -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 -fopenmp-cuda-mode -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
8 #ifndef HEADER
9 #define HEADER
10 
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]
13 
14 // Check that the execution mode of all 3 target regions is set to Spmd Mode.
15 // CHECK-DAG: {{@__omp_offloading_.+l27}}_exec_mode = weak constant i8 0
16 // CHECK-DAG: {{@__omp_offloading_.+l32}}_exec_mode = weak constant i8 0
17 // CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0
18 
19 template<typename tx>
20 tx ftemplate(int n) {
21   int a;
22   short b;
23   tx c;
24   float d;
25   double e;
26 
27   #pragma omp target parallel reduction(+: e)
28   {
29     e += 5;
30   }
31 
32   #pragma omp target parallel reduction(^: c) reduction(*: d)
33   {
34     c ^= 2;
35     d *= 33;
36   }
37 
38   #pragma omp target parallel reduction(|: a) reduction(max: b)
39   {
40     a |= 1;
41     b = 99 > b ? 99 : b;
42   }
43 
44   return a+b+c+d+e;
45 }
46 
47 int bar(int n){
48   int a = 0;
49 
50   a += ftemplate<char>(n);
51 
52   return a;
53 }
54 
55   // CHECK: define {{.*}}void {{@__omp_offloading_.+template.+l27}}(
56   //
57   // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 0)
58   // CHECK: call void @__kmpc_data_sharing_init_stack_spmd
59   // CHECK: br label {{%?}}[[EXECUTE:.+]]
60   //
61   // CHECK: [[EXECUTE]]
62   // CHECK: {{call|invoke}} void [[PFN:@.+]](i32*
63   // CHECK: call void @__kmpc_spmd_kernel_deinit()
64   //
65   //
66   // define internal void [[PFN]](
67   // CHECK: store double {{[0\.e\+]+}}, double* [[E:%.+]], align
68   // CHECK: [[EV:%.+]] = load double, double* [[E]], align
69   // CHECK: [[ADD:%.+]] = fadd double [[EV]], 5
70   // CHECK: store double [[ADD]], double* [[E]], align
71   // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [1 x i8*]* [[RL:%.+]], i{{32|64}} 0, i{{32|64}} 0
72   // CHECK: [[E_CAST:%.+]] = bitcast double* [[E]] to i8*
73   // CHECK: store i8* [[E_CAST]], i8** [[PTR1]], align
74   // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
75   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait(i32 {{.+}}, i32 1, i{{32|64}} {{4|8}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[WARP_COPY_FN:@.+]])
76   // CHECK: switch i32 [[RET]], label {{%?}}[[DEFAULTLABEL:.+]] [
77   // CHECK: i32 1, label {{%?}}[[REDLABEL:.+]]
78 
79   // CHECK: [[REDLABEL]]
80   // CHECK: [[E_INV:%.+]] = load double, double* [[E_IN:%.+]], align
81   // CHECK: [[EV:%.+]] = load double, double* [[E]], align
82   // CHECK: [[ADD:%.+]] = fadd double [[E_INV]], [[EV]]
83   // CHECK: store double [[ADD]], double* [[E_IN]], align
84   // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
85   // CHECK: br label %[[DEFAULTLABEL]]
86   //
87   // CHECK: [[DEFAULTLABEL]]
88   // CHECK: ret
89 
90   //
91   // Reduction function
92   // CHECK: define internal void [[REDUCTION_FUNC:@.+]](i8*, i8*)
93   // CHECK: [[VAR_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i{{32|64}} 0, i{{32|64}} 0
94   // CHECK: [[VAR_RHS_VOID:%.+]] = load i8*, i8** [[VAR_RHS_REF]],
95   // CHECK: [[VAR_RHS:%.+]] = bitcast i8* [[VAR_RHS_VOID]] to double*
96   //
97   // CHECK: [[VAR_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i{{32|64}} 0, i{{32|64}} 0
98   // CHECK: [[VAR_LHS_VOID:%.+]] = load i8*, i8** [[VAR_LHS_REF]],
99   // CHECK: [[VAR_LHS:%.+]] = bitcast i8* [[VAR_LHS_VOID]] to double*
100   //
101   // CHECK: [[VAR_LHS_VAL:%.+]] = load double, double* [[VAR_LHS]],
102   // CHECK: [[VAR_RHS_VAL:%.+]] = load double, double* [[VAR_RHS]],
103   // CHECK: [[RES:%.+]] = fadd double [[VAR_LHS_VAL]], [[VAR_RHS_VAL]]
104   // CHECK: store double [[RES]], double* [[VAR_LHS]],
105   // CHECK: ret void
106 
107   //
108   // Shuffle and reduce function
109   // CHECK: define internal void [[SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
110   // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
111   // CHECK: [[REMOTE_ELT:%.+]] = alloca double
112   //
113   // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
114   // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
115   // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
116   //
117   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
118   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
119   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
120   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
121   //
122   // CHECK: [[ELT_CAST:%.+]] = bitcast double* [[ELT]] to i64*
123   // CHECK: [[REMOTE_ELT_CAST:%.+]] = bitcast double* [[REMOTE_ELT]] to i64*
124   // CHECK: [[ELT_VAL:%.+]] = load i64, i64* [[ELT_CAST]], align
125   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
126   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
127   // CHECK: [[REMOTE_ELT_VAL64:%.+]] = call i64 @__kmpc_shuffle_int64(i64 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
128   //
129   // CHECK: store i64 [[REMOTE_ELT_VAL64]], i64* [[REMOTE_ELT_CAST]], align
130   // CHECK: [[REMOTE_ELT_VOID:%.+]] = bitcast double* [[REMOTE_ELT]] to i8*
131   // CHECK: store i8* [[REMOTE_ELT_VOID]], i8** [[REMOTE_ELT_REF]], align
132   //
133   // Condition to reduce
134   // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
135   //
136   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
137   // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
138   // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
139   //
140   // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
141   // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
142   // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
143   // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
144   // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
145   // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
146   //
147   // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
148   // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
149   // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
150   //
151   // CHECK: [[DO_REDUCE]]
152   // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
153   // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
154   // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
155   // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
156   //
157   // CHECK: [[REDUCE_ELSE]]
158   // CHECK: br label {{%?}}[[REDUCE_CONT]]
159   //
160   // CHECK: [[REDUCE_CONT]]
161   // Now check if we should just copy over the remote reduction list
162   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
163   // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
164   // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
165   // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
166   //
167   // CHECK: [[DO_COPY]]
168   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
169   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
170   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
171   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
172   // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to double*
173   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
174   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load double, double* [[REMOTE_ELT]], align
175   // CHECK: store double [[REMOTE_ELT_VAL]], double* [[ELT]], align
176   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
177   //
178   // CHECK: [[COPY_ELSE]]
179   // CHECK: br label {{%?}}[[COPY_CONT]]
180   //
181   // CHECK: [[COPY_CONT]]
182   // CHECK: void
183 
184   //
185   // Inter warp copy function
186   // CHECK: define internal void [[WARP_COPY_FN]](i8*, i32)
187   // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
188   // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
189   // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
190   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
191   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
192   //
193   // [[DO_COPY]]
194   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
195   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
196   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
197   //
198   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
199   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to double addrspace([[SHARED_ADDRSPACE]])*
200   // CHECK: [[ELT_VAL:%.+]] = load double, double* [[ELT]], align
201   // CHECK: store double [[ELT_VAL]], double addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
202   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
203   //
204   // CHECK: [[COPY_ELSE]]
205   // CHECK: br label {{%?}}[[COPY_CONT]]
206   //
207   // Barrier after copy to shared memory storage medium.
208   // CHECK: [[COPY_CONT]]
209   // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
210   // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
211   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
212   //
213   // Read into warp 0.
214   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
215   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
216   //
217   // CHECK: [[DO_READ]]
218   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
219   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to double addrspace([[SHARED_ADDRSPACE]])*
220   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
221   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
222   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to double*
223   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load double, double addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
224   // CHECK: store double [[MEDIUM_ELT_VAL]], double* [[ELT]], align
225   // CHECK: br label {{%?}}[[READ_CONT:.+]]
226   //
227   // CHECK: [[READ_ELSE]]
228   // CHECK: br label {{%?}}[[READ_CONT]]
229   //
230   // CHECK: [[READ_CONT]]
231   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
232   // CHECK: ret
233 
234 
235 
236 
237 
238 
239 
240 
241 
242 
243   // CHECK: define {{.*}}void {{@__omp_offloading_.+template.+l32}}(
244   //
245   // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 0)
246   // CHECK: call void @__kmpc_data_sharing_init_stack_spmd
247   // CHECK: br label {{%?}}[[EXECUTE:.+]]
248   //
249   // CHECK: [[EXECUTE]]
250   // CHECK: {{call|invoke}} void [[PFN1:@.+]](i32*
251   // CHECK: call void @__kmpc_spmd_kernel_deinit()
252   //
253   //
254   // define internal void [[PFN1]](
255   // CHECK: store float {{1\.[0e\+]+}}, float* [[D:%.+]], align
256   // CHECK: [[C_VAL:%.+]] = load i8, i8* [[C:%.+]], align
257   // CHECK: [[CONV:%.+]] = sext i8 [[C_VAL]] to i32
258   // CHECK: [[XOR:%.+]] = xor i32 [[CONV]], 2
259   // CHECK: [[TRUNC:%.+]] = trunc i32 [[XOR]] to i8
260   // CHECK: store i8 [[TRUNC]], i8* [[C]], align
261   // CHECK: [[DV:%.+]] = load float, float* [[D]], align
262   // CHECK: [[MUL:%.+]] = fmul float [[DV]], {{[0-9e\.\+]+}}
263   // CHECK: store float [[MUL]], float* [[D]], align
264   // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [2 x i8*]* [[RL:%.+]], i{{32|64}} 0, i{{32|64}} 0
265   // CHECK: store i8* [[C]], i8** [[PTR1]], align
266   // CHECK: [[PTR2:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RL]], i{{32|64}} 0, i{{32|64}} 1
267   // CHECK: [[D_CAST:%.+]] = bitcast float* [[D]] to i8*
268   // CHECK: store i8* [[D_CAST]], i8** [[PTR2]], align
269   // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
270   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait(i32 {{.+}}, i32 2, i{{32|64}} {{8|16}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[WARP_COPY_FN:@.+]])
271   // CHECK: switch i32 [[RET]], label {{%?}}[[DEFAULTLABEL:.+]] [
272   // CHECK: i32 1, label {{%?}}[[REDLABEL:.+]]
273 
274   // CHECK: [[REDLABEL]]
275   // CHECK: [[C_INV8:%.+]] = load i8, i8* [[C_IN:%.+]], align
276   // CHECK: [[C_INV:%.+]] = sext i8 [[C_INV8]] to i32
277   // CHECK: [[CV8:%.+]] = load i8, i8* [[C]], align
278   // CHECK: [[CV:%.+]] = sext i8 [[CV8]] to i32
279   // CHECK: [[XOR:%.+]] = xor i32 [[C_INV]], [[CV]]
280   // CHECK: [[TRUNC:%.+]] = trunc i32 [[XOR]] to i8
281   // CHECK: store i8 [[TRUNC]], i8* [[C_IN]], align
282   // CHECK: [[D_INV:%.+]] = load float, float* [[D_IN:%.+]], align
283   // CHECK: [[DV:%.+]] = load float, float* [[D]], align
284   // CHECK: [[MUL:%.+]] = fmul float [[D_INV]], [[DV]]
285   // CHECK: store float [[MUL]], float* [[D_IN]], align
286   // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
287   // CHECK: br label %[[DEFAULTLABEL]]
288   //
289   // CHECK: [[DEFAULTLABEL]]
290   // CHECK: ret
291 
292   //
293   // Reduction function
294   // CHECK: define internal void [[REDUCTION_FUNC:@.+]](i8*, i8*)
295   // CHECK: [[VAR1_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i{{32|64}} 0, i{{32|64}} 0
296   // CHECK: [[VAR1_RHS:%.+]] = load i8*, i8** [[VAR1_RHS_REF]],
297   //
298   // CHECK: [[VAR1_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i{{32|64}} 0, i{{32|64}} 0
299   // CHECK: [[VAR1_LHS:%.+]] = load i8*, i8** [[VAR1_LHS_REF]],
300   //
301   // CHECK: [[VAR2_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS]], i{{32|64}} 0, i{{32|64}} 1
302   // CHECK: [[VAR2_RHS_VOID:%.+]] = load i8*, i8** [[VAR2_RHS_REF]],
303   // CHECK: [[VAR2_RHS:%.+]] = bitcast i8* [[VAR2_RHS_VOID]] to float*
304   //
305   // CHECK: [[VAR2_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS]], i{{32|64}} 0, i{{32|64}} 1
306   // CHECK: [[VAR2_LHS_VOID:%.+]] = load i8*, i8** [[VAR2_LHS_REF]],
307   // CHECK: [[VAR2_LHS:%.+]] = bitcast i8* [[VAR2_LHS_VOID]] to float*
308   //
309   // CHECK: [[VAR1_LHS_VAL8:%.+]] = load i8, i8* [[VAR1_LHS]],
310   // CHECK: [[VAR1_LHS_VAL:%.+]] = sext i8 [[VAR1_LHS_VAL8]] to i32
311   // CHECK: [[VAR1_RHS_VAL8:%.+]] = load i8, i8* [[VAR1_RHS]],
312   // CHECK: [[VAR1_RHS_VAL:%.+]] = sext i8 [[VAR1_RHS_VAL8]] to i32
313   // CHECK: [[XOR:%.+]] = xor i32 [[VAR1_LHS_VAL]], [[VAR1_RHS_VAL]]
314   // CHECK: [[RES:%.+]] = trunc i32 [[XOR]] to i8
315   // CHECK: store i8 [[RES]], i8* [[VAR1_LHS]],
316   //
317   // CHECK: [[VAR2_LHS_VAL:%.+]] = load float, float* [[VAR2_LHS]],
318   // CHECK: [[VAR2_RHS_VAL:%.+]] = load float, float* [[VAR2_RHS]],
319   // CHECK: [[RES:%.+]] = fmul float [[VAR2_LHS_VAL]], [[VAR2_RHS_VAL]]
320   // CHECK: store float [[RES]], float* [[VAR2_LHS]],
321   // CHECK: ret void
322 
323   //
324   // Shuffle and reduce function
325   // CHECK: define internal void [[SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
326   // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
327   // CHECK: [[REMOTE_ELT1:%.+]] = alloca i8
328   // CHECK: [[REMOTE_ELT2:%.+]] = alloca float
329   //
330   // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
331   // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
332   // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
333   //
334   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
335   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
336   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
337   // CHECK: [[ELT_VAL:%.+]] = load i8, i8* [[ELT_VOID]], align
338   //
339   // CHECK: [[ELT_CAST:%.+]] = sext i8 [[ELT_VAL]] to i32
340   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
341   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
342   // CHECK: [[REMOTE_ELT1_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_CAST]], i16 [[LANEOFFSET]], i16 [[WS]])
343   // CHECK: [[REMOTE_ELT1_VAL:%.+]] = trunc i32 [[REMOTE_ELT1_VAL32]] to i8
344   //
345   // CHECK: store i8 [[REMOTE_ELT1_VAL]], i8* [[REMOTE_ELT1]], align
346   // CHECK: store i8* [[REMOTE_ELT1]], i8** [[REMOTE_ELT_REF]], align
347   //
348   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
349   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
350   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
351   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
352   //
353   // CHECK: [[ELT_CAST:%.+]] = bitcast float* [[ELT]] to i32*
354   // CHECK: [[REMOTE_ELT2_CAST:%.+]] = bitcast float* [[REMOTE_ELT2]] to i32*
355   // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT_CAST]], align
356   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
357   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
358   // CHECK: [[REMOTE_ELT2_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
359   //
360   // CHECK: store i32 [[REMOTE_ELT2_VAL32]], i32* [[REMOTE_ELT2_CAST]], align
361   // CHECK: [[REMOTE_ELT2C:%.+]] = bitcast float* [[REMOTE_ELT2]] to i8*
362   // CHECK: store i8* [[REMOTE_ELT2C]], i8** [[REMOTE_ELT_REF]], align
363   //
364   // Condition to reduce
365   // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
366   //
367   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
368   // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
369   // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
370   //
371   // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
372   // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
373   // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
374   // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
375   // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
376   // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
377   //
378   // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
379   // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
380   // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
381   //
382   // CHECK: [[DO_REDUCE]]
383   // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
384   // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
385   // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
386   // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
387   //
388   // CHECK: [[REDUCE_ELSE]]
389   // CHECK: br label {{%?}}[[REDUCE_CONT]]
390   //
391   // CHECK: [[REDUCE_CONT]]
392   // Now check if we should just copy over the remote reduction list
393   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
394   // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
395   // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
396   // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
397   //
398   // CHECK: [[DO_COPY]]
399   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
400   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
401   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
402   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
403   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i8, i8* [[REMOTE_ELT_VOID]], align
404   // CHECK: store i8 [[REMOTE_ELT_VAL]], i8* [[ELT_VOID]], align
405   //
406   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
407   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
408   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
409   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
410   // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to float*
411   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
412   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load float, float* [[REMOTE_ELT]], align
413   // CHECK: store float [[REMOTE_ELT_VAL]], float* [[ELT]], align
414   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
415   //
416   // CHECK: [[COPY_ELSE]]
417   // CHECK: br label {{%?}}[[COPY_CONT]]
418   //
419   // CHECK: [[COPY_CONT]]
420   // CHECK: void
421 
422   //
423   // Inter warp copy function
424   // CHECK: define internal void [[WARP_COPY_FN]](i8*, i32)
425   // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
426   // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
427   // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
428   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
429   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
430   //
431   // [[DO_COPY]]
432   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
433   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
434   //
435   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
436   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i8 addrspace([[SHARED_ADDRSPACE]])*
437   // CHECK: [[ELT_VAL:%.+]] = load i8, i8* [[ELT_VOID]], align
438   // CHECK: store i8 [[ELT_VAL]], i8 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
439   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
440   //
441   // CHECK: [[COPY_ELSE]]
442   // CHECK: br label {{%?}}[[COPY_CONT]]
443   //
444   // Barrier after copy to shared memory storage medium.
445   // CHECK: [[COPY_CONT]]
446   // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
447   // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
448   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
449   //
450   // Read into warp 0.
451   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
452   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
453   //
454   // CHECK: [[DO_READ]]
455   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
456   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i8 addrspace([[SHARED_ADDRSPACE]])*
457   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
458   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
459   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load i8, i8 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
460   // CHECK: store i8 [[MEDIUM_ELT_VAL]], i8* [[ELT_VOID]], align
461   // CHECK: br label {{%?}}[[READ_CONT:.+]]
462   //
463   // CHECK: [[READ_ELSE]]
464   // CHECK: br label {{%?}}[[READ_CONT]]
465   //
466   // CHECK: [[READ_CONT]]
467   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
468   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
469   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
470   //
471   // [[DO_COPY]]
472   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
473   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
474   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
475   //
476   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
477   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to float addrspace([[SHARED_ADDRSPACE]])*
478   // CHECK: [[ELT_VAL:%.+]] = load float, float* [[ELT]], align
479   // CHECK: store float [[ELT_VAL]], float addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
480   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
481   //
482   // CHECK: [[COPY_ELSE]]
483   // CHECK: br label {{%?}}[[COPY_CONT]]
484   //
485   // Barrier after copy to shared memory storage medium.
486   // CHECK: [[COPY_CONT]]
487   // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
488   // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
489   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
490   //
491   // Read into warp 0.
492   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
493   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
494   //
495   // CHECK: [[DO_READ]]
496   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
497   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to float addrspace([[SHARED_ADDRSPACE]])*
498   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 1
499   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
500   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to float*
501   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load float, float addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
502   // CHECK: store float [[MEDIUM_ELT_VAL]], float* [[ELT]], align
503   // CHECK: br label {{%?}}[[READ_CONT:.+]]
504   //
505   // CHECK: [[READ_ELSE]]
506   // CHECK: br label {{%?}}[[READ_CONT]]
507   //
508   // CHECK: [[READ_CONT]]
509   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
510   // CHECK: ret
511 
512 
513 
514 
515 
516 
517 
518 
519 
520 
521   // CHECK: define {{.*}}void {{@__omp_offloading_.+template.+l38}}(
522   //
523   // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 0)
524   // CHECK: call void @__kmpc_data_sharing_init_stack_spmd
525   // CHECK: br label {{%?}}[[EXECUTE:.+]]
526   //
527   // CHECK: [[EXECUTE]]
528   // CHECK: {{call|invoke}} void [[PFN2:@.+]](i32*
529   // CHECK: call void @__kmpc_spmd_kernel_deinit()
530   //
531   //
532   // define internal void [[PFN2]](
533   // CHECK: store i32 0, i32* [[A:%.+]], align
534   // CHECK: store i16 -32768, i16* [[B:%.+]], align
535   // CHECK: [[A_VAL:%.+]] = load i32, i32* [[A:%.+]], align
536   // CHECK: [[OR:%.+]] = or i32 [[A_VAL]], 1
537   // CHECK: store i32 [[OR]], i32* [[A]], align
538   // CHECK: [[BV16:%.+]] = load i16, i16* [[B]], align
539   // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
540   // CHECK: [[CMP:%.+]] = icmp sgt i32 99, [[BV]]
541   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
542   //
543   // CHECK: [[DO_MAX]]
544   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
545   //
546   // CHECK: [[MAX_ELSE]]
547   // CHECK: [[BV:%.+]] = load i16, i16* [[B]], align
548   // CHECK: [[MAX:%.+]] = sext i16 [[BV]] to i32
549   // CHECK: br label {{%?}}[[MAX_CONT]]
550   //
551   // CHECK: [[MAX_CONT]]
552   // CHECK: [[B_LVALUE:%.+]] = phi i32 [ 99, %[[DO_MAX]] ], [ [[MAX]], %[[MAX_ELSE]] ]
553   // CHECK: [[TRUNC:%.+]] = trunc i32 [[B_LVALUE]] to i16
554   // CHECK: store i16 [[TRUNC]], i16* [[B]], align
555   // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [2 x i8*]* [[RL:%.+]], i{{32|64}} 0, i{{32|64}} 0
556   // CHECK: [[A_CAST:%.+]] = bitcast i32* [[A]] to i8*
557   // CHECK: store i8* [[A_CAST]], i8** [[PTR1]], align
558   // CHECK: [[PTR2:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RL]], i{{32|64}} 0, i{{32|64}} 1
559   // CHECK: [[B_CAST:%.+]] = bitcast i16* [[B]] to i8*
560   // CHECK: store i8* [[B_CAST]], i8** [[PTR2]], align
561   // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
562   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait(i32 {{.+}}, i32 2, i{{32|64}} {{8|16}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[WARP_COPY_FN:@.+]])
563   // CHECK: switch i32 [[RET]], label {{%?}}[[DEFAULTLABEL:.+]] [
564   // CHECK: i32 1, label {{%?}}[[REDLABEL:.+]]
565 
566   // CHECK: [[REDLABEL]]
567   // CHECK: [[A_INV:%.+]] = load i32, i32* [[A_IN:%.+]], align
568   // CHECK: [[AV:%.+]] = load i32, i32* [[A]], align
569   // CHECK: [[OR:%.+]] = or i32 [[A_INV]], [[AV]]
570   // CHECK: store i32 [[OR]], i32* [[A_IN]], align
571   // CHECK: [[B_INV16:%.+]] = load i16, i16* [[B_IN:%.+]], align
572   // CHECK: [[B_INV:%.+]] = sext i16 [[B_INV16]] to i32
573   // CHECK: [[BV16:%.+]] = load i16, i16* [[B]], align
574   // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
575   // CHECK: [[CMP:%.+]] = icmp sgt i32 [[B_INV]], [[BV]]
576   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
577   //
578   // CHECK: [[DO_MAX]]
579   // CHECK: [[MAX1:%.+]] = load i16, i16* [[B_IN]], align
580   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
581   //
582   // CHECK: [[MAX_ELSE]]
583   // CHECK: [[MAX2:%.+]] = load i16, i16* [[B]], align
584   // CHECK: br label {{%?}}[[MAX_CONT]]
585   //
586   // CHECK: [[MAX_CONT]]
587   // CHECK: [[B_MAX:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
588   // CHECK: store i16 [[B_MAX]], i16* [[B_IN]], align
589   // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
590   // CHECK: br label %[[DEFAULTLABEL]]
591   //
592   // CHECK: [[DEFAULTLABEL]]
593   // CHECK: ret
594 
595   //
596   // Reduction function
597   // CHECK: define internal void [[REDUCTION_FUNC:@.+]](i8*, i8*)
598   // CHECK: [[VAR1_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i{{32|64}} 0, i{{32|64}} 0
599   // CHECK: [[VAR1_RHS_VOID:%.+]] = load i8*, i8** [[VAR1_RHS_REF]],
600   // CHECK: [[VAR1_RHS:%.+]] = bitcast i8* [[VAR1_RHS_VOID]] to i32*
601   //
602   // CHECK: [[VAR1_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i{{32|64}} 0, i{{32|64}} 0
603   // CHECK: [[VAR1_LHS_VOID:%.+]] = load i8*, i8** [[VAR1_LHS_REF]],
604   // CHECK: [[VAR1_LHS:%.+]] = bitcast i8* [[VAR1_LHS_VOID]] to i32*
605   //
606   // CHECK: [[VAR2_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS]], i{{32|64}} 0, i{{32|64}} 1
607   // CHECK: [[VAR2_RHS_VOID:%.+]] = load i8*, i8** [[VAR2_RHS_REF]],
608   // CHECK: [[VAR2_RHS:%.+]] = bitcast i8* [[VAR2_RHS_VOID]] to i16*
609   //
610   // CHECK: [[VAR2_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS]], i{{32|64}} 0, i{{32|64}} 1
611   // CHECK: [[VAR2_LHS_VOID:%.+]] = load i8*, i8** [[VAR2_LHS_REF]],
612   // CHECK: [[VAR2_LHS:%.+]] = bitcast i8* [[VAR2_LHS_VOID]] to i16*
613   //
614   // CHECK: [[VAR1_LHS_VAL:%.+]] = load i32, i32* [[VAR1_LHS]],
615   // CHECK: [[VAR1_RHS_VAL:%.+]] = load i32, i32* [[VAR1_RHS]],
616   // CHECK: [[OR:%.+]] = or i32 [[VAR1_LHS_VAL]], [[VAR1_RHS_VAL]]
617   // CHECK: store i32 [[OR]], i32* [[VAR1_LHS]],
618   //
619   // CHECK: [[VAR2_LHS_VAL16:%.+]] = load i16, i16* [[VAR2_LHS]],
620   // CHECK: [[VAR2_LHS_VAL:%.+]] = sext i16 [[VAR2_LHS_VAL16]] to i32
621   // CHECK: [[VAR2_RHS_VAL16:%.+]] = load i16, i16* [[VAR2_RHS]],
622   // CHECK: [[VAR2_RHS_VAL:%.+]] = sext i16 [[VAR2_RHS_VAL16]] to i32
623   //
624   // CHECK: [[CMP:%.+]] = icmp sgt i32 [[VAR2_LHS_VAL]], [[VAR2_RHS_VAL]]
625   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
626   //
627   // CHECK: [[DO_MAX]]
628   // CHECK: [[MAX1:%.+]] = load i16, i16* [[VAR2_LHS]], align
629   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
630   //
631   // CHECK: [[MAX_ELSE]]
632   // CHECK: [[MAX2:%.+]] = load i16, i16* [[VAR2_RHS]], align
633   // CHECK: br label {{%?}}[[MAX_CONT]]
634   //
635   // CHECK: [[MAX_CONT]]
636   // CHECK: [[MAXV:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
637   // CHECK: store i16 [[MAXV]], i16* [[VAR2_LHS]],
638   // CHECK: ret void
639 
640   //
641   // Shuffle and reduce function
642   // CHECK: define internal void [[SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
643   // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
644   // CHECK: [[REMOTE_ELT1:%.+]] = alloca i32
645   // CHECK: [[REMOTE_ELT2:%.+]] = alloca i16
646   //
647   // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
648   // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
649   // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
650   //
651   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
652   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
653   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
654   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
655   // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
656   //
657   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
658   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
659   // CHECK: [[REMOTE_ELT1_VAL:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
660   //
661   // CHECK: store i32 [[REMOTE_ELT1_VAL]], i32* [[REMOTE_ELT1]], align
662   // CHECK: [[REMOTE_ELT1C:%.+]] = bitcast i32* [[REMOTE_ELT1]] to i8*
663   // CHECK: store i8* [[REMOTE_ELT1C]], i8** [[REMOTE_ELT_REF]], align
664   //
665   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
666   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
667   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
668   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
669   // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
670   //
671   // CHECK: [[ELT_CAST:%.+]] = sext i16 [[ELT_VAL]] to i32
672   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
673   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
674   // CHECK: [[REMOTE_ELT2_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_CAST]], i16 [[LANEOFFSET]], i16 [[WS]])
675   // CHECK: [[REMOTE_ELT2_VAL:%.+]] = trunc i32 [[REMOTE_ELT2_VAL32]] to i16
676   //
677   // CHECK: store i16 [[REMOTE_ELT2_VAL]], i16* [[REMOTE_ELT2]], align
678   // CHECK: [[REMOTE_ELT2C:%.+]] = bitcast i16* [[REMOTE_ELT2]] to i8*
679   // CHECK: store i8* [[REMOTE_ELT2C]], i8** [[REMOTE_ELT_REF]], align
680   //
681   // Condition to reduce
682   // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
683   //
684   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
685   // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
686   // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
687   //
688   // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
689   // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
690   // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
691   // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
692   // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
693   // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
694   //
695   // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
696   // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
697   // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
698   //
699   // CHECK: [[DO_REDUCE]]
700   // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
701   // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
702   // CHECK: call void [[REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
703   // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
704   //
705   // CHECK: [[REDUCE_ELSE]]
706   // CHECK: br label {{%?}}[[REDUCE_CONT]]
707   //
708   // CHECK: [[REDUCE_CONT]]
709   // Now check if we should just copy over the remote reduction list
710   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
711   // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
712   // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
713   // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
714   //
715   // CHECK: [[DO_COPY]]
716   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
717   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
718   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
719   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
720   // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i32*
721   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
722   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i32, i32* [[REMOTE_ELT]], align
723   // CHECK: store i32 [[REMOTE_ELT_VAL]], i32* [[ELT]], align
724   //
725   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
726   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
727   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
728   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
729   // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i16*
730   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
731   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i16, i16* [[REMOTE_ELT]], align
732   // CHECK: store i16 [[REMOTE_ELT_VAL]], i16* [[ELT]], align
733   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
734   //
735   // CHECK: [[COPY_ELSE]]
736   // CHECK: br label {{%?}}[[COPY_CONT]]
737   //
738   // CHECK: [[COPY_CONT]]
739   // CHECK: void
740 
741   //
742   // Inter warp copy function
743   // CHECK: define internal void [[WARP_COPY_FN]](i8*, i32)
744   // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
745   // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
746   // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
747   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
748   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
749   //
750   // [[DO_COPY]]
751   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 0
752   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
753   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
754   //
755   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
756   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i32 addrspace([[SHARED_ADDRSPACE]])*
757   // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
758   // CHECK: store i32 [[ELT_VAL]], i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
759   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
760   //
761   // CHECK: [[COPY_ELSE]]
762   // CHECK: br label {{%?}}[[COPY_CONT]]
763   //
764   // Barrier after copy to shared memory storage medium.
765   // CHECK: [[COPY_CONT]]
766   // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
767   // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
768   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
769   //
770   // Read into warp 0.
771   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
772   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
773   //
774   // CHECK: [[DO_READ]]
775   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
776   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i32 addrspace([[SHARED_ADDRSPACE]])*
777   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 0
778   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
779   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
780   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load i32, i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
781   // CHECK: store i32 [[MEDIUM_ELT_VAL]], i32* [[ELT]], align
782   // CHECK: br label {{%?}}[[READ_CONT:.+]]
783   //
784   // CHECK: [[READ_ELSE]]
785   // CHECK: br label {{%?}}[[READ_CONT]]
786   //
787   // CHECK: [[READ_CONT]]
788   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
789   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
790   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
791   //
792   // [[DO_COPY]]
793   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i{{32|64}} 0, i{{32|64}} 1
794   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
795   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
796   //
797   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
798   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i16 addrspace([[SHARED_ADDRSPACE]])*
799   // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
800   // CHECK: store i16 [[ELT_VAL]], i16 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
801   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
802   //
803   // CHECK: [[COPY_ELSE]]
804   // CHECK: br label {{%?}}[[COPY_CONT]]
805   //
806   // Barrier after copy to shared memory storage medium.
807   // CHECK: [[COPY_CONT]]
808   // CHECK: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
809   // CHECK: [[ACTIVE_THREADS:%.+]] = mul nsw i32 [[ACTIVE_WARPS:%.+]], [[WS]]
810   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
811   //
812   // Read into warp 0.
813   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
814   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
815   //
816   // CHECK: [[DO_READ]]
817   // CHECK: [[MEDIUM_ELT64:%.+]] = getelementptr inbounds [32 x i64], [32 x i64] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
818   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i64 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT64]] to i16 addrspace([[SHARED_ADDRSPACE]])*
819   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i{{32|64}} 0, i{{32|64}} 1
820   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
821   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
822   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load i16, i16 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
823   // CHECK: store i16 [[MEDIUM_ELT_VAL]], i16* [[ELT]], align
824   // CHECK: br label {{%?}}[[READ_CONT:.+]]
825   //
826   // CHECK: [[READ_ELSE]]
827   // CHECK: br label {{%?}}[[READ_CONT]]
828   //
829   // CHECK: [[READ_CONT]]
830   // CHECK: call void @llvm.nvvm.barrier(i32 1, i32 [[ACTIVE_THREADS]])
831   // CHECK: ret
832 
833 #endif
834