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