From 4483539f019cbc03bf29b8e50cad3b9f844a1e7d Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 11:47:56 +0000 Subject: [PATCH] test: HD=64 random data, 4 K-tiles, accumulate --- tests/unit/test_umma_qk_hd64.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/unit/test_umma_qk_hd64.cu b/tests/unit/test_umma_qk_hd64.cu index 76d88851..66eb50f0 100644 --- a/tests/unit/test_umma_qk_hd64.cu +++ b/tests/unit/test_umma_qk_hd64.cu @@ -29,7 +29,7 @@ test_umma_qk_hd64(const bf16_t* q, const bf16_t* k, { const int tid = threadIdx.x; const int wid = tid / 32, lane = tid % 32; - const int n_ktiles = 1; // hd / MMA_K; // Force 1 K-tile for debugging + const int n_ktiles = hd / MMA_K; // 4 for hd=64 // SMEM: sQ (128, HD) canonical + sK (128, HD) canonical // Each K-tile of (128, 16) = 4096 bytes @@ -194,8 +194,8 @@ int main() { float* h_s_out = (float*)calloc(128 * 16, sizeof(float)); float* h_s_scalar = (float*)calloc(SK, sizeof(float)); - 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); + 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); 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));