diff --git a/tests/unit/test_umma_qk_hd64.cu b/tests/unit/test_umma_qk_hd64.cu index e8d09579..144810d8 100644 --- a/tests/unit/test_umma_qk_hd64.cu +++ b/tests/unit/test_umma_qk_hd64.cu @@ -54,10 +54,18 @@ test_umma_qk_hd64(const bf16_t* q, const bf16_t* k, } __syncthreads(); - // Simple sanity check: thread 0 writes to sQ[0] - if (tid == 0) sQ[0] = f32_to_bf16(1.0f); + // Sanity check: write to sQ[0] + if (tid == 0) { + uint16_t one_bf16 = f32_to_bf16(1.0f); + sQ[0] = one_bf16; + } __syncthreads(); - if (tid == 0) s_out[250] = bf16_to_f32(sQ[0]); // Should be 1.0 + if (tid == 0) { + uint16_t val = sQ[0]; + float fval = bf16_to_f32(val); + s_out[250] = fval; // Should be 1.0 + s_out[251] = (float)one_bf16; // Raw BF16 bits + } __syncthreads(); // Write Q (1, hd) to row 0 of sQ in canonical layout