diff --git a/tests/unit/test_fmha_smem_p.cu b/tests/unit/test_fmha_smem_p.cu
index 4c49fffe..09d1ba87 100644
--- a/tests/unit/test_fmha_smem_p.cu
+++ b/tests/unit/test_fmha_smem_p.cu
@@ -69,7 +69,8 @@ test_fmha_smem_p(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k,
     __syncthreads();
     uint32_t tb = *sTmemBase;
 
-    // ===== STEP 1: QK GEMM =====
+    // ===== STEP 1: QK GEMM — SKIPPED =====
+    /*
     {
         uint64_t dq = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sQ0), BLOCK_MN);
         uint64_t dk = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sK0), BLOCK_MN);
@@ -80,44 +81,19 @@ test_fmha_smem_p(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k,
             __syncthreads();
         }
     }
+    */
 
-    // ===== STEP 2: Softmax — read S from TMEM, write P to SMEM =====
-    if (wid == 0) {
-        float s_vals[SK], row_max = -INFINITY;
-        for (int n = 0; n < SK / 8; n++) {
-            float tmp[8];
-            asm volatile("tcgen05.ld.sync.aligned.32x32b.x8.b32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8];"
-                : "=f"(tmp[0]),"=f"(tmp[1]),"=f"(tmp[2]),"=f"(tmp[3]),
-                  "=f"(tmp[4]),"=f"(tmp[5]),"=f"(tmp[6]),"=f"(tmp[7])
-                : "r"(tb + n*8));
-            asm volatile("tcgen05.wait::ld.sync.aligned;");
-            if (lane == 0) for (int c=0;c<8;c++) {
-                s_vals[n*8+c] = tmp[c] * scale;
-                row_max = fmaxf(row_max, tmp[c] * scale);
-            }
-        }
-        row_max = wmax(row_max);
-        float row_sum = 0.0f;
-        if (lane == 0) for (int j=0;j<SK;j++) {
-            s_vals[j] = expf(s_vals[j] - row_max);
-            row_sum += s_vals[j];
-        }
-        row_sum = wsum(row_sum);
-        if (lane == 0) for (int j=0;j<SK;j++) s_vals[j] /= row_sum;
-
-        // Write P to SMEM in canonical (128, 128) layout
-        // (128, 128): CORES_MN=16, CORES_K=16
-        // Only row 0 has data. Zero all, then write row 0.
-        for (int i = lane; i < 128 * SK; i += 32) sP[i] = 0;
-        if (lane == 0) {
-            for (int j = 0; j < SK; j++) {
-                int core_k = j / 8, lc = j % 8;
-                int dst_idx = core_k * 16 * 64 + 0 * 64 + 0 * 8 + lc;
-                sP[dst_idx] = f32_to_bf16(s_vals[j]);
-            }
+    // ===== STEP 2: Write P directly to SMEM (debug) =====
+    {
+        for (int i = tid; i < 128 * SK; i += 128) sP[i] = 0;
+        __syncthreads();
+        for (int j = tid; j < SK; j += 128) {
+            int core_k = j / 8, lc = j % 8;
+            int dst_idx = core_k * 16 * 64 + 0 * 64 + 0 * 8 + lc;
+            sP[dst_idx] = f32_to_bf16(0.1f);
         }
+        __syncthreads();
     }
-    __syncthreads();
 
     // ===== STEP 3: PV GEMM (SS) =====
     // P(128,128) × V(128,16) → O(128,16) via 8 SS MMA K-tiles