Revert: Move sub_group_id scalarization from CodeAssumption to WIAnalysis.

---------------------------

(cherry picked from commit b059af8)

Author: AnBodrova

IGC/AdaptorOCL/UnifyIROCL.cpp

@@ -374,6 +374,11 @@ static void CommonOCLBasedPasses(
         mpm.add(createBIFTransformsPass());
     }
 
+    if (IGC_IS_FLAG_ENABLED(EnableCodeAssumption))
+    {
+        mpm.add(new CodeAssumption());
+    }
+
     if (pContext->m_instrTypes.hasFRem)
     {
         mpm.add(new HandleFRemInstructions());
@@ -400,11 +405,6 @@ static void CommonOCLBasedPasses(
     mpm.add(createTimeStatsCounterPass(pContext, TIME_Unify_BuiltinImport, STATS_COUNTER_END));
     mpm.add(new BIFFlagCtrlResolution(pContext));
 
-    if (IGC_IS_FLAG_ENABLED(EnableCodeAssumption))
-    {
-        mpm.add(new CodeAssumption());
-    }
-
     if (IGC_GET_FLAG_VALUE(AllowMem2Reg))
     {
         mpm.add(createPromoteMemoryToRegisterPass());

IGC/Compiler/CISACodeGen/WIAnalysis.cpp

@@ -26,13 +26,11 @@ SPDX-License-Identifier: MIT
 #include <stack>
 #include <sstream>
 #include "Probe/Assertion.h"
-#include <llvmWrapper/IR/PatternMatch.h>
 
 using namespace llvm;
 using namespace IGC;
 using namespace IGC::IGCMD;
 using namespace IGC::Debug;
-using namespace llvm::PatternMatch;
 
 static cl::opt<bool> PrintWiaCheck(
     "print-wia-check", cl::init(false), cl::Hidden,
@@ -1257,134 +1255,6 @@ WIAnalysis::WIDependancy WIAnalysisRunner::calculate_dep(const LoadInst* inst)
     return calculate_dep_simple(inst);
 }
 
-// This functions checks if result of inst is sub_group_id value.
-// In other words, it checks the following pattern:
-//
-// %localSizeX = extractelement <3 x i32> %localSize, i32 0
-// %localSizeY = extractelement <3 x i32> %localSize, i32 1
-// %localIdZZext = zext i16 %localIdZ to i32
-// %mul1 = mul i32 %localSizeY, %localIdZzext
-// %localIdYZext = zext i16 %localIdY to i32
-// %add1 = add i32 %mul1, %localIdYZext
-// %mul2 = mul i32 %add1, %localSizeX
-// %add2 = add i32 %mul2, %localIdX4
-// %simdSize = call i32 @llvm.genx.GenISA.simdSize()
-//
-// pattern 1:
-// %lshr1 = lshr i32 %simdSize, 4
-// %lshr2 = lshr i32 %add2, 3
-// %sgid = lshr i32 %lshr2, %lshr1
-//
-// or
-//
-// pattern 2:
-// %sgid = udiv i32 %add2, %simdSize
-//
-// This pattern matching can be removed, when get_subgroup_id call is inlined after applying the WIAnalysis.
-bool WIAnalysisRunner::check_sg_id(const Instruction* inst)
-{
-    const Value *val = cast<Value>(inst);
-    Value *lshrInst1 = nullptr;
-    Value *lshrInst2 = nullptr;
-    Value *addInst1 = nullptr;
-    Value *addInst2 = nullptr;
-    Value *callVal1 = nullptr;
-    Value *extrVal1 = nullptr;
-    Value *extrVal2 = nullptr;
-    Value *mulInst1 = nullptr;
-    Value *mulInst2 = nullptr;
-    Value *localIdX = nullptr;
-    Value *localIdY = nullptr;
-    Value *localIdZ = nullptr;
-
-    FunctionInfoMetaDataHandle funcInfoMD = m_pMdUtils->getFunctionsInfoItem(m_func);
-    SubGroupSizeMetaDataHandle subGroupSize = funcInfoMD->getSubGroupSize();
-    // If subGroupSize is not set in the metadata, then we need to check the pattern with llvm.genx.GenISA.simdSize call.
-    if (subGroupSize->hasValue())
-    {
-        uint32_t simdSize = (uint32_t)subGroupSize->getSIMDSize();
-        IGC_ASSERT(simdSize == 8 || simdSize == 16 || simdSize == 32);
-        uint32_t power = 0;
-        if (simdSize == 32)
-        {
-            power = 5;
-        }
-        else if (simdSize == 16)
-        {
-            power = 4;
-        }
-        else
-        {
-            power = 3;
-        }
-
-        auto lshrPat1 = m_LShr(m_Value(addInst1), m_SpecificInt(power));
-        if (!match(val, lshrPat1))
-            return false;
-    }
-    else
-    {
-        auto lshrPat1 = m_LShr(m_Value(lshrInst1), m_Value(lshrInst2));
-
-        // Check pattern 1.
-        if (match(val, lshrPat1))
-        {
-            auto lshrPat2 = m_LShr(m_Value(addInst1), m_SpecificInt(3));
-            auto lshrPat3 = m_LShr(m_Value(callVal1), m_SpecificInt(4));
-
-            match(lshrInst1, lshrPat2);
-            match(lshrInst2, lshrPat3);
-        }
-
-        // If pattern 1 was not found, check pattern 2.
-        if (!callVal1 || !addInst1)
-        {
-            auto udivPat1 = m_UDiv(m_Value(addInst1), m_Value(callVal1));
-            if (!match(val, udivPat1))
-                return false;
-        }
-
-        GenIntrinsicInst *genInst = dyn_cast<GenIntrinsicInst>(callVal1);
-        if (!genInst || (genInst->getIntrinsicID() != GenISAIntrinsic::GenISA_simdSize))
-            return false;
-    }
-
-    auto addPat11 =  m_Add(m_Value(mulInst1), m_ZExt(m_Value(localIdX)));
-    auto addPat12 =  m_Add(m_ZExt(m_Value(localIdX)), m_Value(mulInst1));
-    if (!match(addInst1, addPat11) && !match(addInst1, addPat12))
-        return false;
-
-    auto mulPat11 = m_Mul(m_ExtractElt(m_Value(extrVal1), m_SpecificInt(0)), m_Value(addInst2));
-    auto mulPat12 = m_Mul(m_Value(addInst2), m_ExtractElt(m_Value(extrVal1), m_SpecificInt(0)));
-    if (!match(mulInst1, mulPat11) && !match(mulInst1, mulPat12))
-        return false;
-
-    auto addPat21 =  m_Add(m_Value(mulInst2), m_ZExt(m_Value(localIdY)));
-    auto addPat22 =  m_Add(m_ZExt(m_Value(localIdY)), m_Value(mulInst2));
-    if (!match(addInst2, addPat21) && !match(addInst2, addPat22))
-        return false;
-
-    auto mulPat21 = m_Mul(m_ExtractElt(m_Value(extrVal2), m_SpecificInt(1)), m_ZExt(m_Value(localIdZ)));
-    auto mulPat22 = m_Mul(m_ZExt(m_Value(localIdZ)), m_ExtractElt(m_Value(extrVal2), m_SpecificInt(1)));
-    if (!match(mulInst2, mulPat21) && !match(mulInst2, mulPat22))
-        return false;
-
-    ImplicitArgs implicitArgs(*m_func, m_pMdUtils);
-    Value *argX = implicitArgs.getImplicitArgValue(*m_func, ImplicitArg::LOCAL_ID_X, m_pMdUtils);
-    Value *argY = implicitArgs.getImplicitArgValue(*m_func, ImplicitArg::LOCAL_ID_Y, m_pMdUtils);
-    Value *argZ = implicitArgs.getImplicitArgValue(*m_func, ImplicitArg::LOCAL_ID_Z, m_pMdUtils);
-    Value *localSize = implicitArgs.getImplicitArgValue(*m_func, ImplicitArg::LOCAL_SIZE, m_pMdUtils);
-    Value *enqLocalSize = implicitArgs.getImplicitArgValue(*m_func, ImplicitArg::ENQUEUED_LOCAL_WORK_SIZE, m_pMdUtils);
-
-    if ((localIdX != argX) || (localIdY != argY) || (localIdZ != argZ))
-        return false;
-
-    if ((localSize != extrVal1) && (enqLocalSize != extrVal1) || (localSize != extrVal2) && (enqLocalSize != extrVal2))
-        return false;
-
-    return true;
-}
-
 WIAnalysis::WIDependancy WIAnalysisRunner::calculate_dep(
     const BinaryOperator* inst)
 {
@@ -1397,11 +1267,6 @@ WIAnalysis::WIDependancy WIAnalysisRunner::calculate_dep(
     WIAnalysis::WIDependancy dep1 = getDependency(op1);
     IGC_ASSERT(dep1 < WIAnalysis::NumDeps);
 
-    if (check_sg_id(inst))
-    {
-        return WIAnalysis::UNIFORM_THREAD;
-    }
-
     // For whatever binary operation,
     // uniform returns uniform
     WIAnalysis::WIDependancy dep = select_conversion[dep0][dep1];

IGC/Compiler/CISACodeGen/WIAnalysis.hpp

@@ -160,8 +160,6 @@ namespace IGC
             bool assume_uniform = false;
         };
 
-        bool check_sg_id(const llvm::Instruction* inst);
-
         /// @brief Update dependency relations between all values
         void updateDeps();
 

IGC/Compiler/Optimizer/CodeAssumption.cpp

@@ -44,7 +44,9 @@ char CodeAssumption::ID = 0;
 
 bool CodeAssumption::runOnModule(Module& M)
 {
-    m_pMDUtils = getAnalysis<MetaDataUtilsWrapper>().getMetaDataUtils();;
+    m_pMDUtils = getAnalysis<MetaDataUtilsWrapper>().getMetaDataUtils();
+    // Add code assist uniform analysis.
+    uniformHelper(&M);
 
     if (IGC_GET_FLAG_VALUE(EnableCodeAssumption) > 1)
     {
@@ -54,6 +56,57 @@ bool CodeAssumption::runOnModule(Module& M)
     return m_changed;
 }
 
+void CodeAssumption::uniformHelper(Module* M)
+{
+    ModuleMetaData* modMD = getAnalysis<MetaDataUtilsWrapper>().getModuleMetaData();
+
+    for (Module::iterator I = M->begin(), E = M->end(); I != E; ++I)
+    {
+        Function* F = &(*I);
+
+        StringRef FN = F->getName();
+
+        // sub_group_id
+        if (!FN.equals("_Z25__spirv_BuiltInSubgroupIdv") &&
+            !FN.equals("__builtin_spirv_BuiltInSubgroupId") &&
+            !FN.equals("_Z16get_sub_group_idv"))
+            continue;
+        // find all the callees
+        for (auto ui = F->use_begin(), ue = F->use_end(); ui != ue; ++ui) {
+            auto CI = dyn_cast<CallInst>(ui->getUser());
+            if (!CI) continue;
+            auto BB = CI->getParent();
+            auto KF = BB->getParent();
+
+            if (!IsSGIdUniform(m_pMDUtils, modMD, KF))
+                continue;
+
+            // The value must be uniform. Using shuffle with index=0 to
+            // enforce it. assuming lane-0 is active
+            Type* int32Ty = Type::getInt32Ty(M->getContext());
+            Value* args[3];
+            args[0] = CI;
+            args[1] = ConstantInt::getNullValue(int32Ty);
+            args[2] = ConstantInt::get(int32Ty, 0);
+
+            Type* ITys[3] = { args[0]->getType(), int32Ty, int32Ty};
+            Function* shuffleIntrin = GenISAIntrinsic::getDeclaration(
+                M,
+                GenISAIntrinsic::GenISA_WaveShuffleIndex,
+                ITys);
+
+            Instruction* shuffleCall = CallInst::Create(shuffleIntrin, args, "sgid", CI->getNextNode());
+
+            shuffleCall->setDebugLoc(CI->getDebugLoc());
+
+            CI->replaceAllUsesWith(shuffleCall);
+            shuffleCall->setOperand(0, CI);
+
+            m_changed = true;
+        }
+    }
+}
+
 void CodeAssumption::addAssumption(Module* M)
 {
     // Do it for 64-bit pointer only

IGC/Compiler/Optimizer/CodeAssumption.hpp

@@ -63,6 +63,9 @@ namespace IGC
 
         IGCMD::MetaDataUtils* m_pMDUtils = nullptr;
 
+        // Simple change to help uniform analysis (later).
+        void uniformHelper(llvm::Module* M);
+
         // Add llvm.assume to assist other optimization such statelessToStateful
         void addAssumption(llvm::Module* M);
 

IGC/Compiler/tests/CodeAssumption/uniform.ll

@@ -0,0 +1,70 @@
+;=========================== begin_copyright_notice ============================
+;
+; Copyright (C) 2022 Intel Corporation
+;
+; SPDX-License-Identifier: MIT
+;
+;============================ end_copyright_notice =============================
+;
+; RUN: igc_opt -enable-debugify -igc-codeassumption -S < %s 2>&1 | FileCheck %s
+; ------------------------------------------------
+; CodeAssumption : Uniform part
+; ------------------------------------------------
+; Was reduced from ocl test kernel:
+;
+; __kernel void bar(int a)
+; {
+;     int subid = get_sub_group_id();
+;     int sum = subid + a;
+; }
+
+; Debug-info related check
+; CHECK-NOT: WARNING
+; CHECK: CheckModuleDebugify: PASS
+
+; Function Attrs: noinline nounwind
+define spir_kernel void @bar(i32 %a) #0 {
+; CHECK-LABEL: @bar(
+; CHECK:  entry:
+; CHECK:    %subid = alloca i32, align 4
+; CHECK:    [[CALL:%.*]] = call spir_func i32 @__builtin_spirv_BuiltInSubgroupId()
+; CHECK:    [[SGID:%.*]] = call i32 @llvm.genx.GenISA.WaveShuffleIndex.i32.i32.i32(i32 [[CALL]], i32 0, i32 0)
+; CHECK:    store i32 [[SGID]], i32* %subid, align 4
+;
+entry:
+  %a.addr = alloca i32, align 4
+  %subid = alloca i32, align 4
+  %sum = alloca i32, align 4
+  store i32 %a, i32* %a.addr, align 4
+  %call = call spir_func i32 @__builtin_spirv_BuiltInSubgroupId() #2
+  store i32 %call, i32* %subid, align 4
+  %0 = load i32, i32* %subid, align 4
+  %1 = load i32, i32* %a.addr, align 4
+  %add = add nsw i32 %0, %1
+  store i32 %add, i32* %sum, align 4
+  ret void
+}
+
+; Function Attrs: nounwind readnone
+declare spir_func i32 @__builtin_spirv_BuiltInSubgroupId() #2
+
+attributes #0 = { noinline nounwind }
+attributes #2 = { nounwind readnone }
+
+!IGCMetadata = !{!2}
+!igc.functions = !{!13}
+
+!2 = !{!"ModuleMD", !3}
+!3 = !{!"FuncMD", !4, !5}
+!4 = !{!"FuncMDMap[0]", void (i32)* @bar}
+!5 = !{!"FuncMDValue[0]", !6, !7, !11, !12}
+!6 = !{!"localOffsets"}
+!7 = !{!"workGroupWalkOrder", !8, !9, !10}
+!8 = !{!"dim0", i32 0}
+!9 = !{!"dim1", i32 1}
+!10 = !{!"dim2", i32 2}
+!11 = !{!"funcArgs"}
+!12 = !{!"functionType", !"KernelFunction"}
+!13 = !{void (i32)* @bar, !14}
+!14 = !{!15}
+!15 = !{!"function_type", i32 0}