From abe1870429316992177fbe03e5efb49a3ffa380d Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 11:44:31 +0000 Subject: [PATCH] test: HD=64 all-ones, expected S[0,j]=64 (unscaled) or 8.0 scaled --- tests/unit/test_umma_qk_hd64.cu | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/tests/unit/test_umma_qk_hd64.cu b/tests/unit/test_umma_qk_hd64.cu index 5726216a..6a671f32 100644 --- a/tests/unit/test_umma_qk_hd64.cu +++ b/tests/unit/test_umma_qk_hd64.cu @@ -130,6 +130,14 @@ test_umma_qk_hd64(const bf16_t* q, const bf16_t* k, bool accumulate = (kt > 0); + // Debug: write descriptor info + if (tid == 0 && kt == 0) { + s_out[240] = (float)q_ktile_addr; + s_out[241] = (float)k_ktile_addr; + s_out[242] = (float)desc_q; + s_out[243] = (float)desc_k; + } + // 4 warp leaders call MMA if (lane == 0) { umma_ss_f16(tb, desc_q, desc_k, idesc, accumulate); @@ -186,9 +194,8 @@ int main() { float* h_s_out = (float*)calloc(128 * 16, sizeof(float)); float* h_s_scalar = (float*)calloc(SK, sizeof(float)); - srand(42); - for (int d = 0; d < HD; d++) h_q[d] = f32_to_bf16_host((float)(rand()%100)/100.0f - 0.5f); - for (int i = 0; i < SK*HD; i++) h_k[i] = f32_to_bf16_host((float)(rand()%100)/100.0f - 0.5f); + for (int d = 0; d < HD; d++) h_q[d] = f32_to_bf16_host(1.0f); + for (int i = 0; i < SK*HD; i++) h_k[i] = f32_to_bf16_host(1.0f); bf16_t *d_q, *d_k; float *d_s_out, *d_s_scalar; cudaMalloc(&d_q, HD*sizeof(bf16_t)); cudaMalloc(&d_k, SK*HD*sizeof(bf16_t));