[AMD][TDM] Minor formatting cleanup

Author: jungpark-mlir

python/src/gluon_ir.cc

@@ -1023,10 +1023,10 @@ void init_gluon_ir(py::module &&m) {
            [](GluonOpBuilder &self, Value descPtr, std::vector<Value> &indices,
               Value result, Value pred, Value barrier,
               std::vector<int64_t> warpBases) {
-             auto warpBasesAttr = warpBases.empty()
-                                      ? DenseI64ArrayAttr()
-                                      : DenseI64ArrayAttr::get(
-                                            self.getContext(), warpBases);
+             auto warpBasesAttr =
+                 warpBases.empty()
+                     ? DenseI64ArrayAttr()
+                     : DenseI64ArrayAttr::get(self.getContext(), warpBases);
              self.create<ttag::AsyncTDMCopyGlobalToLocalOp>(
                  descPtr, indices, result, pred, barrier, warpBasesAttr);
            })

third_party/amd/lib/Dialect/TritonAMDGPU/IR/Dialect.cpp

@@ -39,9 +39,9 @@
 #include "Dialect/TritonAMDGPU/IR/Dialect.cpp.inc"
 // clang-format on
 
+#include "third_party/amd/backend/include/TDMCommon.h"
 #include "third_party/amd/include/Dialect/TritonAMDGPU/Utility/CommonUtils.h"
 #include "third_party/amd/lib/TritonAMDGPUToLLVM/TDMUtility.h"
-#include "third_party/amd/backend/include/TDMCommon.h"
 
 using namespace mlir;
 using namespace mlir::triton::amdgpu;

third_party/amd/python/examples/gluon/f16_gemm_warp_pipeline_gfx1250.py

@@ -104,10 +104,11 @@ def gemm_tdm_pipelined_warp_pipelined_kernel(a_ptr, b_ptr, c_ptr,  #
 # Duplicate warps get pred=0 (hardware no-op), freeing TDM bandwidth.
 # ---------------------------------------------------------------------------
 
+
 @gluon.jit
 def issue_loads_specialized(producer, a_desc, b_desc, off_am, off_bn, a_buffer, b_buffer, BLOCK_K: ttgl.constexpr,
-                            NUM_BUFFERS: ttgl.constexpr, TRANSPOSE_B: ttgl.constexpr,
-                            TDM_WARP_BASES: ttgl.constexpr, pred=1):
+                            NUM_BUFFERS: ttgl.constexpr, TRANSPOSE_B: ttgl.constexpr, TDM_WARP_BASES: ttgl.constexpr,
+                            pred=1):
     pred_i32 = pred.to(ttgl.int32) if hasattr(pred, 'to') else pred
     ttgl.amd.gfx1250.tdm.async_load(a_desc, [off_am, producer * BLOCK_K], a_buffer.index(producer % NUM_BUFFERS),
                                     pred=pred_i32, warp_bases=TDM_WARP_BASES)
@@ -171,8 +172,8 @@ def gemm_tdm_specialized_pipelined_warp_pipelined_kernel(a_ptr, b_ptr, c_ptr,  #
         with ttgl.amd.warp_pipeline_stage("stage0", priority=1):
             consumer, a, b = lds_load(consumer, a_buffer, OPERAND_LAYOUT_A, b_buffer, OPERAND_LAYOUT_B, NUM_BUFFERS,
                                       TRANSPOSE_B)
-            producer = issue_loads_specialized(producer, a_desc, b_desc, 0, 0, a_buffer, b_buffer, BLOCK_K,
-                                               NUM_BUFFERS, TRANSPOSE_B, TDM_WARP_BASES)
+            producer = issue_loads_specialized(producer, a_desc, b_desc, 0, 0, a_buffer, b_buffer, BLOCK_K, NUM_BUFFERS,
+                                               TRANSPOSE_B, TDM_WARP_BASES)
         with ttgl.amd.warp_pipeline_stage("stage1", priority=0):
             accumulator = issue_wmma_compute(a, b, accumulator)
         ttgl.amd.gfx1250.tdm.async_wait(2)
@@ -193,6 +194,7 @@ def gemm_tdm_specialized_pipelined_warp_pipelined_kernel(a_ptr, b_ptr, c_ptr,  #
 # Helper
 # ---------------------------------------------------------------------------
 
+
 def _compute_tdm_warp_bases(block_shape, num_warps, active_warps):
     """Compute warp_bases for partial TDM copy with the given active warp count.
 
@@ -231,6 +233,7 @@ def _compute_tdm_warp_bases(block_shape, num_warps, active_warps):
 # Tests
 # ---------------------------------------------------------------------------
 
+
 @pytest.mark.parametrize("BLOCK_M,BLOCK_N,BLOCK_K", [(256, 256, 64)])
 @pytest.mark.parametrize("NUM_BUFFERS", [3])
 @pytest.mark.parametrize("TRANSPOSE_B", [True])