diff --git a/tests/unit/test_umma_qk.cu b/tests/unit/test_umma_qk.cu index 8d69704d..cc5dd634 100644 --- a/tests/unit/test_umma_qk.cu +++ b/tests/unit/test_umma_qk.cu @@ -106,22 +106,22 @@ test_umma_qk_hd16( s_out[133] = (float)sQ_smem; s_out[134] = (float)sK_smem; s_out[135] = (float)__cvta_generic_to_shared(sQ); + s_out[136] = (float)(sQ_smem >> 4); // start_address for desc + s_out[137] = (float)(sQ_smem & 0xF); // alignment check } __syncthreads(); // ================================================================ - // Call tcgen05.mma SS + // SKIP MMA for now — just test descriptor construction // ================================================================ - if (tid == 0) { - umma_ss_f16(tmem_base, desc_q, desc_k, idesc, /*accumulate=*/false); - } - __syncwarp(); - - // Wait for MMA - if (wid == 0 && lane == 0) { - tmem_fence_store(); - } - __syncthreads(); + // if (tid == 0) { + // umma_ss_f16(tmem_base, desc_q, desc_k, idesc, /*accumulate=*/false); + // } + // __syncwarp(); + // if (wid == 0 && lane == 0) { + // tmem_fence_store(); + // } + // __syncthreads(); // ================================================================ // Read S from TMEM