test/GPGPU/non-read-only-scalars.ll

   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 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \
   6 ; RUN: -disable-output < %s | \
   7 ; RUN: FileCheck %s -check-prefix=KERNEL-IR
   8 ;
   9 ; REQUIRES: pollyacc
  10 ;
  11 ; #include <stdio.h>
  12 ;
  13 ; float foo(float A[]) {
  14 ;   float sum = 0;
  15 ;
  16 ;   for (long i = 0; i < 32; i++)
  17 ;     A[i] = i;
  18 ;
  19 ;   for (long i = 0; i < 32; i++)
  20 ;     A[i] += i;
  21 ;
  22 ;   for (long i = 0; i < 32; i++)
  23 ;     sum += A[i];
  24 ;
  25 ;   return sum;
  26 ; }
  27 ;
  28 ; int main() {
  29 ;   float A[32];
  30 ;   float sum = foo(A);
  31 ;   printf("%f\n", sum);
  32 ; }
  33
  34 ; CODE: Code
  35 ; CODE-NEXT: ====
  36 ; CODE-NEXT: # host
  37 ; CODE-NEXT: {
  38 ; CODE-NEXT:   {
  39 ; CODE-NEXT:     dim3 k0_dimBlock(32);
  40 ; CODE-NEXT:     dim3 k0_dimGrid(1);
  41 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
  42 ; CODE-NEXT:     cudaCheckKernel();
  43 ; CODE-NEXT:   }
  44
  45 ; CODE:   {
  46 ; CODE-NEXT:     dim3 k1_dimBlock;
  47 ; CODE-NEXT:     dim3 k1_dimGrid;
  48 ; CODE-NEXT:     kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_sum_0__phi);
  49 ; CODE-NEXT:     cudaCheckKernel();
  50 ; CODE-NEXT:   }
  51
  52 ; CODE:   for (int c0 = 0; c0 <= 32; c0 += 1) {
  53 ; CODE-NEXT:     {
  54 ; CODE-NEXT:       dim3 k2_dimBlock;
  55 ; CODE-NEXT:       dim3 k2_dimGrid;
  56 ; CODE-NEXT:       kernel2 <<<k2_dimGrid, k2_dimBlock>>> (dev_MemRef_sum_0__phi, dev_MemRef_sum_0, c0);
  57 ; CODE-NEXT:       cudaCheckKernel();
  58 ; CODE-NEXT:     }
  59
  60 ; CODE:     if (c0 <= 31)
  61 ; CODE-NEXT:       {
  62 ; CODE-NEXT:         dim3 k3_dimBlock;
  63 ; CODE-NEXT:         dim3 k3_dimGrid;
  64 ; CODE-NEXT:         kernel3 <<<k3_dimGrid, k3_dimBlock>>> (dev_MemRef_A, dev_MemRef_sum_0__phi, dev_MemRef_sum_0, c0);
  65 ; CODE-NEXT:         cudaCheckKernel();
  66 ; CODE-NEXT:       }
  67
  68 ; CODE:   }
  69 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (32) * sizeof(float), cudaMemcpyDeviceToHost));
  70 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(&MemRef_sum_0__phi, dev_MemRef_sum_0__phi, sizeof(float), cudaMemcpyDeviceToHost));
  71 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(&MemRef_sum_0, dev_MemRef_sum_0, sizeof(float), cudaMemcpyDeviceToHost));
  72 ; CODE-NEXT: }
  73
  74 ; CODE: # kernel0
  75 ; CODE-NEXT: {
  76 ; CODE-NEXT:   Stmt_bb4(t0);
  77 ; CODE-NEXT:   Stmt_bb10(t0);
  78 ; CODE-NEXT: }
  79
  80 ; CODE: # kernel1
  81 ; CODE-NEXT: Stmt_bb17();
  82
  83 ; CODE: # kernel2
  84 ; CODE-NEXT: Stmt_bb18(c0);
  85
  86 ; CODE: # kernel3
  87 ; CODE-NEXT: Stmt_bb20(c0);
  88
  89 ; KERNEL-IR:       store float %p_tmp23, float* %sum.0.phiops
  90 ; KERNEL-IR-NEXT:  [[REGA:%.+]] = addrspacecast i8 addrspace(1)* %MemRef_sum_0__phi to float*
  91 ; KERNEL-IR-NEXT:  [[REGB:%.+]] = load float, float* %sum.0.phiops
  92 ; KERNEL-IR-NEXT:  store float [[REGB]], float* [[REGA]]
  93 ; KERNEL-IR-NEXT:  [[REGC:%.+]] = addrspacecast i8 addrspace(1)* %MemRef_sum_0 to float*
  94 ; KERNEL-IR-NEXT:  [[REGD:%.+]] = load float, float* %sum.0.s2a
  95 ; KERNEL-IR-NEXT:  store float [[REGD]], float* [[REGC]]
  96 ; KERNEL-IR-NEXT:  ret void
  97
  98 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
  99
 100 @.str = private unnamed_addr constant [4 x i8] c"%f\0A\00", align 1
 101
 102 define float @foo(float* %A) {
 103 bb:
 104   br label %bb3
 105
 106 bb3:                                              ; preds = %bb6, %bb
 107   %i.0 = phi i64 [ 0, %bb ], [ %tmp7, %bb6 ]
 108   %exitcond2 = icmp ne i64 %i.0, 32
 109   br i1 %exitcond2, label %bb4, label %bb8
 110
 111 bb4:                                              ; preds = %bb3
 112   %tmp = sitofp i64 %i.0 to float
 113   %tmp5 = getelementptr inbounds float, float* %A, i64 %i.0
 114   store float %tmp, float* %tmp5, align 4
 115   br label %bb6
 116
 117 bb6:                                              ; preds = %bb4
 118   %tmp7 = add nuw nsw i64 %i.0, 1
 119   br label %bb3
 120
 121 bb8:                                              ; preds = %bb3
 122   br label %bb9
 123
 124 bb9:                                              ; preds = %bb15, %bb8
 125   %i1.0 = phi i64 [ 0, %bb8 ], [ %tmp16, %bb15 ]
 126   %exitcond1 = icmp ne i64 %i1.0, 32
 127   br i1 %exitcond1, label %bb10, label %bb17
 128
 129 bb10:                                             ; preds = %bb9
 130   %tmp11 = sitofp i64 %i1.0 to float
 131   %tmp12 = getelementptr inbounds float, float* %A, i64 %i1.0
 132   %tmp13 = load float, float* %tmp12, align 4
 133   %tmp14 = fadd float %tmp13, %tmp11
 134   store float %tmp14, float* %tmp12, align 4
 135   br label %bb15
 136
 137 bb15:                                             ; preds = %bb10
 138   %tmp16 = add nuw nsw i64 %i1.0, 1
 139   br label %bb9
 140
 141 bb17:                                             ; preds = %bb9
 142   br label %bb18
 143
 144 bb18:                                             ; preds = %bb20, %bb17
 145   %sum.0 = phi float [ 0.000000e+00, %bb17 ], [ %tmp23, %bb20 ]
 146   %i2.0 = phi i64 [ 0, %bb17 ], [ %tmp24, %bb20 ]
 147   %exitcond = icmp ne i64 %i2.0, 32
 148   br i1 %exitcond, label %bb19, label %bb25
 149
 150 bb19:                                             ; preds = %bb18
 151   br label %bb20
 152
 153 bb20:                                             ; preds = %bb19
 154   %tmp21 = getelementptr inbounds float, float* %A, i64 %i2.0
 155   %tmp22 = load float, float* %tmp21, align 4
 156   %tmp23 = fadd float %sum.0, %tmp22
 157   %tmp24 = add nuw nsw i64 %i2.0, 1
 158   br label %bb18
 159
 160 bb25:                                             ; preds = %bb18
 161   %sum.0.lcssa = phi float [ %sum.0, %bb18 ]
 162   ret float %sum.0.lcssa
 163 }
 164
 165 define i32 @main() {
 166 bb:
 167   %A = alloca [32 x float], align 16
 168   %tmp = getelementptr inbounds [32 x float], [32 x float]* %A, i64 0, i64 0
 169   %tmp1 = call float @foo(float* %tmp)
 170   %tmp2 = fpext float %tmp1 to double
 171   %tmp3 = call i32 (i8*, ...) @printf(i8* getelementptr inbounds ([4 x i8], [4 x i8]* @.str, i64 0, i64 0), double %tmp2) #2
 172   ret i32 0
 173 }
 174
 175 declare i32 @printf(i8*, ...) #1
 176