biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

revert: back to layout_sf(make_coord(...)) — crd2idx was unnecessary

Browse Source
This commit is contained in:
biondizzle
2026-05-15 21:55:00 +00:00
parent ea678ece64
commit a1fd4d6233
1 changed files with 2 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
src/nvfp4_megamoe_kernel/cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu
View File

@@ -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;