From d5949a23b4a2b909bc9412164e5762c3b4de0b20 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Fri, 15 May 2026 11:39:57 +0000 Subject: [PATCH] =?UTF-8?q?fix:=20use=20cute::crd2idx=20for=20SF=20remap?= =?UTF-8?q?=20=E2=80=94=20layout=5Fsf()=20not=20directly=20callable?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CuTe Layout objects with hierarchical shapes can't be called directly with flat (m, k_sf). Use cute::crd2idx(make_coord(m, k_sf), layout_sf) to convert logical coordinates to physical indices. --- .../cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 a8715762..8b208e1f 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 @@ -122,7 +122,7 @@ __global__ void remap_sf_to_cutlass_kernel( if (m >= MN || k_sf >= K_sf) return; // Compute the CUTLASS physical index from logical (m, k_sf) via the layout - int dst_idx = layout_sf(cute::make_coord(m, k_sf)); + int dst_idx = cute::crd2idx(cute::make_coord(m, k_sf), layout_sf); // Read from source (row-major or col-major) and write to CUTLASS position dst[dst_idx] = col_major_src ? src[k_sf * MN + m] : src[m * K_sf + k_sf];