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