AMD specific scheduling pass for TTGIR instructions (#483)

* Add scheduling pass

* Add LDSWr, LDSRd sinking to scheduler

* Add improvements to AMDReorderInstructions.cpp

* Remove unwanted change in fa tutorial

* Small fixes in fa tutorials

* Address review comments and add some minor fixes

* Enable use of scheduling pass along with amd pipelining

* Don't use reordering together with pipelining

---------

Co-authored-by: Ognjen <oplavsic@luxoft.com>
Co-authored-by: Ognjen Plavsic <ognjen.plavsic@luxoft.com>

include/triton/Dialect/TritonGPU/Transforms/Passes.h

@@ -29,6 +29,8 @@ std::unique_ptr<Pass> createTritonGPUCoalescePass();
 
 std::unique_ptr<Pass> createTritonGPUReorderInstructionsPass();
 
+std::unique_ptr<Pass> createTritonAMDGPUReorderInstructionsPass();
+
 std::unique_ptr<Pass> createTritonGPUDecomposeConversionsPass();
 
 std::unique_ptr<Pass> createTritonGPURemoveLayoutConversionsPass();

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

@@ -214,6 +214,18 @@ def TritonGPUReorderInstructions: Pass<"tritongpu-reorder-instructions", "mlir::
                            "mlir::triton::TritonDialect"];
 }
 
+def TritonAMDGPUReorderInstructions: Pass<"tritonamdgpu-reorder-instructions", "mlir::ModuleOp"> {
+  let summary = "Reorder instructions";
+
+  let description = "This pass reorder instructions so as to be more suitable for good code generation from "
+                    "LLVM backend.";
+
+  let constructor = "mlir::createTritonAMDGPUReorderInstructionsPass()";
+
+  let dependentDialects = ["mlir::triton::gpu::TritonGPUDialect",
+                           "mlir::triton::TritonDialect"];
+}
+
 def TritonGPUDecomposeConversions: Pass<"tritongpu-decompose-conversions", "mlir::ModuleOp"> {
   let summary = "Decompose convert[distributed -> dotOperand] into convert[distributed -> shared -> dotOperand]";