1 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
2 ; RUN: -disable-output < %s | \
3 ; RUN: FileCheck -check-prefix=CODE %s
7 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_global_1, MemRef_global_1, (142) * sizeof(i32), cudaMemcpyHostToDevice));
9 ; CODE-NEXT: dim3 k0_dimBlock(10);
10 ; CODE-NEXT: dim3 k0_dimGrid(1);
11 ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_global_1);
12 ; CODE-NEXT: cudaCheckKernel();
15 ; CODE: cudaCheckReturn(cudaMemcpy(MemRef_global_1, dev_MemRef_global_1, (142) * sizeof(i32), cudaMemcpyDeviceToHost));
16 ; CODE: cudaCheckReturn(cudaFree(dev_MemRef_global_1));
20 ; CODE-NEXT: Stmt_bb33(t0, 0);
23 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
24 target triple = "x86_64-unknown-linux-gnu"
26 %struct.hoge = type { [23 x i16], [22 x i16], [14 x i16], [13 x i16] }
28 @global = external global [9 x %struct.hoge], align 16
29 @global.1 = external global [9 x [152 x i32]], align 16
31 ; Function Attrs: nounwind uwtable
32 define void @widget() #0 {
36 bb1: ; preds = %bb1, %bb
37 br i1 undef, label %bb1, label %bb2
39 bb2: ; preds = %bb2, %bb1
40 br i1 undef, label %bb2, label %bb3
42 bb3: ; preds = %bb3, %bb2
43 br i1 undef, label %bb3, label %bb4
45 bb4: ; preds = %bb4, %bb3
46 br i1 undef, label %bb4, label %bb5
48 bb5: ; preds = %bb5, %bb4
49 br i1 undef, label %bb5, label %bb6
51 bb6: ; preds = %bb6, %bb5
52 br i1 undef, label %bb6, label %bb7
54 bb7: ; preds = %bb7, %bb6
55 br i1 undef, label %bb7, label %bb8
57 bb8: ; preds = %bb8, %bb7
58 br i1 undef, label %bb8, label %bb9
63 bb10: ; preds = %bb12, %bb9
66 bb11: ; preds = %bb11, %bb10
67 br i1 undef, label %bb11, label %bb12
70 br i1 undef, label %bb10, label %bb13
72 bb13: ; preds = %bb18, %bb12
73 br i1 undef, label %bb16, label %bb14
75 bb14: ; preds = %bb16, %bb13
76 br i1 undef, label %bb15, label %bb18
81 bb16: ; preds = %bb16, %bb13
82 br i1 undef, label %bb16, label %bb14
84 bb17: ; preds = %bb17, %bb15
85 br i1 undef, label %bb17, label %bb18
87 bb18: ; preds = %bb17, %bb14
88 br i1 undef, label %bb13, label %bb19
90 bb19: ; preds = %bb25, %bb18
93 bb20: ; preds = %bb24, %bb19
94 br i1 undef, label %bb21, label %bb24
97 br i1 undef, label %bb23, label %bb22
102 bb23: ; preds = %bb21
105 bb24: ; preds = %bb23, %bb22, %bb20
106 br i1 undef, label %bb20, label %bb25
108 bb25: ; preds = %bb24
109 br i1 undef, label %bb19, label %bb26
111 bb26: ; preds = %bb56, %bb25
112 %tmp = phi [9 x [152 x i32]]* [ undef, %bb56 ], [ bitcast (i32* getelementptr inbounds ([9 x [152 x i32]], [9 x [152 x i32]]* @global.1, i64 0, i64 0, i64 32) to [9 x [152 x i32]]*), %bb25 ]
115 bb27: ; preds = %bb27, %bb26
116 br i1 undef, label %bb27, label %bb28
118 bb28: ; preds = %bb27
119 %tmp29 = bitcast [9 x [152 x i32]]* %tmp to i32*
122 bb30: ; preds = %bb38, %bb28
123 %tmp31 = phi i32 [ 3, %bb28 ], [ %tmp40, %bb38 ]
124 %tmp32 = phi i32* [ %tmp29, %bb28 ], [ %tmp39, %bb38 ]
127 bb33: ; preds = %bb33, %bb30
128 %tmp34 = phi i32 [ 0, %bb30 ], [ %tmp37, %bb33 ]
129 %tmp35 = phi i32* [ %tmp32, %bb30 ], [ undef, %bb33 ]
130 %tmp36 = getelementptr inbounds i32, i32* %tmp35, i64 1
131 store i32 undef, i32* %tmp36, align 4, !tbaa !1
132 %tmp37 = add nuw nsw i32 %tmp34, 1
133 br i1 false, label %bb33, label %bb38
135 bb38: ; preds = %bb33
136 %tmp39 = getelementptr i32, i32* %tmp32, i64 12
137 %tmp40 = add nuw nsw i32 %tmp31, 1
138 %tmp41 = icmp ne i32 %tmp40, 13
139 br i1 %tmp41, label %bb30, label %bb42
141 bb42: ; preds = %bb38
142 %tmp43 = getelementptr inbounds [9 x %struct.hoge], [9 x %struct.hoge]* @global, i64 0, i64 0, i32 3, i64 0
145 bb44: ; preds = %bb51, %bb42
146 %tmp45 = phi i32 [ 0, %bb42 ], [ %tmp52, %bb51 ]
147 %tmp46 = phi i16* [ %tmp43, %bb42 ], [ undef, %bb51 ]
148 %tmp47 = load i16, i16* %tmp46, align 2, !tbaa !5
151 bb48: ; preds = %bb48, %bb44
152 %tmp49 = phi i32 [ 0, %bb44 ], [ %tmp50, %bb48 ]
153 %tmp50 = add nuw nsw i32 %tmp49, 1
154 br i1 false, label %bb48, label %bb51
156 bb51: ; preds = %bb48
157 %tmp52 = add nuw nsw i32 %tmp45, 1
158 %tmp53 = icmp ne i32 %tmp52, 13
159 br i1 %tmp53, label %bb44, label %bb54
161 bb54: ; preds = %bb51
164 bb55: ; preds = %bb55, %bb54
165 br i1 undef, label %bb55, label %bb56
167 bb56: ; preds = %bb55
168 br i1 undef, label %bb26, label %bb57
170 bb57: ; preds = %bb60, %bb56
173 bb58: ; preds = %bb58, %bb57
174 br i1 undef, label %bb58, label %bb59
176 bb59: ; preds = %bb59, %bb58
177 br i1 undef, label %bb59, label %bb60
179 bb60: ; preds = %bb59
180 br i1 undef, label %bb57, label %bb61
182 bb61: ; preds = %bb65, %bb60
185 bb62: ; preds = %bb64, %bb61
188 bb63: ; preds = %bb63, %bb62
189 br i1 undef, label %bb63, label %bb64
191 bb64: ; preds = %bb63
192 br i1 undef, label %bb62, label %bb65
194 bb65: ; preds = %bb64
195 br i1 undef, label %bb61, label %bb66
197 bb66: ; preds = %bb70, %bb65
200 bb67: ; preds = %bb69, %bb66
203 bb68: ; preds = %bb68, %bb67
204 br i1 undef, label %bb68, label %bb69
206 bb69: ; preds = %bb68
207 br i1 undef, label %bb67, label %bb70
209 bb70: ; preds = %bb69
210 br i1 undef, label %bb66, label %bb71
212 bb71: ; preds = %bb73, %bb70
215 bb72: ; preds = %bb72, %bb71
216 br i1 undef, label %bb72, label %bb73
218 bb73: ; preds = %bb72
219 br i1 undef, label %bb71, label %bb74
221 bb74: ; preds = %bb80, %bb73
224 bb75: ; preds = %bb79, %bb74
227 bb76: ; preds = %bb78, %bb75
230 bb77: ; preds = %bb77, %bb76
231 br i1 undef, label %bb77, label %bb78
233 bb78: ; preds = %bb77
234 br i1 undef, label %bb76, label %bb79
236 bb79: ; preds = %bb78
237 br i1 undef, label %bb75, label %bb80
239 bb80: ; preds = %bb79
240 br i1 undef, label %bb74, label %bb81
242 bb81: ; preds = %bb85, %bb80
245 bb82: ; preds = %bb84, %bb81
248 bb83: ; preds = %bb83, %bb82
249 br i1 undef, label %bb83, label %bb84
251 bb84: ; preds = %bb83
252 br i1 undef, label %bb82, label %bb85
254 bb85: ; preds = %bb84
255 br i1 undef, label %bb81, label %bb86
257 bb86: ; preds = %bb85
261 attributes #0 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "frame-pointer"="none" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
265 !0 = !{!"clang version 4.0.0"}
266 !1 = !{!2, !2, i64 0}
267 !2 = !{!"int", !3, i64 0}
268 !3 = !{!"omnipotent char", !4, i64 0}
269 !4 = !{!"Simple C/C++ TBAA"}
270 !5 = !{!6, !6, i64 0}
271 !6 = !{!"short", !3, i64 0}