[Polly] [PPCGCodeGeneration] [WIP] Teach `must_kills` to kill scalars that are local the scop.

[bollu]

Summary

• This is useful for more complex examples where we have scalars that

are local to a scop.

• If the local is only used within a scop, we are free to kill it.
• WIP, phi node test seems to fail. Should see why.
• WIP, code is ugly.

[bollu]

• [WIP] ran code through checkpolly, separated out checking for scalars into a pure function.

[bollu]

• update phi-nodes-in-kernel with approximate C source and extra kills.
• All tests pass.

[grosser]

LGTM

[bollu]

• [NFC] Final fixups, reviewed offline. Pushing diff before committing.

[singam-sanjay]

What do you mean by killing a scalar ?

lib/CodeGen/PPCGCodeGeneration.cpp

struct MustKillsInfo {
   /// allow us to easily derive kill information. PHI nodes are not alive
   /// outside the scop and can consequently all be "killed". [params] -> {
   /// [Stmt_phantom[] -> ref_phantom[]] -> phi_ref[] }
   isl::union_map TaggedMustKills;
 
   MustKillsInfo() : KillsSchedule(nullptr), TaggedMustKills(nullptr){};
 };
 
+/// Check if ScopArrayInfo SAI's uses are entirely contained within Scop S
+static bool isScalarUsesContainedInScop(const Scop &S,
+                                        const ScopArrayInfo *SAI) {
+  assert(!SAI->isArrayKind() && "this function only deals with scalars."
+                                " Dealing with arrays required alias analysis");
+
+  const Region &R = S.getRegion();
+  for (User *U : SAI->getBasePtr()->users()) {
+    Instruction *I = dyn_cast<Instruction>(U);
+    assert(I && "invalid user");
+    if (!R.contains(I))
+      return false;
+  }
+  return true;
+}
+
 /// Compute must-kills needed to enable live range reordering with PPCG.
 ///
 /// @params S The Scop to compute live range reordering information
 /// @returns live range reordering information that can be used to setup
 /// PPCG.
 static MustKillsInfo computeMustKillsInfo(const Scop &S) {
   const isl::space ParamSpace(isl::manage(S.getParamSpace()));
   MustKillsInfo Info;
 
-  // 1. Collect phi nodes in scop.
+  // 1. Collect all ScopArrayInfo that satisfy *any* of the criteria:
+  //      1.1 phi nodes in scop.
+  //      1.2 scalars that are only used within the scop
   SmallVector<isl::id, 4> KillMemIds;
   for (ScopArrayInfo *SAI : S.arrays()) {
-    if (!SAI->isPHIKind())
-      continue;
-
+    if (SAI->isPHIKind() ||
+        (!SAI->isArrayKind() && isScalarUsesContainedInScop(S, SAI)))
       KillMemIds.push_back(isl::manage(SAI->getBasePtrId()));
   }
 
   Info.TaggedMustKills = isl::union_map::empty(isl::space(ParamSpace));
 
   // Initialising KillsSchedule to `isl_set_empty` creates an empty node in the
   // schedule:
   //     - filter: "[control] -> { }"
   // So, we choose to not create this to keep the output a little nicer,

test/GPGPU/add-scalars-in-scop-to-kills.ll

+; RUN: opt %loadPolly -analyze -polly-scops < %s | FileCheck %s -check-prefix=SCOP
+; RUN: opt %loadPolly -S -polly-codegen-ppcg < %s | FileCheck %s -check-prefix=HOST-IR
+
+; REQUIRES: pollyacc
+
+; Check that we detect a scop.
+; SCOP:       Function: checkScalarKill
+; SCOP-NEXT: Region: %XLoopInit---%for.end
+; SCOP-NEXT: Max Loop Depth:  1
+
+; Check that we have a scalar that is not a phi node in the scop.
+; i32 MemRef_x_0; // Element size 4
+
+; Check that kernel launch is generated in host IR.
+; the declare would not be generated unless a call to a kernel exists.
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+
+; Check that we add variables that are local to a scop into the kills that we
+; pass to PPCG. This should enable PPCG to codegen this example.
+; void checkScalarKill(int A[], int B[], int C[], const int control1, int control2) {
+; int x;
+; #pragma scop
+;     for(int i = 0; i < 1000; i++) {
+; XLoopInit:        x = 0;
+; 
+;         if (control1 > 2)
+;             C1Add: x += 10;
+;         if (control2 > 3)
+;             C2Add: x += A[i];
+; 
+; BLoopAccumX:        B[i] += x;
+;     }
+; 
+; #pragma endscop
+; }
+; ModuleID = 'test.ll'
+target datalayout = "e-m:o-i64:64-f80:128-n8:16:32:64-S128"
+
+define void @checkScalarKill(i32* %A, i32* %B, i32* %C, i32 %control1, i32 %control2) {
+entry:
+  br label %entry.split
+
+entry.split:                                      ; preds = %entry
+  br label %XLoopInit
+
+XLoopInit:                                        ; preds = %entry.split, %BLoopAccumX
+  %indvars.iv = phi i64 [ 0, %entry.split ], [ %indvars.iv.next, %BLoopAccumX ]
+  %cmp1 = icmp sgt i32 %control1, 2
+  %x.0 = select i1 %cmp1, i32 10, i32 0
+  %cmp2 = icmp sgt i32 %control2, 3
+  br i1 %cmp2, label %C2Add, label %BLoopAccumX
+
+C2Add:                                            ; preds = %XLoopInit
+  %arrayidx = getelementptr inbounds i32, i32* %A, i64 %indvars.iv
+  %tmp6 = load i32, i32* %arrayidx, align 4
+  %add4 = add nsw i32 %tmp6, %x.0
+  br label %BLoopAccumX
+
+BLoopAccumX:                                      ; preds = %XLoopInit, %C2Add
+  %x.1 = phi i32 [ %add4, %C2Add ], [ %x.0, %XLoopInit ]
+  %arrayidx7 = getelementptr inbounds i32, i32* %B, i64 %indvars.iv
+  %tmp11 = load i32, i32* %arrayidx7, align 4
+  %add8 = add nsw i32 %tmp11, %x.1
+  store i32 %add8, i32* %arrayidx7, align 4
+  %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
+  %exitcond = icmp ne i64 %indvars.iv.next, 1000
+  br i1 %exitcond, label %XLoopInit, label %for.end
+
+for.end:                                          ; preds = %BLoopAccumX
+  ret void
+}

test/GPGPU/phi-nodes-in-kernel.ll

 ; 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 -S < %s | \
 ; RUN: FileCheck %s -check-prefix=IR
 
 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-kernel-ir \
 ; RUN: -disable-output < %s | \
 ; RUN: FileCheck %s -check-prefix=KERNEL-IR
 
 ; REQUIRES: pollyacc
 
+; Approximate C source:
+; void kernel_dynprog(int c[50]) {
+;     int iter = 0;
+;     int outl = 0;
+;
+;      while(1) {
+;         for(int indvar = 1 ; indvar <= 49; indvar++) {
+;             c[indvar] = undef;
+;         }
+;         add78 = c[49] + outl;
+;         inc80 = iter + 1;
+;
+;         if (true) break;
+;
+;         outl = add78;
+;         iter = inc80;
+;      }
+;}
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
-; CODE: # host
+; CODE:      # host
 ; CODE-NEXT: {
 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(dev_MemRef_out_l_055__phi, &MemRef_out_l_055__phi, sizeof(i32), cudaMemcpyHostToDevice));
 ; CODE-NEXT:   {
 ; CODE-NEXT:     dim3 k0_dimBlock(32);
 ; CODE-NEXT:     dim3 k0_dimGrid(2);
 ; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_out_l_055__phi, dev_MemRef_out_l_055, dev_MemRef_c);
 ; CODE-NEXT:     cudaCheckKernel();
 ; CODE-NEXT:   }
 
-; CODE:       cudaCheckReturn(cudaMemcpy(&MemRef_out_l_055, dev_MemRef_out_l_055, sizeof(i32), cudaMemcpyDeviceToHost));
-; CODE-NEXT:  cudaCheckReturn(cudaMemcpy(MemRef_c, dev_MemRef_c, (50) * sizeof(i32), cudaMemcpyDeviceToHost));
+; CODE:        cudaCheckReturn(cudaMemcpy(MemRef_c, dev_MemRef_c, (50) * sizeof(i32), cudaMemcpyDeviceToHost));
 ; CODE-NEXT: }
 
 ; CODE: # kernel0
 ; CODE-NEXT: if (32 * b0 + t0 <= 48) {
 ; CODE-NEXT:   if (b0 == 1 && t0 == 16)
 ; CODE-NEXT:     Stmt_for_cond1_preheader(0);
 ; CODE-NEXT:   Stmt_for_body17(0, 32 * b0 + t0);
 ; CODE-NEXT:   if (b0 == 1 && t0 == 16)
 ; CODE-NEXT:     Stmt_for_cond15_for_cond12_loopexit_crit_edge(0);
 ; CODE-NEXT: }
 
 ; IR:      [[REGA:%.+]] = bitcast i32* %out_l.055.phiops to i8*
 ; IR-NEXT: call void @polly_copyFromHostToDevice(i8* [[REGA]], i8* %p_dev_array_MemRef_out_l_055__phi, i64 4)
 
-; IR: [[REGC:%.+]] = bitcast i32* %out_l.055.s2a to i8*
-; IR-NEXT: call void @polly_copyFromDeviceToHost(i8* %p_dev_array_MemRef_out_l_055, i8* [[REGC]], i64 4)
+; %37 = bitcast [50 x i32]* %c to i32*
+; %38 = getelementptr i32, i32* %37, i64 1 
+; IR: [[REGC:%.+]] =   bitcast i32* %38 to i8*
+; IR-NEXT:  call void @polly_copyFromDeviceToHost(i8* %p_dev_array_MemRef_c, i8* [[REGC]], i64 196)
 
 ; KERNEL-IR: entry:
 ; KERNEL-IR-NEXT:   %out_l.055.s2a = alloca i32
 ; KERNEL-IR-NEXT:   %out_l.055.phiops = alloca i32
 ; KERNEL-IR-NEXT:   %1 = addrspacecast i8 addrspace(1)* %MemRef_out_l_055__phi to i32*
 ; KERNEL-IR-NEXT:   %2 = load i32, i32* %1
 ; KERNEL-IR-NEXT:   store i32 %2, i32* %out_l.055.phiops
 ; KERNEL-IR-NEXT:   %3 = addrspacecast i8 addrspace(1)* %MemRef_out_l_055 to i32*