From a1fd4d62330b0ad889bd1246c79705105921693b Mon Sep 17 00:00:00 2001 From: biondizzle Date: Fri, 15 May 2026 21:55:00 +0000 Subject: [PATCH] =?UTF-8?q?revert:=20back=20to=20layout=5Fsf(make=5Fcoord(?= =?UTF-8?q?...))=20=E2=80=94=20crd2idx=20was=20unnecessary?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/src/nvfp4_megamoe_kernel/cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu b/src/nvfp4_megamoe_kernel/cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu index 62089d11..855810cd 100644 --- a/src/nvfp4_megamoe_kernel/cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu +++ b/src/nvfp4_megamoe_kernel/cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu @@ -143,11 +143,9 @@ __global__ void remap_sf_to_cutlass_kernel( // For SFB: logical coord is (n, k_element, l) // k_element = k_sf * InputSFVectorSize (convert SF group to K element) if constexpr (LayoutRank == 2) { - auto logical_coord = cute::make_coord(mn, k_sf * InputSFVectorSize); - dst_idx = cute::crd2idx(logical_coord, layout_sf); + dst_idx = layout_sf(cute::make_coord(mn, k_sf * InputSFVectorSize)); } else if constexpr (LayoutRank == 3) { - auto logical_coord = cute::make_coord(mn, k_sf * InputSFVectorSize, 0); - dst_idx = cute::crd2idx(logical_coord, layout_sf); + dst_idx = layout_sf(cute::make_coord(mn, k_sf * InputSFVectorSize, 0)); } int src_idx = mn * src_stride_mn + k_sf * src_stride_ksf;