blob: 022fa920adb672a36ab030153d014862f9108a0a [file] [log] [blame]
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
; RUN: -disable-output < %s | \
; RUN: FileCheck -check-prefix=CODE %s
; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \
; RUN: -disable-output < %s | \
; RUN: FileCheck %s -check-prefix=KERNEL-IR
;
; REQUIRES: pollyacc
;
; #include <stdio.h>
;
; float foo(float A[]) {
; float sum = 0;
;
; for (long i = 0; i < 32; i++)
; A[i] = i;
;
; for (long i = 0; i < 32; i++)
; A[i] += i;
;
; for (long i = 0; i < 32; i++)
; sum += A[i];
;
; return sum;
; }
;
; int main() {
; float A[32];
; float sum = foo(A);
; printf("%f\n", sum);
; }
; CODE: dim3 k0_dimBlock(32);
; CODE-NEXT: dim3 k0_dimGrid(1);
; CODE-NEXT: kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: {
; CODE-NEXT: dim3 k1_dimBlock;
; CODE-NEXT: dim3 k1_dimGrid;
; CODE-NEXT: kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_sum_0__phi);
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: {
; CODE-NEXT: dim3 k2_dimBlock;
; CODE-NEXT: dim3 k2_dimGrid;
; CODE-NEXT: kernel2 <<<k2_dimGrid, k2_dimBlock>>> (dev_MemRef_A, dev_MemRef_sum_0__phi, dev_MemRef_sum_0);
; CODE-NEXT: cudaCheckKernel();
; CODE-NEXT: }
; CODE: cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (32) * sizeof(float), cudaMemcpyDeviceToHost));
; CODE-NEXT: cudaCheckReturn(cudaMemcpy(&MemRef_sum_0, dev_MemRef_sum_0, sizeof(float), cudaMemcpyDeviceToHost));
; CODE-NEXT: cudaCheckReturn(cudaFree(dev_MemRef_A));
; CODE-NEXT: cudaCheckReturn(cudaFree(dev_MemRef_sum_0__phi));
; CODE-NEXT: cudaCheckReturn(cudaFree(dev_MemRef_sum_0));
; CODE-NEXT: }
; CODE: # kernel0
; CODE-NEXT: {
; CODE-NEXT: Stmt_bb4(t0);
; CODE-NEXT: Stmt_bb10(t0);
; CODE-NEXT: }
; CODE: # kernel1
; CODE-NEXT: Stmt_bb17();
; CODE: # kernel2
; CODE_NEXT: {
; CODE_NEXT: read();
; CODE_NEXT: for (int c0 = 0; c0 <= 32; c0 += 1) {
; CODE_NEXT: Stmt_bb18(c0);
; CODE_NEXT: if (c0 <= 31)
; CODE_NEXT: Stmt_bb20(c0);
; CODE_NEXT: }
; CODE_NEXT: write();
; CODE_NEXT: }
; KERNEL-IR: define ptx_kernel void @FUNC_foo_SCOP_0_KERNEL_1(i8 addrspace(1)* %MemRef_sum_0__phi)
; KERNEL-IR: store float 0.000000e+00, float* %sum.0.phiops
; KERNEL-IR: [[REGA:%.+]] = addrspacecast i8 addrspace(1)* %MemRef_sum_0__phi to float*
; KERNEL-IR: [[REGB:%.+]] = load float, float* %sum.0.phiops
; KERNEL-IR: store float [[REGB]], float* [[REGA]]
; KERNEL-IR: define ptx_kernel void @FUNC_foo_SCOP_0_KERNEL_2(i8 addrspace(1)* %MemRef_A, i8 addrspace(1)* %MemRef_sum_0__phi, i8 addrspace(1)* %MemRef_sum_0)
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
@.str = private unnamed_addr constant [4 x i8] c"%f\0A\00", align 1
define float @foo(float* %A) {
bb:
br label %bb3
bb3: ; preds = %bb6, %bb
%i.0 = phi i64 [ 0, %bb ], [ %tmp7, %bb6 ]
%exitcond2 = icmp ne i64 %i.0, 32
br i1 %exitcond2, label %bb4, label %bb8
bb4: ; preds = %bb3
%tmp = sitofp i64 %i.0 to float
%tmp5 = getelementptr inbounds float, float* %A, i64 %i.0
store float %tmp, float* %tmp5, align 4
br label %bb6
bb6: ; preds = %bb4
%tmp7 = add nuw nsw i64 %i.0, 1
br label %bb3
bb8: ; preds = %bb3
br label %bb9
bb9: ; preds = %bb15, %bb8
%i1.0 = phi i64 [ 0, %bb8 ], [ %tmp16, %bb15 ]
%exitcond1 = icmp ne i64 %i1.0, 32
br i1 %exitcond1, label %bb10, label %bb17
bb10: ; preds = %bb9
%tmp11 = sitofp i64 %i1.0 to float
%tmp12 = getelementptr inbounds float, float* %A, i64 %i1.0
%tmp13 = load float, float* %tmp12, align 4
%tmp14 = fadd float %tmp13, %tmp11
store float %tmp14, float* %tmp12, align 4
br label %bb15
bb15: ; preds = %bb10
%tmp16 = add nuw nsw i64 %i1.0, 1
br label %bb9
bb17: ; preds = %bb9
br label %bb18
bb18: ; preds = %bb20, %bb17
%sum.0 = phi float [ 0.000000e+00, %bb17 ], [ %tmp23, %bb20 ]
%i2.0 = phi i64 [ 0, %bb17 ], [ %tmp24, %bb20 ]
%exitcond = icmp ne i64 %i2.0, 32
br i1 %exitcond, label %bb19, label %bb25
bb19: ; preds = %bb18
br label %bb20
bb20: ; preds = %bb19
%tmp21 = getelementptr inbounds float, float* %A, i64 %i2.0
%tmp22 = load float, float* %tmp21, align 4
%tmp23 = fadd float %sum.0, %tmp22
%tmp24 = add nuw nsw i64 %i2.0, 1
br label %bb18
bb25: ; preds = %bb18
%sum.0.lcssa = phi float [ %sum.0, %bb18 ]
ret float %sum.0.lcssa
}
define i32 @main() {
bb:
%A = alloca [32 x float], align 16
%tmp = getelementptr inbounds [32 x float], [32 x float]* %A, i64 0, i64 0
%tmp1 = call float @foo(float* %tmp)
%tmp2 = fpext float %tmp1 to double
%tmp3 = call i32 (i8*, ...) @printf(i8* getelementptr inbounds ([4 x i8], [4 x i8]* @.str, i64 0, i64 0), double %tmp2) #2
ret i32 0
}
declare i32 @printf(i8*, ...) #1