From 482328160aee9f675ca22c792a73968f94ddbae1 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 13:43:24 +0000 Subject: [PATCH] test: single PV K-tile debug --- tests/unit/test_fmha_ts_hd16.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/unit/test_fmha_ts_hd16.cu b/tests/unit/test_fmha_ts_hd16.cu index d88fd029..4566a045 100644 --- a/tests/unit/test_fmha_ts_hd16.cu +++ b/tests/unit/test_fmha_ts_hd16.cu @@ -137,12 +137,12 @@ test_fmha_ts(const bf16_t* q, const bf16_t* k, const bf16_t* v, // MMA: M=128, N=16. K=16 per MMA call, 8 K-tiles total. uint32_t idesc_pv = make_idesc(BLOCK_MN, HD); - for (int kt = 0; kt < VKT; kt++) { + for (int kt = 0; kt < 1; kt++) { // DEBUG: only first PV K-tile bf16_t* sv = sV_base + kt * V_TILE_SZ; uint64_t dv = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sv), MMA_K_BF16); uint32_t tmem_a = tb + kt * MMA_K_BF16; // P's K-tile columns [16*kt, 16*kt+15] - if (tid == 0) umma_ts_f16(tb_o, tmem_a, dv, idesc_pv, kt > 0); + if (tid == 0) umma_ts_f16(tb_o, tmem_a, dv, idesc_pv, false); // no accumulate for first tile asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); __syncthreads(); }