Merge commit '0e3bf60855809422b5aad7040b47bc8513e4ab96'

Author: whitneywhtsang

include/triton/Dialect/Triton/Transforms/Passes.td

@@ -90,15 +90,4 @@ def TritonLoopAwareCSE : Pass<"triton-loop-aware-cse", "mlir::ModuleOp"> {
   }];
 }
 
-def TritonSCFToCF : Pass</*cli-arg*/"triton-scf-to-cf", /*Op*/"mlir::ModuleOp"> {
-  let summary = "MLIR's SCF To CF plus some extra attributes propagation.";
-  let description = [{
-    This pass uses MLIR's SCF To CF pass as base. Additionally, it propagates
-    some extra attributes to the converted CFG.
-    TODO: upstream the llvm loop attribute propagation and remove this pass.
-  }];
-
-  let dependentDialects = [];
-}
-
 #endif

lib/Dialect/Triton/Transforms/CMakeLists.txt

@@ -13,7 +13,6 @@ add_triton_library(TritonTransforms
   RewriteTensorDescriptorToPointer.cpp
   ArithTypeConversion.cpp
   FunctionTypeConversion.cpp
-  SCFToCF.cpp
 
   DEPENDS
   TritonTransformsIncGen

lib/Dialect/Triton/Transforms/SCFToCF.cpp

@@ -96,7 +96,6 @@ void init_triton_passes_ttgpuir(py::module &&m) {
 
 void init_triton_passes_convert(py::module &&m) {
   using namespace mlir;
-  ADD_PASS_WRAPPER_0("add_triton_scf_to_cf", mlir::triton::createTritonSCFToCF);
   ADD_PASS_WRAPPER_0("add_scf_to_cf", createSCFToControlFlowPass);
   ADD_PASS_WRAPPER_0("add_cf_to_llvmir", createConvertControlFlowToLLVMPass);
   ADD_PASS_WRAPPER_0("add_index_to_llvmir", createConvertIndexToLLVMPass);

python/src/passes.cc

@@ -287,7 +287,7 @@ def make_llir(src, metadata, options):
         # LDS size is determined by provided arch name.
         custom_lds_size = 0
         amd.passes.ttgpuir.add_optimize_lds_usage(pm, options.arch, custom_lds_size)
-        passes.convert.add_triton_scf_to_cf(pm)
+        passes.convert.add_scf_to_cf(pm)
         passes.convert.add_index_to_llvmir(pm)
 
         amd.passes.ttgpuir.add_allocate_shared_memory(pm)

third_party/amd/backend/compiler.py

@@ -348,7 +348,7 @@ def make_llir(self, src, metadata, options, capability):
 
         passes.ttgpuir.add_combine_tensor_select_and_if(pm)
         passes.ttgpuir.add_allocate_warp_groups(pm)
-        passes.convert.add_triton_scf_to_cf(pm)
+        passes.convert.add_scf_to_cf(pm)
         nvidia.passes.ttgpuir.add_allocate_shared_memory_nv(pm, capability, ptx_version)
         nvidia.passes.ttnvgpuir.add_allocate_tensor_memory(pm)
         if knobs.compilation.enable_experimental_consan: