test/GPGPU/region-stmt.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 ; CODE: Code
   9 ; CODE-NEXT: ====
  10 ; CODE-NEXT: # host
  11 ; CODE-NEXT: {
  12 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (128) * sizeof(float), cudaMemcpyHostToDevice));
  13 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (128) * sizeof(float), cudaMemcpyHostToDevice));
  14 ; CODE-NEXT:   {
  15 ; CODE-NEXT:     dim3 k0_dimBlock(32);
  16 ; CODE-NEXT:     dim3 k0_dimGrid(4);
  17 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A, dev_MemRef_B);
  18 ; CODE-NEXT:     cudaCheckKernel();
  19 ; CODE-NEXT:   }
  20
  21 ; CODE:   cudaCheckReturn(cudaMemcpy(MemRef_B, dev_MemRef_B, (128) * sizeof(float), cudaMemcpyDeviceToHost));
  22 ; CODE-NEXT: }
  23
  24 ; CODE: # kernel0
  25 ; CODE-NEXT: Stmt_for_body__TO__if_end(32 * b0 + t0);
  26
  27 ; IR: @polly_initContext
  28
  29 ; KERNEL-IR: kernel_0
  30
  31 ; REQUIRES: pollyacc
  32
  33 ;    void foo(float A[], float B[]) {
  34 ;      for (long i = 0; i < 128; i++)
  35 ;        if (A[i] == 42)
  36 ;          B[i] += 2 * i;
  37 ;        else
  38 ;          B[i] += 4 * i;
  39 ;    }
  40 ;
  41 source_filename = "/tmp/test.c"
  42 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
  43
  44 define void @foo(float* %A, float* %B) {
  45 entry:
  46   br label %for.cond
  47
  48 for.cond:                                         ; preds = %for.inc, %entry
  49   %i.0 = phi i64 [ 0, %entry ], [ %inc, %for.inc ]
  50   %exitcond = icmp ne i64 %i.0, 128
  51   br i1 %exitcond, label %for.body, label %for.end
  52
  53 for.body:                                         ; preds = %for.cond
  54   %arrayidx = getelementptr inbounds float, float* %A, i64 %i.0
  55   %tmp = load float, float* %arrayidx, align 4
  56   %cmp1 = fcmp oeq float %tmp, 4.200000e+01
  57   br i1 %cmp1, label %if.then, label %if.else
  58
  59 if.then:                                          ; preds = %for.body
  60   %mul = shl nsw i64 %i.0, 1
  61   %conv = sitofp i64 %mul to float
  62   %arrayidx2 = getelementptr inbounds float, float* %B, i64 %i.0
  63   %tmp1 = load float, float* %arrayidx2, align 4
  64   %add = fadd float %tmp1, %conv
  65   store float %add, float* %arrayidx2, align 4
  66   br label %if.end
  67
  68 if.else:                                          ; preds = %for.body
  69   %mul3 = shl nsw i64 %i.0, 2
  70   %conv4 = sitofp i64 %mul3 to float
  71   %arrayidx5 = getelementptr inbounds float, float* %B, i64 %i.0
  72   %tmp2 = load float, float* %arrayidx5, align 4
  73   %add6 = fadd float %tmp2, %conv4
  74   store float %add6, float* %arrayidx5, align 4
  75   br label %if.end
  76
  77 if.end:                                           ; preds = %if.else, %if.then
  78   br label %for.inc
  79
  80 for.inc:                                          ; preds = %if.end
  81   %inc = add nuw nsw i64 %i.0, 1
  82   br label %for.cond
  83
  84 for.end:                                          ; preds = %for.cond
  85   ret void
  86 }