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

debug: single-thread SF layout dump at specific indices

Browse Source
This commit is contained in:
biondizzle
2026-05-14 16:13:05 +00:00
parent 17e6033ade
commit c4b5b52a33
1 changed files with 22 additions and 13 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -131,20 +131,29 @@ __global__ void remap_sf_to_cutlass_kernel(
printf("[remap] flat_rank=%d MN=%d K_sf=%d total=%d\n", R, MN, K_sf, total);
}
// Debug: print specific indices to understand the coordinate decomposition
// Print idx 0, 1, 4, 16, 128, 512, 512*448, 128*448
if (dst_idx == 0 || dst_idx == 1 || dst_idx == 4 || dst_idx == 16 ||
dst_idx == 128 || dst_idx == 512 || dst_idx == 512*448 ||
dst_idx == 128*448 || dst_idx == 4 || dst_idx == 3) {
printf("[remap] idx=%d", dst_idx);
if constexpr (R >= 8) {
printf(" f=%d,%d,%d,%d,%d,%d,%d,%d",
int(cute::get<0>(flat)), int(cute::get<1>(flat)),
int(cute::get<2>(flat)), int(cute::get<3>(flat)),
int(cute::get<4>(flat)), int(cute::get<5>(flat)),
int(cute::get<6>(flat)), int(cute::get<7>(flat)));
// Debug: single-thread print at specific indices
// Only thread 0 prints, and only at specific dst_idx values
// This avoids garbled output from concurrent threads
if (threadIdx.x == 0 && blockIdx.x == 0) {
// Print a series of indices to understand the decomposition
int debug_indices[] = {0, 1, 2, 3, 4, 15, 16, 31, 32, 63, 64, 127, 128, 255, 256, 511, 512, 512*448-1};
int n_debug = sizeof(debug_indices) / sizeof(debug_indices[0]);
for (int di = 0; di < n_debug; di++) {
int didx = debug_indices[di];
if (didx < total) {
auto dcoord = cute::idx2crd(didx, layout_sf.shape(), layout_sf.stride());
auto dflat = cute::flatten(dcoord);
printf("[remap] idx=%d f=", didx);
if constexpr (R >= 8) {
printf("%d,%d,%d,%d,%d,%d,%d,%d",
int(cute::get<0>(dflat)), int(cute::get<1>(dflat)),
int(cute::get<2>(dflat)), int(cute::get<3>(dflat)),
int(cute::get<4>(dflat)), int(cute::get<5>(dflat)),
int(cute::get<6>(dflat)), int(cute::get<7>(dflat)));
}
printf("\n");
}
}
printf(" m=%d k=%d\n", m, k_sf);
}
int m = 0, k_sf = 0;