[Backend][NFC] Expose DecomposeScaledBlocked Pass (triton-lang#8078)

Exposed and Refactored DecomposeScaledBlocked pass so that it can be
used by other backends.

Modifications include:
- Moved the pass definition to a header file.
- Made `cvtDotOperand` a member method to be reused later.
- Made `getTransposeOrder` a static method.

This is one of a series of PRs to decompose scaled dot on AMD backend.

Author: knwng

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

@@ -1,7 +1,40 @@
 #include "mlir/IR/PatternMatch.h"
+#include "triton/Dialect/Triton/IR/Dialect.h"
 
 namespace mlir::triton::gpu {
 
+class DecomposeScaledBlocked : public OpRewritePattern<DotScaledOp> {
+public:
+  DecomposeScaledBlocked(MLIRContext *context, PatternBenefit benefit)
+      : OpRewritePattern<DotScaledOp>(context, benefit) {}
+
+  LogicalResult matchAndRewrite(DotScaledOp scaledDotOp,
+                                PatternRewriter &rewriter) const override;
+
+protected:
+  FloatType getComputeType(ScaleDotElemType aType, ScaleDotElemType bType,
+                           PatternRewriter &rewriter) const;
+  TypedValue<RankedTensorType> scaleTo16(PatternRewriter &rewriter,
+                                         TypedValue<RankedTensorType> scale,
+                                         FloatType computeType) const;
+  TypedValue<RankedTensorType>
+  broadcastScale(PatternRewriter &rewriter, DotScaledOp scaledDotOp,
+                 ModuleOp mod, TypedValue<RankedTensorType> scale,
+                 int dim) const;
+  TypedValue<RankedTensorType> maskNan(PatternRewriter &rewriter,
+                                       DotScaledOp scaledDotOp, ModuleOp mod,
+                                       TypedValue<RankedTensorType> mxfp,
+                                       TypedValue<RankedTensorType> scale,
+                                       int dim) const;
+  TypedValue<RankedTensorType> scaleArg(PatternRewriter &rewriter,
+                                        DotScaledOp scaledDotOp, int opIdx,
+                                        FloatType computeType) const;
+  TypedValue<RankedTensorType>
+  cvtDotOperand(PatternRewriter &rewriter, DotScaledOp scaledDotOp, int opIdx,
+                TypedValue<RankedTensorType> v) const;
+  static SmallVector<int, 2> getTransposeOrder(int rank);
+};
+
 void populateDecomposeScaledBlockedPatterns(mlir::RewritePatternSet &patterns,
                                             int benefit);