From fb390b24e2665800625e30b9c4047688ce6a124f Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 14 May 2026 15:24:18 +0000
Subject: [PATCH] debug: add printf to SF remap kernel to check flat_rank and
 layout shape

---
 .../cutlass_nvfp4_gemm/cutlass_nvfp4_gemm.cu              | 8 ++++++++
 1 file changed, 8 insertions(+)

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 6367e297..6c0178bf 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
@@ -117,6 +117,14 @@ __global__ void remap_sf_to_cutlass_kernel(
 
     constexpr int flat_rank = cute::rank_v<decltype(flat)>;
 
+    // Debug: print flat_rank and first few coordinates (only for idx 0)
+    if (dst_idx == 0) {
+        printf("[remap] flat_rank=%d, MN=%d, K_sf=%d, total=%d\n", flat_rank, MN, K_sf, total);
+        printf("[remap] layout shape: ");
+        cute::print(layout_sf.shape());
+        printf("\n");
+    }
+
     int m, k_sf_out;
 
     if constexpr (flat_rank == 6) {