test/GPGPU/double-parallel-loop.ll

   1 ; RUN: opt %loadPolly -polly-scops -analyze < %s | FileCheck %s
   2 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-schedule \
   3 ; RUN: -disable-output < %s | \
   4 ; RUN: FileCheck -check-prefix=SCHED %s
   5
   6 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
   7 ; RUN: -disable-output < %s | \
   8 ; RUN: FileCheck -check-prefix=CODE %s
   9
  10 ; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \
  11 ; RUN: FileCheck %s -check-prefix=IR
  12
  13 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \
  14 ; RUN: -disable-output < %s | \
  15 ; RUN: FileCheck %s -check-prefix=KERNEL-IR
  16
  17 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-asm \
  18 ; RUN: -disable-output < %s | \
  19 ; RUN: FileCheck %s -check-prefix=KERNEL-ASM
  20
  21 ; REQUIRES: pollyacc
  22
  23 ; CHECK: Stmt_bb5
  24 ; CHECK-NEXT:       Domain :=
  25 ; CHECK-NEXT:           { Stmt_bb5[i0, i1] : 0 <= i0 <= 1023 and 0 <= i1 <= 1023 };
  26 ; CHECK-NEXT:       Schedule :=
  27 ; CHECK-NEXT:           { Stmt_bb5[i0, i1] -> [i0, i1] };
  28 ; CHECK-NEXT:       ReadAccess :=       [Reduction Type: NONE] [Scalar: 0]
  29 ; CHECK-NEXT:           { Stmt_bb5[i0, i1] -> MemRef_A[i0, i1] };
  30 ; CHECK-NEXT:       MustWriteAccess :=  [Reduction Type: NONE] [Scalar: 0]
  31 ; CHECK-NEXT:           { Stmt_bb5[i0, i1] -> MemRef_A[i0, i1] };
  32
  33 ; SCHED: domain: "{ Stmt_bb5[i0, i1] : 0 <= i0 <= 1023 and 0 <= i1 <= 1023 }"
  34 ; SCHED-NEXT: child:
  35 ; SCHED-NEXT:   context: "{ [] }"
  36 ; SCHED-NEXT:   child:
  37 ; SCHED-NEXT:     extension: "{ [] -> from_device_MemRef_A[]; [] -> to_device_MemRef_A[] }"
  38 ; SCHED-NEXT:     child:
  39 ; SCHED-NEXT:       sequence:
  40 ; SCHED-NEXT:       - filter: "{ to_device_MemRef_A[] }"
  41 ; SCHED-NEXT:         child:
  42 ; SCHED-NEXT:           set:
  43 ; SCHED-NEXT:           - filter: "{ to_device_MemRef_A[] }"
  44 ; SCHED-NEXT:             child:
  45 ; SCHED-NEXT:               guard: "{ [] }"
  46 ; SCHED-NEXT:       - filter: "{ Stmt_bb5[i0, i1] }"
  47 ; SCHED-NEXT:         child:
  48 ; SCHED-NEXT:           guard: "{ [] }"
  49 ; SCHED-NEXT:           child:
  50 ; SCHED-NEXT:             mark: "kernel"
  51 ; SCHED-NEXT:             child:
  52 ; SCHED-NEXT:               context: "[b0, b1, t0, t1] -> { [] : 0 <= b0 <= 31 and 0 <= b1 <= 31 and 0 <= t0 <= 31 and 0 <= t1 <= 15 }"
  53 ; SCHED-NEXT:               child:
  54 ; SCHED-NEXT:                 filter: "[b0, b1] -> { Stmt_bb5[i0, i1] : -31 - 32b0 + i0 <= 8192*floor((i0)/8192) <= -32b0 + i0 and -31 - 32b1 + i1 <= 8192*floor((i1)/8192) <= -32b1 + i1 }"
  55 ; SCHED-NEXT:                 child:
  56 ; SCHED-NEXT:                   schedule: "[{ Stmt_bb5[i0, i1] -> [(floor((i0)/8192))] }, { Stmt_bb5[i0, i1] -> [(floor((i1)/8192))] }]"
  57 ; SCHED-NEXT:                   permutable: 1
  58 ; SCHED-NEXT:                   coincident: [ 1, 1 ]
  59 ; SCHED-NEXT:                   child:
  60 ; SCHED-NEXT:                     filter: "[t0, t1] -> { Stmt_bb5[i0, i1] : 32*floor((-t0 + i0)/32) = -t0 + i0 and 16*floor((-t1 + i1)/16) = -t1 + i1 and 0 <= t0 <= 31 and 0 <= t1 <= 15 }"
  61 ; SCHED-NEXT:                     child:
  62 ; SCHED-NEXT:                       schedule: "[{ Stmt_bb5[i0, i1] -> [(0)] }, { Stmt_bb5[i0, i1] -> [(floor((i1)/16) - 2*floor((i1)/32))] }]"
  63 ; SCHED-NEXT:                       permutable: 1
  64 ; SCHED-NEXT:                       coincident: [ 1, 1 ]
  65 ; SCHED-NEXT:       - filter: "{ from_device_MemRef_A[] }"
  66 ; SCHED-NEXT:         child:
  67 ; SCHED-NEXT:           set:
  68 ; SCHED-NEXT:           - filter: "{ from_device_MemRef_A[] }"
  69 ; SCHED-NEXT:             child:
  70 ; SCHED-NEXT:               guard: "{ [] }"
  71
  72 ; CODE: Code
  73 ; CODE-NEXT: ====
  74 ; CODE-NEXT: # host
  75 ; CODE-NEXT: {
  76 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (1024) * (1024) * sizeof(float), cudaMemcpyHostToDevice));
  77 ; CODE-NEXT:   {
  78 ; CODE-NEXT:     dim3 k0_dimBlock(16, 32);
  79 ; CODE-NEXT:     dim3 k0_dimGrid(32, 32);
  80 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
  81 ; CODE-NEXT:     cudaCheckKernel();
  82 ; CODE-NEXT:   }
  83
  84 ; CODE:   cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (1024) * (1024) * sizeof(float), cudaMemcpyDeviceToHost));
  85 ; CODE-NEXT: }
  86
  87 ; CODE: # kernel0
  88 ; CODE-NEXT: for (int c3 = 0; c3 <= 1; c3 += 1)
  89 ; CODE-NEXT:   Stmt_bb5(32 * b0 + t0, 32 * b1 + t1 + 16 * c3);
  90
  91 ; IR: polly.split_new_and_old:
  92 ; IR-NEXT:    br i1 true, label %polly.start, label %bb2
  93
  94 ; IR: polly.start:
  95 ; IR-NEXT:    [[GPUContext:%.*]] = call i8* @polly_initContext()
  96 ; IR-NEXT:    %p_dev_array_MemRef_A = call i8* @polly_allocateMemoryForDevice(i64 4194304)
  97 ; IR-NEXT:    call void @polly_freeDeviceMemory(i8* %p_dev_array_MemRef_A)
  98 ; IR-NEXT:    call void @polly_freeContext(i8* [[GPUContext]])
  99 ; IR-NEXT:    br label %polly.exiting
 100
 101 ; IR: polly.exiting:
 102 ; IR-NEXT:    br label %polly.merge_new_and_old
 103
 104 ; KERNEL-IR-LABEL: define ptx_kernel void @kernel_0(i8* %MemRef_A) {
 105 ; KERNEL-IR-NEXT: entry:
 106 ; KERNEL-IR-NEXT:   %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
 107 ; KERNEL-IR-NEXT:   %b0 = zext i32 %0 to i64
 108 ; KERNEL-IR-NEXT:   %1 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.y()
 109 ; KERNEL-IR-NEXT:   %b1 = zext i32 %1 to i64
 110 ; KERNEL-IR-NEXT:   %2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 111 ; KERNEL-IR-NEXT:   %t0 = zext i32 %2 to i64
 112 ; KERNEL-IR-NEXT:   %3 = call i32 @llvm.nvvm.read.ptx.sreg.tid.y()
 113 ; KERNEL-IR-NEXT:   %t1 = zext i32 %3 to i64
 114 ; KERNEL-IR-NEXT:   br label %polly.loop_preheader
 115
 116 ; KERNEL-IR-LABEL: polly.loop_exit:                                  ; preds = %polly.stmt.bb5
 117 ; KERNEL-IR-NEXT:   ret void
 118
 119 ; KERNEL-IR-LABEL: polly.loop_header:                                ; preds = %polly.stmt.bb5, %polly.loop_preheader
 120 ; KERNEL-IR-NEXT:   %polly.indvar = phi i64 [ 0, %polly.loop_preheader ], [ %polly.indvar_next, %polly.stmt.bb5 ]
 121 ; KERNEL-IR-NEXT:   %4 = mul nsw i64 32, %b0
 122 ; KERNEL-IR-NEXT:   %5 = add nsw i64 %4, %t0
 123 ; KERNEL-IR-NEXT:   %6 = mul nsw i64 32, %b1
 124 ; KERNEL-IR-NEXT:   %7 = add nsw i64 %6, %t1
 125 ; KERNEL-IR-NEXT:   %8 = mul nsw i64 16, %polly.indvar
 126 ; KERNEL-IR-NEXT:   %9 = add nsw i64 %7, %8
 127 ; KERNEL-IR-NEXT:   br label %polly.stmt.bb5
 128
 129 ; KERNEL-IR-LABEL: polly.stmt.bb5:                                   ; preds = %polly.loop_header
 130 ; KERNEL-IR-NEXT:   %10 = mul i64 %9, %5
 131 ; KERNEL-IR-NEXT:   %p_tmp6 = sitofp i64 %10 to float
 132 ; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A = bitcast i8* %MemRef_A to float*
 133 ; KERNEL-IR-NEXT:   %11 = mul nsw i64 32, %b0
 134 ; KERNEL-IR-NEXT:   %12 = add nsw i64 %11, %t0
 135 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A = mul nsw i64 %12, 1024
 136 ; KERNEL-IR-NEXT:   %13 = mul nsw i64 32, %b1
 137 ; KERNEL-IR-NEXT:   %14 = add nsw i64 %13, %t1
 138 ; KERNEL-IR-NEXT:   %15 = mul nsw i64 16, %polly.indvar
 139 ; KERNEL-IR-NEXT:   %16 = add nsw i64 %14, %15
 140 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A = add nsw i64 %polly.access.mul.MemRef_A, %16
 141 ; KERNEL-IR-NEXT:   %polly.access.MemRef_A = getelementptr float, float* %polly.access.cast.MemRef_A, i64 %polly.access.add.MemRef_A
 142 ; KERNEL-IR-NEXT:   %tmp8_p_scalar_ = load float, float* %polly.access.MemRef_A, align 4
 143 ; KERNEL-IR-NEXT:   %p_tmp9 = fadd float %tmp8_p_scalar_, %p_tmp6
 144 ; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A1 = bitcast i8* %MemRef_A to float*
 145 ; KERNEL-IR-NEXT:   %17 = mul nsw i64 32, %b0
 146 ; KERNEL-IR-NEXT:   %18 = add nsw i64 %17, %t0
 147 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A2 = mul nsw i64 %18, 1024
 148 ; KERNEL-IR-NEXT:   %19 = mul nsw i64 32, %b1
 149 ; KERNEL-IR-NEXT:   %20 = add nsw i64 %19, %t1
 150 ; KERNEL-IR-NEXT:   %21 = mul nsw i64 16, %polly.indvar
 151 ; KERNEL-IR-NEXT:   %22 = add nsw i64 %20, %21
 152 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A3 = add nsw i64 %polly.access.mul.MemRef_A2, %22
 153 ; KERNEL-IR-NEXT:   %polly.access.MemRef_A4 = getelementptr float, float* %polly.access.cast.MemRef_A1, i64 %polly.access.add.MemRef_A3
 154 ; KERNEL-IR-NEXT:   store float %p_tmp9, float* %polly.access.MemRef_A4, align 4
 155 ; KERNEL-IR-NEXT:   %polly.indvar_next = add nsw i64 %polly.indvar, 1
 156 ; KERNEL-IR-NEXT:   %polly.loop_cond = icmp sle i64 %polly.indvar, 0
 157 ; KERNEL-IR-NEXT:   br i1 %polly.loop_cond, label %polly.loop_header, label %polly.loop_exit
 158
 159 ; KERNEL-IR-LABEL: polly.loop_preheader:                             ; preds = %entry
 160 ; KERNEL-IR-NEXT:   br label %polly.loop_header
 161
 162
 163 ; KERNEL-ASM: .version 3.2
 164 ; KERNEL-ASM-NEXT: .target sm_30
 165 ; KERNEL-ASM-NEXT: .address_size 64
 166
 167 ; KERNEL-ASM:   // .globl     kernel_0
 168
 169 ; KERNEL-ASM: .visible .entry kernel_0(
 170 ; KERNEL-ASM-NEXT:   .param .u64 kernel_0_param_0
 171 ; KERNEL-ASM-NEXT: )
 172
 173 ;    void double_parallel_loop(float A[][1024]) {
 174 ;      for (long i = 0; i < 1024; i++)
 175 ;        for (long j = 0; j < 1024; j++)
 176 ;          A[i][j] += i * j;
 177 ;    }
 178 ;
 179 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 180
 181 define void @double_parallel_loop([1024 x float]* %A) {
 182 bb:
 183   br label %bb2
 184
 185 bb2:                                              ; preds = %bb13, %bb
 186   %i.0 = phi i64 [ 0, %bb ], [ %tmp14, %bb13 ]
 187   %exitcond1 = icmp ne i64 %i.0, 1024
 188   br i1 %exitcond1, label %bb3, label %bb15
 189
 190 bb3:                                              ; preds = %bb2
 191   br label %bb4
 192
 193 bb4:                                              ; preds = %bb10, %bb3
 194   %j.0 = phi i64 [ 0, %bb3 ], [ %tmp11, %bb10 ]
 195   %exitcond = icmp ne i64 %j.0, 1024
 196   br i1 %exitcond, label %bb5, label %bb12
 197
 198 bb5:                                              ; preds = %bb4
 199   %tmp = mul nuw nsw i64 %i.0, %j.0
 200   %tmp6 = sitofp i64 %tmp to float
 201   %tmp7 = getelementptr inbounds [1024 x float], [1024 x float]* %A, i64 %i.0, i64 %j.0
 202   %tmp8 = load float, float* %tmp7, align 4
 203   %tmp9 = fadd float %tmp8, %tmp6
 204   store float %tmp9, float* %tmp7, align 4
 205   br label %bb10
 206
 207 bb10:                                             ; preds = %bb5
 208   %tmp11 = add nuw nsw i64 %j.0, 1
 209   br label %bb4
 210
 211 bb12:                                             ; preds = %bb4
 212   br label %bb13
 213
 214 bb13:                                             ; preds = %bb12
 215   %tmp14 = add nuw nsw i64 %i.0, 1
 216   br label %bb2
 217
 218 bb15:                                             ; preds = %bb2
 219   ret void
 220 }