diff --git a/tests/unit/test_tma_kqk.cu b/tests/unit/test_tma_kqk.cu index d0d95610..e8acc343 100644 --- a/tests/unit/test_tma_kqk.cu +++ b/tests/unit/test_tma_kqk.cu @@ -164,13 +164,13 @@ int main() { float* h_out = (float*)malloc(SK * sizeof(float)); cudaMemcpy(h_out, d_out, SK * sizeof(float), cudaMemcpyDeviceToHost); - float scale = 1.0f / sqrtf((float)HD); + // Reference: raw dot product (MMA is unscaled, scale applied during softmax) int fail = 0; float max_rel = 0; for (int j = 0; j < SK; j++) { float dot = 0; for (int d = 0; d < HD; d++) dot += bf16_to_f32_host(h_q[d]) * bf16_to_f32_host(h_k[j * HD + d]); - float ref = dot * scale; + float ref = dot; // NO scale — MMA output is unscaled float got = h_out[j]; float rel = fabsf(ref) > 1e-4f ? fabsf(got - ref) / fabsf(ref) : fabsf(got - ref); if (rel > max_rel) max_rel = rel;