[Polly] [PPCGCodeGeneration] [WIP] Do not consider writes to killable scalars as scalar stores

[bollu]

Summary

• If a scalar can be killed, then the order of writes into it does not matter, since we don't read it later anyway. So, do not track these as scalar stores

• TODO 1. How do I test this? there's no direct way to see if this works or not. Perhaps I could observe that there is no br i1 false at the RTC check. However, that is too indirect for my taste.

• TODO 2. Refactor the SAI check into a separate function, since it is used in two places now.

[bollu]

@grosser - does the change make sense to you? I'll clean it once I know that this is a sane change.

[grosser]

Thanks for pushing forward here. One thing I would like to check if the scalars that are private are passed into the kernel or not. AFAIU ppcg is not passing such "private" scalars to the kernel, but Polly-ACC still does this today. Can you check if/what is going on here?

Best,
Tobias

[bollu]

From what I can tell, PPCG does seem to pass "private" scalars:

test.c

void checkPrivatization(int A[], int B[], int C[], const int control) {
int x = 0;
#pragma scop
    for(int i = 0; i < 1000; i++) {

ifcond:        if (control != 0)
              xinc: x = C[i];

xuse:        B[i] = x * A[i];
    }
#pragma endscop
}

testbed_kernel.cu

/* PPCG output */
#include "testbed_kernel.hu"
__global__ void kernel0(int *A, int *B, int *C, int *x, int control)
{
    int b0 = blockIdx.x;
    int t0 = threadIdx.x;
    __shared__ int shared_x;

    {
      if (control == 0 && t0 == 0)
        shared_x = *x;
      __syncthreads();
      if ((control <= -1 && 32 * b0 + t0 <= 999) || (control >= 1 && 32 * b0 + t0 <= 999))
        shared_x = C[32 * b0 + t0];
      if (32 * b0 + t0 <= 999)
        B[32 * b0 + t0] = (shared_x * A[32 * b0 + t0]);
      __syncthreads();
    }
}

I would have expected it to not take x as a parameter, but it seems to do so.
Are you referring to some other behaviour I am unaware of?

[bollu]

• Cleanup code.

[bollu]

@grosser - with this change, we break the test test/GPGPU/invariant-load-hoisting-with-failing-scop.ll which gets triggered when we have a BuildSuccessful = 0 with invariant load hoisting. However, the only way to trigger this is either

1. screw up verifyModule

or

2. have a test case where storedScalar is true.

I'm not able to generate an example after this patch such that storedScalar = true.

Any ideas for this?

[grosser]

Seems inconsistent, I get for the following input:

float X;                                                                         
void f(float A[]) {                                                              
#pragma scop                                                                     
  for (long i = 0; i < 1024; i++) {                                              
    X = i * i;                                                                   
    A[i] = X;                                                                    
  }                                                                              
 X = 100;                                                                        
#pragma endscop                                                                  
}

the following output

__global__ void kernel0(float *A)
{
    int b0 = blockIdx.x;
    int t0 = threadIdx.x;
    float private_X;

    {
      private_X = ((32 * b0 + t0) * (32 * b0 + t0));
      A[32 * b0 + t0] = private_X;
    }
}
__global__ void kernel1(float *X)
{

    X[0] = 100;
}

[bollu]

• [WIP] previous commmit makes no sense, that checks if we care about the value at the *end* of a scop, not in between loop iterations

[bollu]

Abandoned because rL311259 contains changes that this patch was supposed to perform.

lib/CodeGen/PPCGCodeGeneration.cpp

for (User *U : SAI->getBasePtr()->users()) {
     Instruction *I = dyn_cast<Instruction>(U);
     assert(I && "invalid user of scop array info");
     if (!R.contains(I))
       return false;
   }
   return true;
 }
 
+static bool canKillSAI(const Scop &S, const ScopArrayInfo *SAI) {
+  return (SAI->isPHIKind() ||
+          (SAI->isValueKind() && isScalarUsesContainedInScop(S, SAI)));
+}
+
 /// 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 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() ||
-        (SAI->isValueKind() && isScalarUsesContainedInScop(S, SAI)))
+    if (canKillSAI(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] -> { }"

test/GPGPU/allow-scalar-stores.ll

+; RUN: opt %loadPolly -analyze -polly-scops < %s | FileCheck %s -check-prefix=SCOP
+; RUN: opt %loadPolly -S -polly-codegen-ppcg -polly-acc-fail-on-verify-module-failure < %s | FileCheck %s -check-prefix=HOST-IR
+
+; REQUIRES: pollyacc
+
+; SCOP:      Function: f
+; SCOP-NEXT: Region: %for.body---%for.end
+; SCOP-NEXT: Max Loop Depth:  1
+
+; 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*)
+
+; void f(int A[], int B[], int control, int C[]) {
+;     int x;
+; #pragma scop
+;     for(int i = 0; i < 1000; i ++) {
+;         x = 0;
+;         if(control) x = C[i];
+;         B[i] = x * A[i];
+;
+;     }
+; #pragma endscop
+; }
+
+target datalayout = "e-m:o-i64:64-f80:128-n8:16:32:64-S128"
+
+define void @f(i32* %A, i32* %B, i32 %control, i32* %C) {
+entry:
+  br label %entry.split
+
+entry.split:                                      ; preds = %entry
+  br label %for.body
+
+for.body:                                         ; preds = %entry.split, %if.end
+  %indvars.iv = phi i64 [ 0, %entry.split ], [ %indvars.iv.next, %if.end ]
+  %tobool = icmp eq i32 %control, 0
+  br i1 %tobool, label %if.end, label %if.then
+
+if.then:                                          ; preds = %for.body
+  %arrayidx = getelementptr inbounds i32, i32* %C, i64 %indvars.iv
+  %tmp4 = load i32, i32* %arrayidx, align 4
+  br label %if.end
+
+if.end:                                           ; preds = %for.body, %if.then
+  %x.0 = phi i32 [ %tmp4, %if.then ], [ 0, %for.body ]
+  %arrayidx2 = getelementptr inbounds i32, i32* %A, i64 %indvars.iv
+  %tmp8 = load i32, i32* %arrayidx2, align 4
+  %mul = mul nsw i32 %tmp8, %x.0
+  %arrayidx4 = getelementptr inbounds i32, i32* %B, i64 %indvars.iv
+  store i32 %mul, i32* %arrayidx4, align 4
+  %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
+  %exitcond = icmp ne i64 %indvars.iv.next, 1000
+  br i1 %exitcond, label %for.body, label %for.end
+
+for.end:                                          ; preds = %if.end
+  ret void
+}