test/GPGPU/non-zero-array-offset.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 -S < %s | \
   6 ; RUN: FileCheck %s -check-prefix=IR
   7 ;
   8 ; REQUIRES: pollyacc
   9
  10 ; CODE:      cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (16) * sizeof(float), cudaMemcpyHostToDevice));
  11 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (8) * sizeof(float), cudaMemcpyHostToDevice));
  12
  13 ; CODE:          dim3 k0_dimBlock(8);
  14 ; CODE-NEXT:     dim3 k0_dimGrid(1);
  15 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
  16 ; CODE-NEXT:     cudaCheckKernel();
  17 ; CODE-NEXT:   }
  18
  19 ; CODE:        {
  20 ; CODE-NEXT:     dim3 k1_dimBlock(8);
  21 ; CODE-NEXT:     dim3 k1_dimGrid(1);
  22 ; CODE-NEXT:     kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_B);
  23 ; CODE-NEXT:     cudaCheckKernel();
  24 ; CODE-NEXT:   }
  25
  26 ; CODE:   cudaCheckReturn(cudaMemcpy(MemRef_B, dev_MemRef_B, (16) * sizeof(float), cudaMemcpyDeviceToHost));
  27 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (8) * sizeof(float), cudaMemcpyDeviceToHost));
  28
  29 ; CODE: # kernel0
  30 ; CODE-NEXT: Stmt_bb11(t0);
  31
  32 ; CODE: # kernel1
  33 ; CODE-NEXT: Stmt_bb3(t0);
  34
  35 ; IR:       %p_dev_array_MemRef_B = call i8* @polly_allocateMemoryForDevice(i64 32)
  36 ; IR-NEXT:  %p_dev_array_MemRef_A = call i8* @polly_allocateMemoryForDevice(i64 32)
  37 ; IR-NEXT:  [[REG0:%.+]] = getelementptr float, float* %B, i64 8
  38 ; IR-NEXT:  [[REG1:%.+]] = bitcast float* [[REG0]] to i8*
  39 ; IR-NEXT:  call void @polly_copyFromHostToDevice(i8* [[REG1]], i8* %p_dev_array_MemRef_B, i64 32)
  40
  41 ; IR:      [[REGA:%.+]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_B)
  42 ; IR-NEXT: [[REGB:%.+]]  = bitcast i8* [[REGA]] to float*
  43 ; IR-NEXT: [[REGC:%.+]]  = getelementptr float, float* [[REGB]], i64 -8
  44 ; IR-NEXT: [[REGD:%.+]]  = bitcast float* [[REGC]] to i8*
  45
  46 ;    void foo(float A[], float B[]) {
  47 ;      for (long i = 0; i < 8; i++)
  48 ;        B[i + 8] *= 4;
  49 ;
  50 ;      for (long i = 0; i < 8; i++)
  51 ;        A[i] *= 12;
  52 ;    }
  53 ;
  54 ;    #ifdef OUTPUT
  55 ;    int main() {
  56 ;      float A[16];
  57 ;
  58 ;      for (long i = 0; i < 16; i++) {
  59 ;        __sync_synchronize();
  60 ;        A[i] = i;
  61 ;      }
  62 ;
  63 ;      foo(A, A);
  64 ;
  65 ;      float sum = 0;
  66 ;      for (long i = 0; i < 16; i++) {
  67 ;        __sync_synchronize();
  68 ;        sum += A[i];
  69 ;      }
  70 ;
  71 ;      printf("%f\n", sum);
  72 ;    }
  73 ;    #endif
  74 ;
  75 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
  76
  77 define void @foo(float* %A, float* %B) {
  78 bb:
  79   br label %bb2
  80
  81 bb2:                                              ; preds = %bb7, %bb
  82   %i.0 = phi i64 [ 0, %bb ], [ %tmp8, %bb7 ]
  83   %exitcond1 = icmp ne i64 %i.0, 8
  84   br i1 %exitcond1, label %bb3, label %bb9
  85
  86 bb3:                                              ; preds = %bb2
  87   %tmp = add nuw nsw i64 %i.0, 8
  88   %tmp4 = getelementptr inbounds float, float* %B, i64 %tmp
  89   %tmp5 = load float, float* %tmp4, align 4
  90   %tmp6 = fmul float %tmp5, 4.000000e+00
  91   store float %tmp6, float* %tmp4, align 4
  92   br label %bb7
  93
  94 bb7:                                              ; preds = %bb3
  95   %tmp8 = add nuw nsw i64 %i.0, 1
  96   br label %bb2
  97
  98 bb9:                                              ; preds = %bb2
  99   br label %bb10
 100
 101 bb10:                                             ; preds = %bb15, %bb9
 102   %i1.0 = phi i64 [ 0, %bb9 ], [ %tmp16, %bb15 ]
 103   %exitcond = icmp ne i64 %i1.0, 8
 104   br i1 %exitcond, label %bb11, label %bb17
 105
 106 bb11:                                             ; preds = %bb10
 107   %tmp12 = getelementptr inbounds float, float* %A, i64 %i1.0
 108   %tmp13 = load float, float* %tmp12, align 4
 109   %tmp14 = fmul float %tmp13, 1.200000e+01
 110   store float %tmp14, float* %tmp12, align 4
 111   br label %bb15
 112
 113 bb15:                                             ; preds = %bb11
 114   %tmp16 = add nuw nsw i64 %i1.0, 1
 115   br label %bb10
 116
 117 bb17:                                             ; preds = %bb10
 118   ret void
 119 }