Enable the SLPVectorizer on Triton side with no scheduling part. Remove the SLPVectorizer once the IGCVectorizer works well.

Author: chengjunlu

third_party/intel/backend/compiler.py

@@ -310,7 +310,7 @@ def make_llir(src, metadata, options):
             paths = [path for (name, path) in options.extern_libs]
             llvm.link_extern_libs(llvm_mod, paths)
         intel.optimize_module(llvm_mod, llvm.OPTIMIZE_O3)
-        if os.getenv("TRITON_INTEL_ENABLE_POST_PROCESS_LLIR", "0") == "1":
+        if os.getenv("TRITON_INTEL_ENABLE_POST_PROCESS_LLIR", "1") == "1":
             intel.post_process_llir(llvm_mod)
 
         # Get some metadata
@@ -335,6 +335,9 @@ def make_spv(src, metadata, options):
         else:
             metadata["build_flags"] = ""
 
+        if os.getenv("TRITON_INTEL_ENABLE_POST_PROCESS_LLIR", "1") == "1":
+            metadata["build_flags"] += " -igc_opts 'DisablePHIScalarization=1'"
+
         if options.generate_native_code:
             with tempfile.NamedTemporaryFile(delete=False, mode='wb', suffix='.spv') as fsrc, \
                 tempfile.NamedTemporaryFile(delete=False, mode='r', suffix='.log') as flog:

third_party/intel/lib/Target/LLVMIR/SLPVectorizer.cpp

@@ -13408,10 +13408,6 @@ Value *BoUpSLP::vectorizeTree(
     const ExtraValueToDebugLocsMap &ExternallyUsedValues,
     SmallVectorImpl<std::pair<Value *, Value *>> &ReplacedExternals,
     Instruction *ReductionRoot) {
-  // All blocks must be scheduled before any instructions are inserted.
-  for (auto &BSIter : BlocksSchedules) {
-    scheduleBlock(BSIter.second.get());
-  }
   // Clean Entry-to-LastInstruction table. It can be affected after scheduling,
   // need to rebuild it.
   EntryToLastInstruction.clear();
@@ -15269,7 +15265,11 @@ bool SLPVectorizerPass::runImpl(Function &F, ScalarEvolution *SE_,
     R.clearReductionData();
 
     // Vectorize trees that end at reductions.
+    //    llvm::outs() << "BB befoer vectorize:\n" << *BB << "\n";
+    //    llvm::outs().flush();
     Changed |= vectorizeChainsInBlock(BB, R);
+    //    llvm::outs() << "BB after vectorize:\n" << *BB << "\n";
+    //    llvm::outs().flush();
   }
 
   if (Changed) {
@@ -15436,14 +15436,25 @@ bool SLPVectorizerPass::vectorizeInsertElementInst(InsertElementInst *IEI,
     return false;
 
   LLVM_DEBUG(dbgs() << "SLP: array mappable to vector: " << *IEI << "\n");
-  return tryToVectorizeList(BuildVectorInsts, R);
+  //  llvm::outs() << "Before Vectorize the instruction Function:\n" <<
+  //  *BB->getParent() << "\n";
+  auto changed = tryToVectorizeList(BuildVectorInsts, R);
+  //  llvm::outs() << "After Vectorize the instruction BB:\n" <<
+  //  *BB->getParent() << "\n"; llvm::outs().flush();
+
+  return changed;
 }
 
 bool SLPVectorizerPass::vectorizeInserts(InstSetVector &Instructions,
                                          BasicBlock *BB, BoUpSLP &R) {
   assert(all_of(Instructions, IsaPred<InsertElementInst>) &&
          "This function only accepts Insert instructions");
   bool OpsChanged = false;
+
+  //  for (auto *I : Instructions) {
+  //    llvm::outs() << "Vectorize insert instructions: " << *I << "\n";
+  //    llvm::outs().flush();
+  //  }
   // Try to match and vectorize a buildvector sequence.
   for (auto *I : reverse(Instructions)) {
     if (R.isDeleted(I))
@@ -15545,8 +15556,11 @@ void mlir::triton::intel::SLPVectorizer(llvm::Module &mod, bool trace) {
   FunctionPassManager FPM;
   FPM.addPass(SLPVectorizerPass(trace));
 
+  //  ::llvm::setCurrentDebugType("SLP");
+  //  ::llvm::DebugFlag = true;
   for (llvm::Function &function : mod.functions()) {
     if (isCandidate(function))
       FPM.run(function, FAM);
   }
+  //  ::llvm::DebugFlag = false;
 }