1; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \ 2; RUN: -disable-output < %s | \ 3; RUN: FileCheck -check-prefix=CODE %s 4 5; REQUIRES: pollyacc 6 7; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_global_1, MemRef_global_1, (142) * sizeof(i32), cudaMemcpyHostToDevice)); 8; CODE-NEXT: { 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(); 13; CODE-NEXT: } 14 15; CODE: cudaCheckReturn(cudaMemcpy(MemRef_global_1, dev_MemRef_global_1, (142) * sizeof(i32), cudaMemcpyDeviceToHost)); 16; CODE: cudaCheckReturn(cudaFree(dev_MemRef_global_1)); 17; CODE-NEXT: } 18 19; CODE: # kernel0 20; CODE-NEXT: Stmt_bb33(t0, 0); 21 22 23target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" 24target triple = "x86_64-unknown-linux-gnu" 25 26%struct.hoge = type { [23 x i16], [22 x i16], [14 x i16], [13 x i16] } 27 28@global = external global [9 x %struct.hoge], align 16 29@global.1 = external global [9 x [152 x i32]], align 16 30 31; Function Attrs: nounwind uwtable 32define void @widget() #0 { 33bb: 34 br label %bb1 35 36bb1: ; preds = %bb1, %bb 37 br i1 undef, label %bb1, label %bb2 38 39bb2: ; preds = %bb2, %bb1 40 br i1 undef, label %bb2, label %bb3 41 42bb3: ; preds = %bb3, %bb2 43 br i1 undef, label %bb3, label %bb4 44 45bb4: ; preds = %bb4, %bb3 46 br i1 undef, label %bb4, label %bb5 47 48bb5: ; preds = %bb5, %bb4 49 br i1 undef, label %bb5, label %bb6 50 51bb6: ; preds = %bb6, %bb5 52 br i1 undef, label %bb6, label %bb7 53 54bb7: ; preds = %bb7, %bb6 55 br i1 undef, label %bb7, label %bb8 56 57bb8: ; preds = %bb8, %bb7 58 br i1 undef, label %bb8, label %bb9 59 60bb9: ; preds = %bb8 61 br label %bb10 62 63bb10: ; preds = %bb12, %bb9 64 br label %bb11 65 66bb11: ; preds = %bb11, %bb10 67 br i1 undef, label %bb11, label %bb12 68 69bb12: ; preds = %bb11 70 br i1 undef, label %bb10, label %bb13 71 72bb13: ; preds = %bb18, %bb12 73 br i1 undef, label %bb16, label %bb14 74 75bb14: ; preds = %bb16, %bb13 76 br i1 undef, label %bb15, label %bb18 77 78bb15: ; preds = %bb14 79 br label %bb17 80 81bb16: ; preds = %bb16, %bb13 82 br i1 undef, label %bb16, label %bb14 83 84bb17: ; preds = %bb17, %bb15 85 br i1 undef, label %bb17, label %bb18 86 87bb18: ; preds = %bb17, %bb14 88 br i1 undef, label %bb13, label %bb19 89 90bb19: ; preds = %bb25, %bb18 91 br label %bb20 92 93bb20: ; preds = %bb24, %bb19 94 br i1 undef, label %bb21, label %bb24 95 96bb21: ; preds = %bb20 97 br i1 undef, label %bb23, label %bb22 98 99bb22: ; preds = %bb21 100 br label %bb24 101 102bb23: ; preds = %bb21 103 br label %bb24 104 105bb24: ; preds = %bb23, %bb22, %bb20 106 br i1 undef, label %bb20, label %bb25 107 108bb25: ; preds = %bb24 109 br i1 undef, label %bb19, label %bb26 110 111bb26: ; 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 ] 113 br label %bb27 114 115bb27: ; preds = %bb27, %bb26 116 br i1 undef, label %bb27, label %bb28 117 118bb28: ; preds = %bb27 119 %tmp29 = bitcast [9 x [152 x i32]]* %tmp to i32* 120 br label %bb30 121 122bb30: ; preds = %bb38, %bb28 123 %tmp31 = phi i32 [ 3, %bb28 ], [ %tmp40, %bb38 ] 124 %tmp32 = phi i32* [ %tmp29, %bb28 ], [ %tmp39, %bb38 ] 125 br label %bb33 126 127bb33: ; 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 134 135bb38: ; 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 140 141bb42: ; preds = %bb38 142 %tmp43 = getelementptr inbounds [9 x %struct.hoge], [9 x %struct.hoge]* @global, i64 0, i64 0, i32 3, i64 0 143 br label %bb44 144 145bb44: ; 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 149 br label %bb48 150 151bb48: ; 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 155 156bb51: ; 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 160 161bb54: ; preds = %bb51 162 br label %bb55 163 164bb55: ; preds = %bb55, %bb54 165 br i1 undef, label %bb55, label %bb56 166 167bb56: ; preds = %bb55 168 br i1 undef, label %bb26, label %bb57 169 170bb57: ; preds = %bb60, %bb56 171 br label %bb58 172 173bb58: ; preds = %bb58, %bb57 174 br i1 undef, label %bb58, label %bb59 175 176bb59: ; preds = %bb59, %bb58 177 br i1 undef, label %bb59, label %bb60 178 179bb60: ; preds = %bb59 180 br i1 undef, label %bb57, label %bb61 181 182bb61: ; preds = %bb65, %bb60 183 br label %bb62 184 185bb62: ; preds = %bb64, %bb61 186 br label %bb63 187 188bb63: ; preds = %bb63, %bb62 189 br i1 undef, label %bb63, label %bb64 190 191bb64: ; preds = %bb63 192 br i1 undef, label %bb62, label %bb65 193 194bb65: ; preds = %bb64 195 br i1 undef, label %bb61, label %bb66 196 197bb66: ; preds = %bb70, %bb65 198 br label %bb67 199 200bb67: ; preds = %bb69, %bb66 201 br label %bb68 202 203bb68: ; preds = %bb68, %bb67 204 br i1 undef, label %bb68, label %bb69 205 206bb69: ; preds = %bb68 207 br i1 undef, label %bb67, label %bb70 208 209bb70: ; preds = %bb69 210 br i1 undef, label %bb66, label %bb71 211 212bb71: ; preds = %bb73, %bb70 213 br label %bb72 214 215bb72: ; preds = %bb72, %bb71 216 br i1 undef, label %bb72, label %bb73 217 218bb73: ; preds = %bb72 219 br i1 undef, label %bb71, label %bb74 220 221bb74: ; preds = %bb80, %bb73 222 br label %bb75 223 224bb75: ; preds = %bb79, %bb74 225 br label %bb76 226 227bb76: ; preds = %bb78, %bb75 228 br label %bb77 229 230bb77: ; preds = %bb77, %bb76 231 br i1 undef, label %bb77, label %bb78 232 233bb78: ; preds = %bb77 234 br i1 undef, label %bb76, label %bb79 235 236bb79: ; preds = %bb78 237 br i1 undef, label %bb75, label %bb80 238 239bb80: ; preds = %bb79 240 br i1 undef, label %bb74, label %bb81 241 242bb81: ; preds = %bb85, %bb80 243 br label %bb82 244 245bb82: ; preds = %bb84, %bb81 246 br label %bb83 247 248bb83: ; preds = %bb83, %bb82 249 br i1 undef, label %bb83, label %bb84 250 251bb84: ; preds = %bb83 252 br i1 undef, label %bb82, label %bb85 253 254bb85: ; preds = %bb84 255 br i1 undef, label %bb81, label %bb86 256 257bb86: ; preds = %bb85 258 ret void 259} 260 261attributes #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" } 262 263!llvm.ident = !{!0} 264 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} 272