test/GPGPU/host-statement.ll

   1 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
   2 ; RUN: -polly-invariant-load-hoisting=false \
   3 ; RUN: -disable-output < %s | \
   4 ; RUN: FileCheck -check-prefix=CODE %s
   5
   6 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \
   7 ; RUN: -polly-invariant-load-hoisting=false \
   8 ; RUN: -disable-output < %s | \
   9 ; RUN: FileCheck -check-prefix=KERNEL-IR %s
  10
  11 ; REQUIRES: pollyacc
  12
  13 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
  14 target triple = "x86_64-unknown-linux-gnu"
  15
  16 declare void @llvm.lifetime.start(i64, i8* nocapture) #0
  17
  18 ; This test case tests that we can correctly handle a ScopStmt that is
  19 ; scheduled on the host, instead of within a kernel.
  20
  21 ; CODE-LABEL: Code
  22 ; CODE-NEXT: ====
  23 ; CODE-NEXT: # host
  24 ; CODE-NEXT: {
  25 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (512) * (512) * sizeof(double), cudaMemcpyHostToDevice));
  26 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_R, MemRef_R, (p_0 + 1) * (512) * sizeof(double), cudaMemcpyHostToDevice));
  27 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_Q, MemRef_Q, (512) * (512) * sizeof(double), cudaMemcpyHostToDevice));
  28 ; CODE-NEXT:   {
  29 ; CODE-NEXT:     dim3 k0_dimBlock(32);
  30 ; CODE-NEXT:     dim3 k0_dimGrid(16);
  31 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p_0, p_1);
  32 ; CODE-NEXT:     cudaCheckKernel();
  33 ; CODE-NEXT:   }
  34
  35 ; CODE:   if (p_0 <= 510 && p_1 <= 510) {
  36 ; CODE-NEXT:     {
  37 ; CODE-NEXT:       dim3 k1_dimBlock(32);
  38 ; CODE-NEXT:       dim3 k1_dimGrid(p_1 <= -1048034 ? 32768 : -p_1 + floord(31 * p_1 + 30, 32) + 16);
  39 ; CODE-NEXT:       kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p_0, p_1);
  40 ; CODE-NEXT:       cudaCheckKernel();
  41 ; CODE-NEXT:     }
  42
  43 ; CODE:     {
  44 ; CODE-NEXT:       dim3 k2_dimBlock(16, 32);
  45 ; CODE-NEXT:       dim3 k2_dimGrid(16, p_1 <= -7650 ? 256 : -p_1 + floord(31 * p_1 + 30, 32) + 16);
  46 ; CODE-NEXT:       kernel2 <<<k2_dimGrid, k2_dimBlock>>> (dev_MemRef_A, dev_MemRef_R, dev_MemRef_Q, p_0, p_1);
  47 ; CODE-NEXT:       cudaCheckKernel();
  48 ; CODE-NEXT:     }
  49
  50 ; CODE:   }
  51 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (512) * (512) * sizeof(double), cudaMemcpyDeviceToHost));
  52 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(MemRef_R, dev_MemRef_R, (p_0 + 1) * (512) * sizeof(double), cudaMemcpyDeviceToHost));
  53 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(MemRef_Q, dev_MemRef_Q, (512) * (512) * sizeof(double), cudaMemcpyDeviceToHost));
  54 ; CODE-NEXT:     Stmt_for_cond33_preheader();
  55
  56 ; CODE: }
  57
  58 ; CODE: # kernel0
  59 ; CODE-NEXT: Stmt_for_body16(32 * b0 + t0);
  60
  61 ; CODE: # kernel1
  62 ; CODE-NEXT: for (int c0 = 0; c0 <= (-p_1 - 32 * b0 + 510) / 1048576; c0 += 1)
  63 ; CODE-NEXT:   for (int c1 = 0; c1 <= 15; c1 += 1) {
  64 ; CODE-NEXT:     if (p_1 + 32 * b0 + t0 + 1048576 * c0 <= 510 && c1 == 0)
  65 ; CODE-NEXT:       Stmt_for_body35(32 * b0 + t0 + 1048576 * c0);
  66 ; CODE-NEXT:     if (p_1 + 32 * b0 + t0 + 1048576 * c0 <= 510)
  67 ; CODE-NEXT:       for (int c3 = 0; c3 <= 31; c3 += 1)
  68 ; CODE-NEXT:         Stmt_for_body42(32 * b0 + t0 + 1048576 * c0, 32 * c1 + c3);
  69 ; CODE-NEXT:     sync0();
  70 ; CODE-NEXT:   }
  71
  72 ; CODE: # kernel2
  73 ; CODE-NEXT: for (int c0 = 0; c0 <= (-p_1 - 32 * b0 + 510) / 8192; c0 += 1)
  74 ; CODE-NEXT:   if (p_1 + 32 * b0 + t0 + 8192 * c0 <= 510)
  75 ; CODE-NEXT:     for (int c3 = 0; c3 <= 1; c3 += 1)
  76 ; CODE-NEXT:       Stmt_for_body62(32 * b0 + t0 + 8192 * c0, 32 * b1 + t1 + 16 * c3);
  77
  78 ; KERNEL-IR: call void @llvm.nvvm.barrier0()
  79
  80 ; Function Attrs: nounwind uwtable
  81 define internal void @kernel_gramschmidt(i32 %ni, i32 %nj, [512 x double]* %A, [512 x double]* %R, [512 x double]* %Q) #1 {
  82 entry:
  83   br label %entry.split
  84
  85 entry.split:                                      ; preds = %entry
  86   br label %for.cond1.preheader
  87
  88 for.cond1.preheader:                              ; preds = %entry.split, %for.inc86
  89   %indvars.iv24 = phi i64 [ 0, %entry.split ], [ %indvars.iv.next25, %for.inc86 ]
  90   %indvars.iv19 = phi i64 [ 1, %entry.split ], [ %indvars.iv.next20, %for.inc86 ]
  91   br label %for.inc
  92
  93 for.inc:                                          ; preds = %for.cond1.preheader, %for.inc
  94   %indvars.iv = phi i64 [ 0, %for.cond1.preheader ], [ %indvars.iv.next, %for.inc ]
  95   %nrm.02 = phi double [ 0.000000e+00, %for.cond1.preheader ], [ %add, %for.inc ]
  96   %arrayidx5 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv, i64 %indvars.iv24
  97   %tmp = load double, double* %arrayidx5, align 8, !tbaa !1
  98   %arrayidx9 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv, i64 %indvars.iv24
  99   %tmp27 = load double, double* %arrayidx9, align 8, !tbaa !1
 100   %mul = fmul double %tmp, %tmp27
 101   %add = fadd double %nrm.02, %mul
 102   %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
 103   %exitcond = icmp ne i64 %indvars.iv.next, 512
 104   br i1 %exitcond, label %for.inc, label %for.end
 105
 106 for.end:                                          ; preds = %for.inc
 107   %add.lcssa = phi double [ %add, %for.inc ]
 108   %call = tail call double @sqrt(double %add.lcssa) #2
 109   %arrayidx13 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv24
 110   store double %call, double* %arrayidx13, align 8, !tbaa !1
 111   br label %for.body16
 112
 113 for.cond33.preheader:                             ; preds = %for.body16
 114   %indvars.iv.next25 = add nuw nsw i64 %indvars.iv24, 1
 115   %cmp347 = icmp slt i64 %indvars.iv.next25, 512
 116   br i1 %cmp347, label %for.body35.lr.ph, label %for.inc86
 117
 118 for.body35.lr.ph:                                 ; preds = %for.cond33.preheader
 119   br label %for.body35
 120
 121 for.body16:                                       ; preds = %for.end, %for.body16
 122   %indvars.iv10 = phi i64 [ 0, %for.end ], [ %indvars.iv.next11, %for.body16 ]
 123   %arrayidx20 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv10, i64 %indvars.iv24
 124   %tmp28 = load double, double* %arrayidx20, align 8, !tbaa !1
 125   %arrayidx24 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv24
 126   %tmp29 = load double, double* %arrayidx24, align 8, !tbaa !1
 127   %div = fdiv double %tmp28, %tmp29
 128   %arrayidx28 = getelementptr inbounds [512 x double], [512 x double]* %Q, i64 %indvars.iv10, i64 %indvars.iv24
 129   store double %div, double* %arrayidx28, align 8, !tbaa !1
 130   %indvars.iv.next11 = add nuw nsw i64 %indvars.iv10, 1
 131   %exitcond12 = icmp ne i64 %indvars.iv.next11, 512
 132   br i1 %exitcond12, label %for.body16, label %for.cond33.preheader
 133
 134 for.cond33.loopexit:                              ; preds = %for.body62
 135   %indvars.iv.next22 = add nuw nsw i64 %indvars.iv21, 1
 136   %lftr.wideiv = trunc i64 %indvars.iv.next22 to i32
 137   %exitcond23 = icmp ne i32 %lftr.wideiv, 512
 138   br i1 %exitcond23, label %for.body35, label %for.cond33.for.inc86_crit_edge
 139
 140 for.body35:                                       ; preds = %for.body35.lr.ph, %for.cond33.loopexit
 141   %indvars.iv21 = phi i64 [ %indvars.iv19, %for.body35.lr.ph ], [ %indvars.iv.next22, %for.cond33.loopexit ]
 142   %arrayidx39 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv21
 143   store double 0.000000e+00, double* %arrayidx39, align 8, !tbaa !1
 144   br label %for.body42
 145
 146 for.cond60.preheader:                             ; preds = %for.body42
 147   br label %for.body62
 148
 149 for.body42:                                       ; preds = %for.body35, %for.body42
 150   %indvars.iv13 = phi i64 [ 0, %for.body35 ], [ %indvars.iv.next14, %for.body42 ]
 151   %arrayidx46 = getelementptr inbounds [512 x double], [512 x double]* %Q, i64 %indvars.iv13, i64 %indvars.iv24
 152   %tmp30 = load double, double* %arrayidx46, align 8, !tbaa !1
 153   %arrayidx50 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv13, i64 %indvars.iv21
 154   %tmp31 = load double, double* %arrayidx50, align 8, !tbaa !1
 155   %mul51 = fmul double %tmp30, %tmp31
 156   %arrayidx55 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv21
 157   %tmp32 = load double, double* %arrayidx55, align 8, !tbaa !1
 158   %add56 = fadd double %tmp32, %mul51
 159   store double %add56, double* %arrayidx55, align 8, !tbaa !1
 160   %indvars.iv.next14 = add nuw nsw i64 %indvars.iv13, 1
 161   %exitcond15 = icmp ne i64 %indvars.iv.next14, 512
 162   br i1 %exitcond15, label %for.body42, label %for.cond60.preheader
 163
 164 for.body62:                                       ; preds = %for.cond60.preheader, %for.body62
 165   %indvars.iv16 = phi i64 [ 0, %for.cond60.preheader ], [ %indvars.iv.next17, %for.body62 ]
 166   %arrayidx66 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv16, i64 %indvars.iv21
 167   %tmp33 = load double, double* %arrayidx66, align 8, !tbaa !1
 168   %arrayidx70 = getelementptr inbounds [512 x double], [512 x double]* %Q, i64 %indvars.iv16, i64 %indvars.iv24
 169   %tmp34 = load double, double* %arrayidx70, align 8, !tbaa !1
 170   %arrayidx74 = getelementptr inbounds [512 x double], [512 x double]* %R, i64 %indvars.iv24, i64 %indvars.iv21
 171   %tmp35 = load double, double* %arrayidx74, align 8, !tbaa !1
 172   %mul75 = fmul double %tmp34, %tmp35
 173   %sub = fsub double %tmp33, %mul75
 174   %arrayidx79 = getelementptr inbounds [512 x double], [512 x double]* %A, i64 %indvars.iv16, i64 %indvars.iv21
 175   store double %sub, double* %arrayidx79, align 8, !tbaa !1
 176   %indvars.iv.next17 = add nuw nsw i64 %indvars.iv16, 1
 177   %exitcond18 = icmp ne i64 %indvars.iv.next17, 512
 178   br i1 %exitcond18, label %for.body62, label %for.cond33.loopexit
 179
 180 for.cond33.for.inc86_crit_edge:                   ; preds = %for.cond33.loopexit
 181   br label %for.inc86
 182
 183 for.inc86:                                        ; preds = %for.cond33.for.inc86_crit_edge, %for.cond33.preheader
 184   %indvars.iv.next20 = add nuw nsw i64 %indvars.iv19, 1
 185   %exitcond26 = icmp ne i64 %indvars.iv.next25, 512
 186   br i1 %exitcond26, label %for.cond1.preheader, label %for.end88
 187
 188 for.end88:                                        ; preds = %for.inc86
 189   ret void
 190 }
 191
 192 ; Function Attrs: argmemonly nounwind
 193 declare void @llvm.lifetime.end(i64, i8* nocapture) #0
 194
 195 ; Function Attrs: nounwind
 196 declare double @sqrt(double) #2
 197
 198 attributes #0 = { argmemonly nounwind }
 199 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" }
 200 attributes #2 = { nounwind }
 201
 202 !llvm.ident = !{!0}
 203
 204 !0 = !{!"clang version 3.9.0 (trunk 275267) (llvm/trunk 275268)"}
 205 !1 = !{!2, !2, i64 0}
 206 !2 = !{!"double", !3, i64 0}
 207 !3 = !{!"omnipotent char", !4, i64 0}
 208 !4 = !{!"Simple C/C++ TBAA"}