From c4b5b52a33816985d772130f423bf3b9166be56a Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 14 May 2026 16:13:05 +0000
Subject: [PATCH] debug: single-thread SF layout dump at specific indices

---
 .../cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu  | 35 ++++++++++++-------
 1 file changed, 22 insertions(+), 13 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 4053279e..9a7cf49f 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
@@ -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;