clean up

Author: masahi

tests/python/unittest/test_mma_16x8x16_4k_tune.py

@@ -285,7 +285,6 @@ def mma_fill_impl(a: T.handle) -> None:
 workload = te.create_prim_func(te_workload.matmul_fp16(n=N, m=M, k=K))
 
 tune = False
-use_ldmatrix = True
 
 
 def schedule(sch: tir.Schedule):
@@ -408,25 +407,8 @@ def shared_16x16_to_ldmatrix_32x8_layout(i, j):
     sch.transform_layout(B_warp, 0, "write", index_map=shared_16x16_to_ldmatrix_32x8_layout)
     sch.transform_layout(C_warp, 0, "read", index_map=shared_16x16_to_ldmatrix_32x8_layout)
 
-    # return
-
-    if use_ldmatrix:
-        sch.tensorize(loop_a, "mma.ldmatrix_a")
-        sch.tensorize(loop_b, "mma.ldmatrix_b")
-    else:
-        warp_loop1, warp_loop2 = sch.get_loops(A_warp)[-2:]
-        f_0, f_1 = sch.split(warp_loop1, factors=[None, 8])
-        f_2, f_3 = sch.split(warp_loop2, factors=[None, 2])
-        sch.reorder(f_1, f_2, f_0, f_3)
-        fused_1 = sch.fuse(f_1, f_2)
-        fused_2 = sch.fuse(f_0, f_3)
-        sch.bind(fused_1, "threadIdx.x")
-
-        warp_loop1, warp_loop2 = sch.get_loops(B_warp)[-2:]
-        f_0, f_1 = sch.split(warp_loop1, factors=[4, 2])
-        sch.reorder(warp_loop2, f_0, f_1)
-        fused_1 = sch.fuse(warp_loop2, f_0)
-        sch.bind(fused_1, "threadIdx.x")
+    sch.tensorize(loop_a, "mma.ldmatrix_a")
+    sch.tensorize(loop_b, "mma.ldmatrix_b")
 
     mma_loop = sch.get_loops(block_inner)[-3]
     sch.tensorize(mma_loop, "mma_sync")