| ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \ |
| ; RUN: -polly-invariant-load-hoisting=false \ |
| ; RUN: -disable-output < %s | \ |
| ; RUN: FileCheck -check-prefix=CODE %s |
| |
| ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \ |
| ; RUN: -polly-invariant-load-hoisting=false \ |
| ; RUN: -disable-output < %s | \ |
| ; RUN: FileCheck -check-prefix=KERNEL-IR %s |
| |
| ; REQUIRES: pollyacc |
| |
| target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" |
| target triple = "x86_64-unknown-linux-gnu" |
| |
| declare void @llvm.lifetime.start(i64, i8* nocapture) #0 |
| |
| ; This test case tests that we can correctly handle a ScopStmt that is |
| ; scheduled on the host, instead of within a kernel. |
| |
| ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (512) * (512) * sizeof(double), cudaMemcpyHostToDevice)); |
| ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_R, MemRef_R, (p_0 + 1) * (512) * sizeof(double), cudaMemcpyHostToDevice)); |
| ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_Q, MemRef_Q, (512) * (512) * sizeof(double), cudaMemcpyHostToDevice)); |
| ; CODE-NEXT: { |
| ; CODE-NEXT: dim3 k0_dimBlock(32); |
| ; CODE-NEXT: dim3 k0_dimGrid(16); |
| ; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p_0, p_1); |
| ; CODE-NEXT: cudaCheckKernel(); |
| ; CODE-NEXT: } |
| |
| ; CODE: if (p_0 <= 510 && p_1 <= 510) { |
| ; CODE-NEXT: { |
| ; CODE-NEXT: dim3 k1_dimBlock(32); |
| ; CODE-NEXT: dim3 k1_dimGrid(p_1 <= -1048034 ? 32768 : -p_1 + floord(31 * p_1 + 30, 32) + 16); |
| ; CODE-NEXT: kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p_0, p_1); |
| ; CODE-NEXT: cudaCheckKernel(); |
| ; CODE-NEXT: } |
| |
| ; CODE: { |
| ; CODE-NEXT: dim3 k2_dimBlock(16, 32); |
| ; CODE-NEXT: dim3 k2_dimGrid(16, p_1 <= -7650 ? 256 : -p_1 + floord(31 * p_1 + 30, 32) + 16); |
| ; CODE-NEXT: kernel2 <<<k2_dimGrid, k2_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p_0, p_1); |
| ; CODE-NEXT: cudaCheckKernel(); |
| ; CODE-NEXT: } |
| |
| ; CODE: } |
| ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (512) * (512) * sizeof(double), cudaMemcpyDeviceToHost)); |
| ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(MemRef_R, dev_MemRef_R, (p_0 + 1) * (512) * sizeof(double), cudaMemcpyDeviceToHost)); |
| ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(MemRef_Q, dev_MemRef_Q, (512) * (512) * sizeof(double), cudaMemcpyDeviceToHost)); |
| ; CODE-NEXT: Stmt_for_cond33_preheader_last(); |
| |
| ; CODE: } |
| |
| ; CODE: # kernel0 |
| ; CODE-NEXT: Stmt_for_body16(32 * b0 + t0); |
| |
| ; CODE: # kernel1 |
| ; CODE-NEXT: for (int c0 = 0; c0 <= (-p_1 - 32 * b0 + 510) / 1048576; c0 += 1) |
| ; CODE-NEXT: for (int c1 = 0; c1 <= 15; c1 += 1) { |
| ; CODE-NEXT: if (p_1 + 32 * b0 + t0 + 1048576 * c0 <= 510 && c1 == 0) |
| ; CODE-NEXT: Stmt_for_body35(32 * b0 + t0 + 1048576 * c0); |
| ; CODE-NEXT: if (p_1 + 32 * b0 + t0 + 1048576 * c0 <= 510) |
| ; CODE-NEXT: for (int c3 = 0; c3 <= 31; c3 += 1) |
| ; CODE-NEXT: Stmt_for_body42(32 * b0 + t0 + 1048576 * c0, 32 * c1 + c3); |
| ; CODE-NEXT: sync0(); |
| ; CODE-NEXT: } |
| |
| ; CODE: # kernel2 |
| ; CODE-NEXT: for (int c0 = 0; c0 <= (-p_1 - 32 * b0 + 510) / 8192; c0 += 1) |
| ; CODE-NEXT: if (p_1 + 32 * b0 + t0 + 8192 * c0 <= 510) |
| ; CODE-NEXT: for (int c3 = 0; c3 <= 1; c3 += 1) |
| ; CODE-NEXT: Stmt_for_body62(32 * b0 + t0 + 8192 * c0, 32 * b1 + t1 + 16 * c3); |
| |
| ; KERNEL-IR: call void @llvm.nvvm.barrier0() |
| |
| ; Function Attrs: nounwind uwtable |
| define internal void @kernel_gramschmidt(i32 %ni, i32 %nj, [512 x double]* %A, [512 x double]* %R, [512 x double]* %Q) #1 { |
| entry: |
| br label %entry.split |
| |
| entry.split: ; preds = %entry |
| br label %for.cond1.preheader |
| |
| for.cond1.preheader: ; preds = %entry.split, %for.inc86 |
| %indvars.iv24 = phi i64 [ 0, %entry.split ], [ %indvars.iv.next25, %for.inc86 ] |
| %indvars.iv19 = phi i64 [ 1, %entry.split ], [ %indvars.iv.next20, %for.inc86 ] |
| br label %for.inc |
| |
| for.inc: ; preds = %for.cond1.preheader, %for.inc |
| %indvars.iv = phi i64 [ 0, %for.cond1.preheader ], [ %indvars.iv.next, %for.inc ] |
| %nrm.02 = phi double [ 0.000000e+00, %for.cond1.preheader ], [ %add, %for.inc ] |
| %arrayidx5 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv, i64 %indvars.iv24 |
| %tmp = load double, double* %arrayidx5, align 8, !tbaa !1 |
| %arrayidx9 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv, i64 %indvars.iv24 |
| %tmp27 = load double, double* %arrayidx9, align 8, !tbaa !1 |
| %mul = fmul double %tmp, %tmp27 |
| %add = fadd double %nrm.02, %mul |
| %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1 |
| %exitcond = icmp ne i64 %indvars.iv.next, 512 |
| br i1 %exitcond, label %for.inc, label %for.end |
| |
| for.end: ; preds = %for.inc |
| %add.lcssa = phi double [ %add, %for.inc ] |
| %call = tail call double @sqrt(double %add.lcssa) #2 |
| %arrayidx13 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv24 |
| store double %call, double* %arrayidx13, align 8, !tbaa !1 |
| br label %for.body16 |
| |
| for.cond33.preheader: ; preds = %for.body16 |
| %indvars.iv.next25 = add nuw nsw i64 %indvars.iv24, 1 |
| %cmp347 = icmp slt i64 %indvars.iv.next25, 512 |
| br i1 %cmp347, label %for.body35.lr.ph, label %for.inc86 |
| |
| for.body35.lr.ph: ; preds = %for.cond33.preheader |
| br label %for.body35 |
| |
| for.body16: ; preds = %for.end, %for.body16 |
| %indvars.iv10 = phi i64 [ 0, %for.end ], [ %indvars.iv.next11, %for.body16 ] |
| %arrayidx20 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv10, i64 %indvars.iv24 |
| %tmp28 = load double, double* %arrayidx20, align 8, !tbaa !1 |
| %arrayidx24 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv24 |
| %tmp29 = load double, double* %arrayidx24, align 8, !tbaa !1 |
| %div = fdiv double %tmp28, %tmp29 |
| %arrayidx28 = getelementptr inbounds [512 x double], [512 x double]* %Q, i64 %indvars.iv10, i64 %indvars.iv24 |
| store double %div, double* %arrayidx28, align 8, !tbaa !1 |
| %indvars.iv.next11 = add nuw nsw i64 %indvars.iv10, 1 |
| %exitcond12 = icmp ne i64 %indvars.iv.next11, 512 |
| br i1 %exitcond12, label %for.body16, label %for.cond33.preheader |
| |
| for.cond33.loopexit: ; preds = %for.body62 |
| %indvars.iv.next22 = add nuw nsw i64 %indvars.iv21, 1 |
| %lftr.wideiv = trunc i64 %indvars.iv.next22 to i32 |
| %exitcond23 = icmp ne i32 %lftr.wideiv, 512 |
| br i1 %exitcond23, label %for.body35, label %for.cond33.for.inc86_crit_edge |
| |
| for.body35: ; preds = %for.body35.lr.ph, %for.cond33.loopexit |
| %indvars.iv21 = phi i64 [ %indvars.iv19, %for.body35.lr.ph ], [ %indvars.iv.next22, %for.cond33.loopexit ] |
| %arrayidx39 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv21 |
| store double 0.000000e+00, double* %arrayidx39, align 8, !tbaa !1 |
| br label %for.body42 |
| |
| for.cond60.preheader: ; preds = %for.body42 |
| br label %for.body62 |
| |
| for.body42: ; preds = %for.body35, %for.body42 |
| %indvars.iv13 = phi i64 [ 0, %for.body35 ], [ %indvars.iv.next14, %for.body42 ] |
| %arrayidx46 = getelementptr inbounds [512 x double], [512 x double]* %Q, i64 %indvars.iv13, i64 %indvars.iv24 |
| %tmp30 = load double, double* %arrayidx46, align 8, !tbaa !1 |
| %arrayidx50 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv13, i64 %indvars.iv21 |
| %tmp31 = load double, double* %arrayidx50, align 8, !tbaa !1 |
| %mul51 = fmul double %tmp30, %tmp31 |
| %arrayidx55 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv21 |
| %tmp32 = load double, double* %arrayidx55, align 8, !tbaa !1 |
| %add56 = fadd double %tmp32, %mul51 |
| store double %add56, double* %arrayidx55, align 8, !tbaa !1 |
| %indvars.iv.next14 = add nuw nsw i64 %indvars.iv13, 1 |
| %exitcond15 = icmp ne i64 %indvars.iv.next14, 512 |
| br i1 %exitcond15, label %for.body42, label %for.cond60.preheader |
| |
| for.body62: ; preds = %for.cond60.preheader, %for.body62 |
| %indvars.iv16 = phi i64 [ 0, %for.cond60.preheader ], [ %indvars.iv.next17, %for.body62 ] |
| %arrayidx66 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv16, i64 %indvars.iv21 |
| %tmp33 = load double, double* %arrayidx66, align 8, !tbaa !1 |
| %arrayidx70 = getelementptr inbounds [512 x double], [512 x double]* %Q, i64 %indvars.iv16, i64 %indvars.iv24 |
| %tmp34 = load double, double* %arrayidx70, align 8, !tbaa !1 |
| %arrayidx74 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv21 |
| %tmp35 = load double, double* %arrayidx74, align 8, !tbaa !1 |
| %mul75 = fmul double %tmp34, %tmp35 |
| %sub = fsub double %tmp33, %mul75 |
| %arrayidx79 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv16, i64 %indvars.iv21 |
| store double %sub, double* %arrayidx79, align 8, !tbaa !1 |
| %indvars.iv.next17 = add nuw nsw i64 %indvars.iv16, 1 |
| %exitcond18 = icmp ne i64 %indvars.iv.next17, 512 |
| br i1 %exitcond18, label %for.body62, label %for.cond33.loopexit |
| |
| for.cond33.for.inc86_crit_edge: ; preds = %for.cond33.loopexit |
| br label %for.inc86 |
| |
| for.inc86: ; preds = %for.cond33.for.inc86_crit_edge, %for.cond33.preheader |
| %indvars.iv.next20 = add nuw nsw i64 %indvars.iv19, 1 |
| %exitcond26 = icmp ne i64 %indvars.iv.next25, 512 |
| br i1 %exitcond26, label %for.cond1.preheader, label %for.end88 |
| |
| for.end88: ; preds = %for.inc86 |
| ret void |
| } |
| |
| ; Function Attrs: argmemonly nounwind |
| declare void @llvm.lifetime.end(i64, i8* nocapture) #0 |
| |
| ; Function Attrs: nounwind |
| declare double @sqrt(double) #2 |
| |
| attributes #0 = { argmemonly nounwind } |
| attributes #1 = { nounwind uwtable "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-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" } |
| attributes #2 = { nounwind } |
| |
| !llvm.ident = !{!0} |
| |
| !0 = !{!"clang version 3.9.0 (trunk 275267) (llvm/trunk 275268)"} |
| !1 = !{!2, !2, i64 0} |
| !2 = !{!"double", !3, i64 0} |
| !3 = !{!"omnipotent char", !4, i64 0} |
| !4 = !{!"Simple C/C++ TBAA"} |
