Limit PromoteToPredicatedMemoryAccess

Limite PromoteToPredicatedMemoryAccess to:
- simple load/store
- load/stores using only legal types
Early exit if found other instructions which read from memory or
calls which are convergent, since it is not safe to
remove branch in this case.

Author: YuriPlyakhin

IGC/Compiler/Legalizer/TypeLegalizer.h

@@ -1,6 +1,6 @@
 /*========================== begin_copyright_notice ============================
 
-Copyright (C) 2017-2021 Intel Corporation
+Copyright (C) 2017 Intel Corporation
 
 SPDX-License-Identifier: MIT
 
@@ -110,6 +110,21 @@ namespace IGC {
 
             bool runOnFunction(Function& F) override;
 
+            static bool isLegalInteger(unsigned width)
+            {
+                switch (width)
+                {
+                case 8:
+                case 16:
+                case 32:
+                case 64:
+                    return true;
+                default:
+                    break;
+                }
+                return false;
+            }
+
         private:
             void getAnalysisUsage(AnalysisUsage& AU) const override;
 
@@ -215,21 +230,6 @@ namespace IGC {
                 return Width;
             }
 
-            bool isLegalInteger(unsigned width) const
-            {
-                switch (width)
-                {
-                case 8:
-                case 16:
-                case 32:
-                case 64:
-                    return true;
-                default:
-                    break;
-                }
-                return false;
-            }
-
             /// getLargestLegalIntTypeSize() - Return the size of the largest legal
             /// integer type with size not bigger than Width bits.
             unsigned getLargestLegalIntTypeSize(unsigned Width) const {

IGC/Compiler/Optimizer/PromoteToPredicatedMemoryAccess.cpp

@@ -18,12 +18,14 @@ SPDX-License-Identifier: MIT
 #include "Compiler/CISACodeGen/OpenCLKernelCodeGen.hpp"
 #include "Compiler/CodeGenPublic.h"
 #include "Compiler/IGCPassSupport.h"
+#include "Compiler/Legalizer/TypeLegalizer.h"
 #include "Probe/Assertion.h"
 
 #define DEBUG_TYPE "PromoteToPredicatedMemoryAccess"
 
 using namespace llvm;
 using namespace IGC;
+using namespace IGC::Legalizer;
 
 static cl::opt<int> PredicatedMemOpIfConvertMaxInstrs(
     "igc-predmem-ifconv-max-instrs", cl::init(5), cl::Hidden,
@@ -108,6 +110,16 @@ void PromoteToPredicatedMemoryAccess::fixPhiNode(PHINode &Phi, BasicBlock &Prede
   Phi.eraseFromParent();
 }
 
+namespace {
+  bool IsTypeLegal(Type *Ty) {
+    if (!Ty->isIntOrIntVectorTy())
+      return true;
+
+    unsigned Width = Ty->getScalarSizeInBits();
+    return TypeLegalizer::isLegalInteger(Width) || Width == 1;
+  }
+} // namespace
+
 bool PromoteToPredicatedMemoryAccess::trySingleBlockIfConv(Value &Cond, BasicBlock &BranchBB,
                                                 BasicBlock &ConvBB, BasicBlock &SuccBB,
                                                 bool Inverse) {
@@ -141,26 +153,34 @@ bool PromoteToPredicatedMemoryAccess::trySingleBlockIfConv(Value &Cond, BasicBlo
     auto *Inst = dyn_cast<Instruction>(Phi.getIncomingValue(Idx));
     if (!Inst)
       return false;
-    if (!isa<LoadInst>(Inst))
+    if (auto *LI = dyn_cast<LoadInst>(Inst); !LI || !LI->isSimple())
+      return false;
+
+    if (!IsTypeLegal(Inst->getType()))
       return false;
 
     Insts[Inst] = Phi.getIncomingValueForBlock(&BranchBB);
   }
 
-  // Collect all the void
+  // Collect the rest of the instructions
   for (auto &I : ConvBB) {
     // Check if this load is handled in the previous loop
     if (isa<LoadInst>(&I) && Insts.count(&I))
         continue;
 
     // Store
-    if(isa<StoreInst>(&I)) {
+    if(auto *SI = dyn_cast<StoreInst>(&I)) {
+        if (!SI->isSimple() || !IsTypeLegal(SI->getValueOperand()->getType()))
+            return false;
         Insts[&I] = nullptr;
         continue;
     }
 
-    if (I.mayHaveSideEffects())
+    if (I.mayHaveSideEffects() || I.mayReadFromMemory())
       return false;
+
+    if (CallInst *CI = dyn_cast<CallInst>(&I); CI && CI->isConvergent())
+        return false;
   }
 
   LLVM_DEBUG(dbgs() << "Found if-convertible block:\n"

IGC/Compiler/Optimizer/PromoteToPredicatedMemoryAccess.hpp

@@ -27,7 +27,8 @@ namespace IGC
     //     %17 = call <4 x float> @llvm.genx.GenISA.PredicatedLoad.v4f32.p1v4f32.v4f32(<4 x float> addrspace(1)* %bitc0, i64 16, i1 %pred, <4 x float> %mergeValue)
     // if found in specific pattern and then performs if-conversion.
     //
-    // The pass looks for conditional branches that can be if-converted. The only "hammock" form
+    // Constraints:
+    // 1. The pass looks for conditional branches that can be if-converted. The only "hammock" form
     // of the control flow is supported, i.e. the true block has a single
     // predecessor and the false block has two predecessors. The true block must
     // have a single successor that is the false block.
@@ -37,9 +38,11 @@ namespace IGC
     //   }
     //   false block
     //
-    // All the instructions in the true block must be safe to execute in the false
+    // 2. All the instructions in the true block must be safe to execute in the false
     // block. The pass makes the instructions in the true block to be executed
     // conditionally and replaces branch to unconditional one.
+    // 3. Load/Store instructions should be simple.
+    // 4. Load/Store instructions should use legal data types.
     //
     // The pass expects that the simplifycfg pass will be run after it to clean up
     // the CFG.

IGC/Compiler/tests/PromoteToPredicatedMemoryAccess/load.ll

@@ -95,3 +95,105 @@ exit:
   %res = phi <64 x i32> [ %load, %then ], [ %data, %if ], [ zeroinitializer, %entry ]
   ret <64 x i32> %res
 }
+
+; Test exit if load is not simple
+; CHECK-LABEL: @test5(
+define <64 x i32> @test5(<64 x i32> addrspace(1)* %src, i1 %pred, <64 x i32> %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: %load = load volatile <64 x i32>, <64 x i32> addrspace(1)* %src, align 4
+  %load = load volatile <64 x i32>, <64 x i32> addrspace(1)* %src, align 4
+  br label %exit
+
+exit:
+; CHECK: %res = phi <64 x i32> [ %load, %st ], [ %data, %entry ]
+; CHECK: ret <64 x i32> %res
+  %res = phi <64 x i32> [ %load, %st ], [ %data, %entry ]
+  ret <64 x i32> %res
+}
+
+; Test exit if integer size is not legal
+; CHECK-LABEL: @test6(
+define i42 @test6(i42 addrspace(1)* %src, i1 %pred, i42 %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: %load = load i42, i42 addrspace(1)* %src, align 4
+  %load = load i42, i42 addrspace(1)* %src, align 4
+  br label %exit
+
+exit:
+; CHECK: %res = phi i42 [ %load, %st ], [ %data, %entry ]
+; CHECK: ret i42 %res
+  %res = phi i42 [ %load, %st ], [ %data, %entry ]
+  ret i42 %res
+}
+
+; Call instruction with side effects
+; CHECK-LABEL: @test7(
+define i32 @test7(i32 addrspace(1)* %src, i1 %pred, i32 %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: %load = load i32, i32 addrspace(1)* %src, align 4
+  %load = load i32, i32 addrspace(1)* %src, align 4
+  call void @side_effect_function()
+  br label %exit
+
+exit:
+; CHECK: %res = phi i32 [ %load, %st ], [ %data, %entry ]
+; CHECK: ret i32 %res
+  %res = phi i32 [ %load, %st ], [ %data, %entry ]
+  ret i32 %res
+}
+
+declare void @side_effect_function()
+
+; Convergent call
+; CHECK-LABEL: @test8(
+define i32 @test8(i32 addrspace(1)* %src, i1 %pred, i32 %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: %load = load i32, i32 addrspace(1)* %src, align 4
+  %load = load i32, i32 addrspace(1)* %src, align 4
+  call void @convergent_function()
+  br label %exit
+
+exit:
+; CHECK: %res = phi i32 [ %load, %st ], [ %data, %entry ]
+; CHECK: ret i32 %res
+  %res = phi i32 [ %load, %st ], [ %data, %entry ]
+  ret i32 %res
+}
+
+declare void @convergent_function() nounwind readnone willreturn convergent
+
+; Instruction may read from memory
+; CHECK-LABEL: @test9(
+define i32 @test9(i32 addrspace(1)* %src, i32 addrspace(1)* %src1, i1 %pred, i32 %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: %load = load i32, i32 addrspace(1)* %src, align 4
+  %load = load i32, i32 addrspace(1)* %src, align 4
+  %load1 = load i32, i32 addrspace(1)* %src1, align 4
+  br label %exit
+
+exit:
+; CHECK: %res = phi i32 [ %load, %st ], [ %data, %entry ]
+; CHECK: ret i32 %res
+  %res = phi i32 [ %load, %st ], [ %data, %entry ]
+  ret i32 %res
+}

IGC/Compiler/tests/PromoteToPredicatedMemoryAccess/store.ll

@@ -44,3 +44,35 @@ st:
 exit:
   ret void
 }
+
+; early exit, not simple
+; CHECK-LABEL: @test3(
+define void @test3(<64 x i32> addrspace(1)* %dst, i1 %pred, <64 x i32> %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: store volatile <64 x i32> %data, <64 x i32> addrspace(1)* %dst, align 4
+  store volatile <64 x i32> %data, <64 x i32> addrspace(1)* %dst, align 4
+  br label %exit
+
+exit:
+  ret void
+}
+
+; illegal int type, exit
+; CHECK-LABEL: @test4(
+define void @test4(<64 x i33> addrspace(1)* %dst, i1 %pred, <64 x i33> %data) {
+entry:
+; CHECK: br i1 %pred, label %st, label %exit
+  br i1 %pred, label %st, label %exit
+
+st:
+; CHECK: store <64 x i33> %data, <64 x i33> addrspace(1)* %dst, align 4
+  store <64 x i33> %data, <64 x i33> addrspace(1)* %dst, align 4
+  br label %exit
+
+exit:
+  ret void
+}

IGC/ocloc_tests/optimizations/PromoteToPredicatedMemoryAccess/LoadNonUniformBuffer.cl

@@ -24,9 +24,10 @@ __kernel void test(__global const float* buffer1,
     __global const float* buffers[2] = {buffer1, buffer2};
     int gid = get_global_id(0);
     int bufferIndex = gid % 2;
+    __global const float* buffer = buffers[bufferIndex];
 
     float val = 0;
     if (gid < predicate)
-        val = buffers[bufferIndex][gid];
+        val = buffer[gid];
     outputBuffer[gid] = val;
 }

IGC/ocloc_tests/optimizations/PromoteToPredicatedMemoryAccess/StoreSubDW.cl

@@ -16,6 +16,7 @@ SPDX-License-Identifier: MIT
 
 __kernel void basic(__global const char* in, __global char* out, const int predicate) {
     int gid = get_global_id(0);
+    char val = in[gid];
     if (gid < predicate)
-        out[gid] = in[gid];
+        out[gid] = val;
 }

IGC/ocloc_tests/optimizations/PromoteToPredicatedMemoryAccess/StoreUniform.cl

@@ -14,7 +14,7 @@ SPDX-License-Identifier: MIT
 
 // CHECK-ASM:     (W&f{{[0-9\.]+}}) store.ugm.d32x1t.a64 (1|M0)  [r{{[0-9:]+}}]    r{{[0-9:]+}}
 
-__kernel void uniform_store(__global const float* in, __global float* out, const int predicate) {
+__kernel void uniform_store(__global float* out, const int predicate) {
     if (predicate)
-        out[0] = in[0];
+        out[0] = 3;
 }