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
8 #ifndef HEADER
9 #define HEADER
10 
11 // CHECK: [[MAP_TY:%.+]] = type { [128 x i8] }
12 
13 // CHECK-DAG: [[KERNEL_PTR:@.+]] = internal addrspace(3) global i8* null
14 // CHECK-DAG: [[KERNEL_SHARED1:@.+]] = internal unnamed_addr constant i16 1
15 // CHECK-DAG: [[KERNEL_SHARED2:@.+]] = internal unnamed_addr constant i16 1
16 // CHECK-DAG: [[KERNEL_SHARED3:@.+]] = internal unnamed_addr constant i16 1
17 // CHECK-DAG: [[KERNEL_SIZE1:@.+]] = internal unnamed_addr constant i{{64|32}} {{16|8}}
18 // CHECK-DAG: [[KERNEL_SIZE2:@.+]] = internal unnamed_addr constant i{{64|32}} 16
19 // CHECK-DAG: [[KERNEL_SIZE3:@.+]] = internal unnamed_addr constant i{{64|32}} 8
20 
21 // Check for the data transfer medium in shared memory to transfer the reduction list to the first warp.
22 // CHECK-DAG: [[TRANSFER_STORAGE:@.+]] = common addrspace([[SHARED_ADDRSPACE:[0-9]+]]) global [32 x i32]
23 
24 // Check that the execution mode of 2 target regions is set to Non-SPMD and the 3rd is in SPMD.
25 // CHECK-DAG: {{@__omp_offloading_.+l37}}_exec_mode = weak constant i8 1
26 // CHECK-DAG: {{@__omp_offloading_.+l43}}_exec_mode = weak constant i8 1
27 // CHECK-DAG: {{@__omp_offloading_.+l50}}_exec_mode = weak constant i8 0
28 
29 template<typename tx>
30 tx ftemplate(int n) {
31   int a;
32   short b;
33   tx c;
34   float d;
35   double e;
36 
37   #pragma omp target
38   #pragma omp teams reduction(+: e)
39   {
40     e += 5;
41   }
42 
43   #pragma omp target
44   #pragma omp teams reduction(^: c) reduction(*: d)
45   {
46     c ^= 2;
47     d *= 33;
48   }
49 
50   #pragma omp target
51   #pragma omp teams reduction(|: a) reduction(max: b)
52   #pragma omp parallel reduction(|: a) reduction(max: b)
53   {
54     a |= 1;
55     b = 99 > b ? 99 : b;
56   }
57 
58   return a+b+c+d+e;
59 }
60 
61 int bar(int n){
62   int a = 0;
63 
64   a += ftemplate<char>(n);
65 
66   return a;
67 }
68 
69   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l37}}_worker()
70 
71   // CHECK: define {{.*}}void [[T1:@__omp_offloading_.+template.+l37]](
72   //
73   // CHECK: {{call|invoke}} void [[T1]]_worker()
74   //
75   // CHECK: call void @__kmpc_kernel_init(
76   //
77   // CHECK: store double {{[0\.e\+]+}}, double* [[E:%.+]], align
78   // CHECK: [[EV:%.+]] = load double, double* [[E]], align
79   // CHECK: [[ADD:%.+]] = fadd double [[EV]], 5
80   // CHECK: store double [[ADD]], double* [[E]], align
81   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_teams_reduce_nowait_simple(%struct.ident_t* [[LOC:@.+]], i32 [[GTID:%.+]], [8 x i32]* [[LOCK:@.+]])
82   // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
83   // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
84   //
85   // CHECK: [[IFLABEL]]
86   // CHECK: [[E_INV:%.+]] = load double, double* [[E_IN:%.+]], align
87   // CHECK: [[EV:%.+]] = load double, double* [[E]], align
88   // CHECK: [[ADD:%.+]] = fadd double [[E_INV]], [[EV]]
89   // CHECK: store double [[ADD]], double* [[E_IN]], align
90   // CHECK: call void @__kmpc_nvptx_teams_end_reduce_nowait_simple(%struct.ident_t* [[LOC]], i32 [[GTID]], [8 x i32]* [[LOCK]])
91   // CHECK: br label %[[EXIT]]
92   //
93   // CHECK: [[EXIT]]
94   // CHECK: call void @__kmpc_kernel_deinit(
95 
96   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l43}}_worker()
97 
98   // CHECK: define {{.*}}void [[T2:@__omp_offloading_.+template.+l43]](
99   //
100   // CHECK: {{call|invoke}} void [[T2]]_worker()
101 
102   //
103   // CHECK: call void @__kmpc_kernel_init(
104   //
105   // CHECK: store float {{1\.[0e\+]+}}, float* [[D:%.+]], align
106   // CHECK: [[C_VAL:%.+]] = load i8, i8* [[C:%.+]], align
107   // CHECK: [[CONV:%.+]] = sext i8 [[C_VAL]] to i32
108   // CHECK: [[XOR:%.+]] = xor i32 [[CONV]], 2
109   // CHECK: [[TRUNC:%.+]] = trunc i32 [[XOR]] to i8
110   // CHECK: store i8 [[TRUNC]], i8* [[C]], align
111   // CHECK: [[DV:%.+]] = load float, float* [[D]], align
112   // CHECK: [[MUL:%.+]] = fmul float [[DV]], {{[0-9e\.\+]+}}
113   // CHECK: store float [[MUL]], float* [[D]], align
114   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_teams_reduce_nowait_simple(%struct.ident_t* [[LOC:@.+]], i32 [[GTID:%.+]], [8 x i32]* [[LOCK:@.+]])
115   // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
116   // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
117   //
118   // CHECK: [[IFLABEL]]
119   // CHECK: [[C_INV8:%.+]] = load i8, i8* [[C_IN:%.+]], align
120   // CHECK: [[C_INV:%.+]] = sext i8 [[C_INV8]] to i32
121   // CHECK: [[CV8:%.+]] = load i8, i8* [[C]], align
122   // CHECK: [[CV:%.+]] = sext i8 [[CV8]] to i32
123   // CHECK: [[XOR:%.+]] = xor i32 [[C_INV]], [[CV]]
124   // CHECK: [[TRUNC:%.+]] = trunc i32 [[XOR]] to i8
125   // CHECK: store i8 [[TRUNC]], i8* [[C_IN]], align
126   // CHECK: [[D_INV:%.+]] = load float, float* [[D_IN:%.+]], align
127   // CHECK: [[DV:%.+]] = load float, float* [[D]], align
128   // CHECK: [[MUL:%.+]] = fmul float [[D_INV]], [[DV]]
129   // CHECK: store float [[MUL]], float* [[D_IN]], align
130   // CHECK: call void @__kmpc_nvptx_teams_end_reduce_nowait_simple(%struct.ident_t* [[LOC]], i32 [[GTID]], [8 x i32]* [[LOCK]])
131   // CHECK: br label %[[EXIT]]
132   //
133   // CHECK: [[EXIT]]
134   // CHECK: call void @__kmpc_kernel_deinit(
135 
136   // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l50}}(
137   //
138   // CHECK: call void @__kmpc_spmd_kernel_init(
139   // CHECK: call void @__kmpc_data_sharing_init_stack_spmd()
140   // CHECK: call void @__kmpc_get_team_static_memory(i16 1, i8* addrspacecast (i8 addrspace(3)* getelementptr inbounds ([[MEM_TY:%.+]], %{{.+}} addrspace(3)* [[KERNEL_RD:@.+]], i32 0, i32 0, i32 0) to i8*), i{{64|32}} {{8|16}}, i16 1, i8** addrspacecast (i8* addrspace(3)* [[KERNEL_PTR:@.+]] to i8**))
141   // CHECK: [[PTR:%.+]] = load i8*, i8* addrspace(3)* [[KERNEL_PTR]],
142   // CHECK: [[GLOBAL_REC:%.+]] = bitcast i8* [[PTR]] to [[GLOB_REC_TY:%.+]]*
143   // CHECK-DAG: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOB_REC_TY]], [[GLOB_REC_TY]]* [[GLOBAL_REC]], i32 0, i32 0
144   // CHECK-DAG: [[B_ADDR:%.+]] = getelementptr inbounds [[GLOB_REC_TY]], [[GLOB_REC_TY]]* [[GLOBAL_REC]], i32 0, i32 1
145   // CHECK: store i32 0, i32* [[A_ADDR]],
146   // CHECK: store i16 -32768, i16* [[B_ADDR]],
147   // CHECK: call void [[OUTLINED:@.+]](i32* {{.+}}, i32* {{.+}}, i32* [[A_ADDR]], i16* [[B_ADDR]])
148   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_teams_reduce_nowait_simple(%struct.ident_t* [[LOC:@.+]], i32 [[GTID:%.+]], [8 x i32]* [[LOCK:@.+]])
149   // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
150   // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
151   //
152   // CHECK: [[IFLABEL]]
153   // CHECK: [[A_INV:%.+]] = load i32, i32* [[A_IN:%.+]], align
154   // CHECK: [[AV:%.+]] = load i32, i32* [[A_ADDR]], align
155   // CHECK: [[OR:%.+]] = or i32 [[A_INV]], [[AV]]
156   // CHECK: store i32 [[OR]], i32* [[A_IN]], align
157   // CHECK: [[B_INV16:%.+]] = load i16, i16* [[B_IN:%.+]], align
158   // CHECK: [[B_INV:%.+]] = sext i16 [[B_INV16]] to i32
159   // CHECK: [[BV16:%.+]] = load i16, i16* [[B_ADDR]], align
160   // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
161   // CHECK: [[CMP:%.+]] = icmp sgt i32 [[B_INV]], [[BV]]
162   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
163   //
164   // CHECK: [[DO_MAX]]
165   // CHECK: [[MAX1:%.+]] = load i16, i16* [[B_IN]], align
166   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
167   //
168   // CHECK: [[MAX_ELSE]]
169   // CHECK: [[MAX2:%.+]] = load i16, i16* [[B_ADDR]], align
170   // CHECK: br label {{%?}}[[MAX_CONT]]
171   //
172   // CHECK: [[MAX_CONT]]
173   // CHECK: [[B_MAX:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
174   // CHECK: store i16 [[B_MAX]], i16* [[B_IN]], align
175   // CHECK: call void @__kmpc_nvptx_teams_end_reduce_nowait_simple(%struct.ident_t* [[LOC]], i32 [[GTID]], [8 x i32]* [[LOCK]])
176   // CHECK: br label %[[EXIT]]
177   //
178   // CHECK: [[EXIT]]
179   // CHECK: call void @__kmpc_restore_team_static_memory(i16 1, i16 1)
180   // CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
181 
182   // CHECK: define internal void [[OUTLINED]](i32* noalias %{{.+}}, i32* noalias %{{.+}}, i32* dereferenceable{{.+}}, i16* dereferenceable{{.+}})
183   //
184   // CHECK: store i32 0, i32* [[A:%.+]], align
185   // CHECK: store i16 -32768, i16* [[B:%.+]], align
186   // CHECK: [[A_VAL:%.+]] = load i32, i32* [[A:%.+]], align
187   // CHECK: [[OR:%.+]] = or i32 [[A_VAL]], 1
188   // CHECK: store i32 [[OR]], i32* [[A]], align
189   // CHECK: [[BV16:%.+]] = load i16, i16* [[B]], align
190   // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
191   // CHECK: [[CMP:%.+]] = icmp sgt i32 99, [[BV]]
192   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
193   //
194   // CHECK: [[DO_MAX]]
195   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
196   //
197   // CHECK: [[MAX_ELSE]]
198   // CHECK: [[BV:%.+]] = load i16, i16* [[B]], align
199   // CHECK: [[MAX:%.+]] = sext i16 [[BV]] to i32
200   // CHECK: br label {{%?}}[[MAX_CONT]]
201   //
202   // CHECK: [[MAX_CONT]]
203   // CHECK: [[B_LVALUE:%.+]] = phi i32 [ 99, %[[DO_MAX]] ], [ [[MAX]], %[[MAX_ELSE]] ]
204   // CHECK: [[TRUNC:%.+]] = trunc i32 [[B_LVALUE]] to i16
205   // CHECK: store i16 [[TRUNC]], i16* [[B]], align
206   // CHECK: [[PTR1:%.+]] = getelementptr inbounds [[RLT:.+]], [2 x i8*]* [[RL:%.+]], i{{.+}} 0, i[[SZ:.+]] 0
207   // CHECK: [[A_CAST:%.+]] = bitcast i32* [[A]] to i8*
208   // CHECK: store i8* [[A_CAST]], i8** [[PTR1]], align
209   // CHECK: [[PTR2:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RL]], i[[SZ]] 0, i[[SZ]] 1
210   // CHECK: [[B_CAST:%.+]] = bitcast i16* [[B]] to i8*
211   // CHECK: store i8* [[B_CAST]], i8** [[PTR2]], align
212   // CHECK: [[ARG_RL:%.+]] = bitcast [[RLT]]* [[RL]] to i8*
213   // CHECK: [[RET:%.+]] = call i32 @__kmpc_nvptx_parallel_reduce_nowait_v2(%struct.ident_t* [[LOC]], i32 {{.+}}, i32 2, i[[SZ]] {{8|16}}, i8* [[ARG_RL]], void (i8*, i16, i16, i16)* [[PAR_SHUFFLE_REDUCE_FN:@.+]], void (i8*, i32)* [[PAR_WARP_COPY_FN:@.+]])
214   // CHECK: [[COND:%.+]] = icmp eq i32 [[RET]], 1
215   // CHECK: br i1 [[COND]], label {{%?}}[[IFLABEL:.+]], label {{%?}}[[EXIT:.+]]
216   //
217   // CHECK: [[IFLABEL]]
218   // CHECK: [[A_INV:%.+]] = load i32, i32* [[A_IN:%.+]], align
219   // CHECK: [[AV:%.+]] = load i32, i32* [[A]], align
220   // CHECK: [[OR:%.+]] = or i32 [[A_INV]], [[AV]]
221   // CHECK: store i32 [[OR]], i32* [[A_IN]], align
222   // CHECK: [[B_INV16:%.+]] = load i16, i16* [[B_IN:%.+]], align
223   // CHECK: [[B_INV:%.+]] = sext i16 [[B_INV16]] to i32
224   // CHECK: [[BV16:%.+]] = load i16, i16* [[B]], align
225   // CHECK: [[BV:%.+]] = sext i16 [[BV16]] to i32
226   // CHECK: [[CMP:%.+]] = icmp sgt i32 [[B_INV]], [[BV]]
227   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
228   //
229   // CHECK: [[DO_MAX]]
230   // CHECK: [[MAX1:%.+]] = load i16, i16* [[B_IN]], align
231   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
232   //
233   // CHECK: [[MAX_ELSE]]
234   // CHECK: [[MAX2:%.+]] = load i16, i16* [[B]], align
235   // CHECK: br label {{%?}}[[MAX_CONT]]
236   //
237   // CHECK: [[MAX_CONT]]
238   // CHECK: [[B_MAX:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
239   // CHECK: store i16 [[B_MAX]], i16* [[B_IN]], align
240   // CHECK: call void @__kmpc_nvptx_end_reduce_nowait(
241   // CHECK: br label %[[EXIT]]
242   //
243   // CHECK: [[EXIT]]
244   // CHECK: ret void
245 
246   //
247   // Reduction function
248   // CHECK: define internal void [[PAR_REDUCTION_FUNC:@.+]](i8*, i8*)
249   // CHECK: [[VAR1_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
250   // CHECK: [[VAR1_RHS_VOID:%.+]] = load i8*, i8** [[VAR1_RHS_REF]],
251   // CHECK: [[VAR1_RHS:%.+]] = bitcast i8* [[VAR1_RHS_VOID]] to i32*
252   //
253   // CHECK: [[VAR1_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS:%.+]], i[[SZ]] 0, i[[SZ]] 0
254   // CHECK: [[VAR1_LHS_VOID:%.+]] = load i8*, i8** [[VAR1_LHS_REF]],
255   // CHECK: [[VAR1_LHS:%.+]] = bitcast i8* [[VAR1_LHS_VOID]] to i32*
256   //
257   // CHECK: [[VAR2_RHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_RHS]], i[[SZ]] 0, i[[SZ]] 1
258   // CHECK: [[VAR2_RHS_VOID:%.+]] = load i8*, i8** [[VAR2_RHS_REF]],
259   // CHECK: [[VAR2_RHS:%.+]] = bitcast i8* [[VAR2_RHS_VOID]] to i16*
260   //
261   // CHECK: [[VAR2_LHS_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST_LHS]], i[[SZ]] 0, i[[SZ]] 1
262   // CHECK: [[VAR2_LHS_VOID:%.+]] = load i8*, i8** [[VAR2_LHS_REF]],
263   // CHECK: [[VAR2_LHS:%.+]] = bitcast i8* [[VAR2_LHS_VOID]] to i16*
264   //
265   // CHECK: [[VAR1_LHS_VAL:%.+]] = load i32, i32* [[VAR1_LHS]],
266   // CHECK: [[VAR1_RHS_VAL:%.+]] = load i32, i32* [[VAR1_RHS]],
267   // CHECK: [[OR:%.+]] = or i32 [[VAR1_LHS_VAL]], [[VAR1_RHS_VAL]]
268   // CHECK: store i32 [[OR]], i32* [[VAR1_LHS]],
269   //
270   // CHECK: [[VAR2_LHS_VAL16:%.+]] = load i16, i16* [[VAR2_LHS]],
271   // CHECK: [[VAR2_LHS_VAL:%.+]] = sext i16 [[VAR2_LHS_VAL16]] to i32
272   // CHECK: [[VAR2_RHS_VAL16:%.+]] = load i16, i16* [[VAR2_RHS]],
273   // CHECK: [[VAR2_RHS_VAL:%.+]] = sext i16 [[VAR2_RHS_VAL16]] to i32
274   //
275   // CHECK: [[CMP:%.+]] = icmp sgt i32 [[VAR2_LHS_VAL]], [[VAR2_RHS_VAL]]
276   // CHECK: br i1 [[CMP]], label {{%?}}[[DO_MAX:.+]], label {{%?}}[[MAX_ELSE:.+]]
277   //
278   // CHECK: [[DO_MAX]]
279   // CHECK: [[MAX1:%.+]] = load i16, i16* [[VAR2_LHS]], align
280   // CHECK: br label {{%?}}[[MAX_CONT:.+]]
281   //
282   // CHECK: [[MAX_ELSE]]
283   // CHECK: [[MAX2:%.+]] = load i16, i16* [[VAR2_RHS]], align
284   // CHECK: br label {{%?}}[[MAX_CONT]]
285   //
286   // CHECK: [[MAX_CONT]]
287   // CHECK: [[MAXV:%.+]] = phi i16 [ [[MAX1]], %[[DO_MAX]] ], [ [[MAX2]], %[[MAX_ELSE]] ]
288   // CHECK: store i16 [[MAXV]], i16* [[VAR2_LHS]],
289   // CHECK: ret void
290   //
291   // Shuffle and reduce function
292   // CHECK: define internal void [[PAR_SHUFFLE_REDUCE_FN]](i8*, i16 {{.*}}, i16 {{.*}}, i16 {{.*}})
293   // CHECK: [[REMOTE_RED_LIST:%.+]] = alloca [[RLT]], align
294   // CHECK: [[REMOTE_ELT1:%.+]] = alloca i32
295   // CHECK: [[REMOTE_ELT2:%.+]] = alloca i16
296   //
297   // CHECK: [[LANEID:%.+]] = load i16, i16* {{.+}}, align
298   // CHECK: [[LANEOFFSET:%.+]] = load i16, i16* {{.+}}, align
299   // CHECK: [[ALGVER:%.+]] = load i16, i16* {{.+}}, align
300   //
301   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
302   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
303   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
304   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
305   // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
306   //
307   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
308   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
309   // CHECK: [[REMOTE_ELT1_VAL:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_VAL]], i16 [[LANEOFFSET]], i16 [[WS]])
310   //
311   // CHECK: store i32 [[REMOTE_ELT1_VAL]], i32* [[REMOTE_ELT1]], align
312   // CHECK: [[REMOTE_ELT1C:%.+]] = bitcast i32* [[REMOTE_ELT1]] to i8*
313   // CHECK: store i8* [[REMOTE_ELT1C]], i8** [[REMOTE_ELT_REF]], align
314   //
315   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
316   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
317   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
318   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
319   // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
320   //
321   // CHECK: [[ELT_CAST:%.+]] = sext i16 [[ELT_VAL]] to i32
322   // CHECK: [[WS32:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
323   // CHECK: [[WS:%.+]] = trunc i32 [[WS32]] to i16
324   // CHECK: [[REMOTE_ELT2_VAL32:%.+]] = call i32 @__kmpc_shuffle_int32(i32 [[ELT_CAST]], i16 [[LANEOFFSET]], i16 [[WS]])
325   // CHECK: [[REMOTE_ELT2_VAL:%.+]] = trunc i32 [[REMOTE_ELT2_VAL32]] to i16
326   //
327   // CHECK: store i16 [[REMOTE_ELT2_VAL]], i16* [[REMOTE_ELT2]], align
328   // CHECK: [[REMOTE_ELT2C:%.+]] = bitcast i16* [[REMOTE_ELT2]] to i8*
329   // CHECK: store i8* [[REMOTE_ELT2C]], i8** [[REMOTE_ELT_REF]], align
330   //
331   // Condition to reduce
332   // CHECK: [[CONDALG0:%.+]] = icmp eq i16 [[ALGVER]], 0
333   //
334   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
335   // CHECK: [[COND2:%.+]] = icmp ult i16 [[LANEID]], [[LANEOFFSET]]
336   // CHECK: [[CONDALG1:%.+]] = and i1 [[COND1]], [[COND2]]
337   //
338   // CHECK: [[COND3:%.+]] = icmp eq i16 [[ALGVER]], 2
339   // CHECK: [[COND4:%.+]] = and i16 [[LANEID]], 1
340   // CHECK: [[COND5:%.+]] = icmp eq i16 [[COND4]], 0
341   // CHECK: [[COND6:%.+]] = and i1 [[COND3]], [[COND5]]
342   // CHECK: [[COND7:%.+]] = icmp sgt i16 [[LANEOFFSET]], 0
343   // CHECK: [[CONDALG2:%.+]] = and i1 [[COND6]], [[COND7]]
344   //
345   // CHECK: [[COND8:%.+]] = or i1 [[CONDALG0]], [[CONDALG1]]
346   // CHECK: [[SHOULD_REDUCE:%.+]] = or i1 [[COND8]], [[CONDALG2]]
347   // CHECK: br i1 [[SHOULD_REDUCE]], label {{%?}}[[DO_REDUCE:.+]], label {{%?}}[[REDUCE_ELSE:.+]]
348   //
349   // CHECK: [[DO_REDUCE]]
350   // CHECK: [[RED_LIST1_VOID:%.+]] = bitcast [[RLT]]* [[RED_LIST]] to i8*
351   // CHECK: [[RED_LIST2_VOID:%.+]] = bitcast [[RLT]]* [[REMOTE_RED_LIST]] to i8*
352   // CHECK: call void [[PAR_REDUCTION_FUNC]](i8* [[RED_LIST1_VOID]], i8* [[RED_LIST2_VOID]])
353   // CHECK: br label {{%?}}[[REDUCE_CONT:.+]]
354   //
355   // CHECK: [[REDUCE_ELSE]]
356   // CHECK: br label {{%?}}[[REDUCE_CONT]]
357   //
358   // CHECK: [[REDUCE_CONT]]
359   // Now check if we should just copy over the remote reduction list
360   // CHECK: [[COND1:%.+]] = icmp eq i16 [[ALGVER]], 1
361   // CHECK: [[COND2:%.+]] = icmp uge i16 [[LANEID]], [[LANEOFFSET]]
362   // CHECK: [[SHOULD_COPY:%.+]] = and i1 [[COND1]], [[COND2]]
363   // CHECK: br i1 [[SHOULD_COPY]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
364   //
365   // CHECK: [[DO_COPY]]
366   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
367   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
368   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
369   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
370   // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i32*
371   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
372   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i32, i32* [[REMOTE_ELT]], align
373   // CHECK: store i32 [[REMOTE_ELT_VAL]], i32* [[ELT]], align
374   //
375   // CHECK: [[REMOTE_ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[REMOTE_RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
376   // CHECK: [[REMOTE_ELT_VOID:%.+]] = load i8*, i8** [[REMOTE_ELT_REF]],
377   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
378   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
379   // CHECK: [[REMOTE_ELT:%.+]] = bitcast i8* [[REMOTE_ELT_VOID]] to i16*
380   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
381   // CHECK: [[REMOTE_ELT_VAL:%.+]] = load i16, i16* [[REMOTE_ELT]], align
382   // CHECK: store i16 [[REMOTE_ELT_VAL]], i16* [[ELT]], align
383   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
384   //
385   // CHECK: [[COPY_ELSE]]
386   // CHECK: br label {{%?}}[[COPY_CONT]]
387   //
388   // CHECK: [[COPY_CONT]]
389   // CHECK: void
390 
391   //
392   // Inter warp copy function
393   // CHECK: define internal void [[PAR_WARP_COPY_FN]](i8*, i32)
394   // CHECK-DAG: [[LANEID:%.+]] = and i32 {{.+}}, 31
395   // CHECK-DAG: [[WARPID:%.+]] = ashr i32 {{.+}}, 5
396   // CHECK-DAG: [[RED_LIST:%.+]] = bitcast i8* {{.+}} to [[RLT]]*
397   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
398   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
399   //
400   // [[DO_COPY]]
401   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 0
402   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
403   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
404   //
405   // CHECK: [[MEDIUM_ELT:%.+]] = getelementptr inbounds [32 x i32], [32 x i32] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
406   // CHECK: [[ELT_VAL:%.+]] = load i32, i32* [[ELT]], align
407   // CHECK: store volatile i32 [[ELT_VAL]], i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
408   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
409   //
410   // CHECK: [[COPY_ELSE]]
411   // CHECK: br label {{%?}}[[COPY_CONT]]
412   //
413   // Barrier after copy to shared memory storage medium.
414   // CHECK: [[COPY_CONT]]
415   // CHECK: call void @__kmpc_barrier(%struct.ident_t* @
416   // CHECK: [[ACTIVE_WARPS:%.+]] = load i32, i32*
417   //
418   // Read into warp 0.
419   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
420   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
421   //
422   // CHECK: [[DO_READ]]
423   // CHECK: [[MEDIUM_ELT:%.+]] = getelementptr inbounds [32 x i32], [32 x i32] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
424   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 0
425   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
426   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i32*
427   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load volatile i32, i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
428   // CHECK: store i32 [[MEDIUM_ELT_VAL]], i32* [[ELT]], align
429   // CHECK: br label {{%?}}[[READ_CONT:.+]]
430   //
431   // CHECK: [[READ_ELSE]]
432   // CHECK: br label {{%?}}[[READ_CONT]]
433   //
434   // CHECK: [[READ_CONT]]
435   // CHECK: call void @__kmpc_barrier(%struct.ident_t* @
436   // CHECK: [[IS_WARP_MASTER:%.+]] = icmp eq i32 [[LANEID]], 0
437   // CHECK: br i1 [[IS_WARP_MASTER]], label {{%?}}[[DO_COPY:.+]], label {{%?}}[[COPY_ELSE:.+]]
438   //
439   // [[DO_COPY]]
440   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST]], i[[SZ]] 0, i[[SZ]] 1
441   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
442   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
443   //
444   // CHECK: [[MEDIUM_ELT32:%.+]] = getelementptr inbounds [32 x i32], [32 x i32] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[WARPID]]
445   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT32]] to i16 addrspace([[SHARED_ADDRSPACE]])*
446   // CHECK: [[ELT_VAL:%.+]] = load i16, i16* [[ELT]], align
447   // CHECK: store volatile i16 [[ELT_VAL]], i16 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
448   // CHECK: br label {{%?}}[[COPY_CONT:.+]]
449   //
450   // CHECK: [[COPY_ELSE]]
451   // CHECK: br label {{%?}}[[COPY_CONT]]
452   //
453   // Barrier after copy to shared memory storage medium.
454   // CHECK: [[COPY_CONT]]
455   // CHECK: call void @__kmpc_barrier(%struct.ident_t* @
456   // CHECK: [[ACTIVE_WARPS:%.+]] = load i32, i32*
457   //
458   // Read into warp 0.
459   // CHECK: [[IS_W0_ACTIVE_THREAD:%.+]] = icmp ult i32 [[TID:%.+]], [[ACTIVE_WARPS]]
460   // CHECK: br i1 [[IS_W0_ACTIVE_THREAD]], label {{%?}}[[DO_READ:.+]], label {{%?}}[[READ_ELSE:.+]]
461   //
462   // CHECK: [[DO_READ]]
463   // CHECK: [[MEDIUM_ELT32:%.+]] = getelementptr inbounds [32 x i32], [32 x i32] addrspace([[SHARED_ADDRSPACE]])* [[TRANSFER_STORAGE]], i64 0, i32 [[TID]]
464   // CHECK: [[MEDIUM_ELT:%.+]] = bitcast i32 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT32]] to i16 addrspace([[SHARED_ADDRSPACE]])*
465   // CHECK: [[ELT_REF:%.+]] = getelementptr inbounds [[RLT]], [[RLT]]* [[RED_LIST:%.+]], i[[SZ]] 0, i[[SZ]] 1
466   // CHECK: [[ELT_VOID:%.+]] = load i8*, i8** [[ELT_REF]],
467   // CHECK: [[ELT:%.+]] = bitcast i8* [[ELT_VOID]] to i16*
468   // CHECK: [[MEDIUM_ELT_VAL:%.+]] = load volatile i16, i16 addrspace([[SHARED_ADDRSPACE]])* [[MEDIUM_ELT]], align
469   // CHECK: store i16 [[MEDIUM_ELT_VAL]], i16* [[ELT]], align
470   // CHECK: br label {{%?}}[[READ_CONT:.+]]
471   //
472   // CHECK: [[READ_ELSE]]
473   // CHECK: br label {{%?}}[[READ_CONT]]
474   //
475   // CHECK: [[READ_CONT]]
476   // CHECK: ret
477 
478 #endif
479