From d840fbbf857ce8dea2222ec8b251e74a38f51bba Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 23:10:49 +0000 Subject: [PATCH] test: clean multirow test with proper SMEM calc --- tests/unit/test_fmha_6warp_multirow.cu | 344 ++++++++++--------------- 1 file changed, 137 insertions(+), 207 deletions(-) diff --git a/tests/unit/test_fmha_6warp_multirow.cu b/tests/unit/test_fmha_6warp_multirow.cu index 2abcf2f6..38c3f44a 100644 --- a/tests/unit/test_fmha_6warp_multirow.cu +++ b/tests/unit/test_fmha_6warp_multirow.cu @@ -1,7 +1,6 @@ /** - * Debug test: verify fmha_6warp_multirow works for T=1 (decode regression). - * Uses printf to trace kernel progress. - * Compile: -DHD_VAL=64 + * Test multi-row FMHA kernel (6-warp, T>1 prefill). + * Compile with -DHD_VAL=64 etc. */ #include @@ -24,225 +23,156 @@ static float bf16_to_f32_host(bf16_t h) { uint32_t u=(uint32_t)h<<16; float f; m constexpr int HD = HD_VAL; constexpr int SK = 128; +constexpr int MAX_T = 128; -// Simplified kernel — T=1 only, no multi-head -__global__ void __launch_bounds__(192) -test_decode_kernel(const bf16_t* __restrict__ q, const bf16_t* __restrict__ k, - const bf16_t* __restrict__ v, bf16_t* __restrict__ o, float* lse_out, - int s_k, float scale) { - const int tid = threadIdx.x; - const int wid = tid / 32; - const int lane = tid % 32; - const int T = 1; +#include "dsv4/kernels/attention/fmha_6warp_multirow.cuh" - static constexpr int NKT_QK = HD / MMA_K_BF16; - static constexpr int NKT_PV = SK / MMA_K_BF16; - static constexpr int N_NSUB = HD / 16; - static constexpr int TILE_SZ = 128 * MMA_K_BF16; - static constexpr int V_SUB_SZ = 16 * MMA_K_BF16; - static constexpr int TMEM_N = (HD <= 128) ? 128 : 256; - static constexpr int CORES_MN = 128 / 8; - - extern __shared__ char sbuf[]; - uint32_t* sTmemBase = (uint32_t*)sbuf; - float* sRowMax = (float*)(sbuf + 8); - float* sRowSum = sRowMax + 1; - bf16_t* sQ0 = (bf16_t*)(((uintptr_t)(sRowSum + 1) + 127) & ~(uintptr_t)127); - bf16_t* sK0 = sQ0 + TILE_SZ; - bf16_t* sPk = (bf16_t*)(((uintptr_t)(sK0 + TILE_SZ) + 127) & ~(uintptr_t)127); - bf16_t* sV = (bf16_t*)(((uintptr_t)(sPk + TILE_SZ) + 127) & ~(uintptr_t)127); - - if (wid == 4) { - uint32_t sp = __cvta_generic_to_shared(sTmemBase); - tmem_alloc(sp, TMEM_N); - } - __syncthreads(); - uint32_t tb = *sTmemBase; - - // QK GEMM - for (int kt = 0; kt < NKT_QK; kt++) { - if (wid == 5) { - for (int i = lane; i < TILE_SZ; i += 32) sQ0[i] = 0; - for (int d = lane; d < MMA_K_BF16; d += 32) { - int full_d = kt * MMA_K_BF16 + d; - if (full_d < HD) { - int ck = d / 8, lc = d % 8; - sQ0[ck * CORES_MN * 64 + 0 * 64 + 0 * 8 + lc] = q[full_d]; - } - } - for (int i = lane; i < TILE_SZ; i += 32) sK0[i] = 0; - for (int r = 0; r < s_k; r++) { - for (int d = lane; d < MMA_K_BF16; d += 32) { - int full_d = kt * MMA_K_BF16 + d; - if (full_d < HD) { - int ck = d / 8, lc = d % 8; - int tmn = r / 8, lr = r % 8; - sK0[ck * CORES_MN * 64 + tmn * 64 + lr * 8 + lc] = k[r * HD + full_d]; - } - } - } - } - __syncthreads(); - if (wid == 4) { - uint32_t idesc = make_idesc(128, 128); - uint64_t dq = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sQ0), 128); - uint64_t dk = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sK0), 128); - if (tid == 128) umma_ss_f16(tb, dq, dk, idesc, kt > 0); - asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); - } - __syncthreads(); - } - - // Softmax — T=1, only warp 0, lane 0 - float row_max = -INFINITY, row_sum = 0.0f; - float s_vals[SK]; - if (wid == 0) { - for (int n = 0; n < SK / 8; n++) { - float tmp[8]; - asm volatile("tcgen05.ld.sync.aligned.32x32b.x8.b32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8];" - : "=f"(tmp[0]),"=f"(tmp[1]),"=f"(tmp[2]),"=f"(tmp[3]), - "=f"(tmp[4]),"=f"(tmp[5]),"=f"(tmp[6]),"=f"(tmp[7]) - : "r"(tb + n*8)); - asm volatile("tcgen05.wait::ld.sync.aligned;"); - if (lane == 0) for (int c=0;c<8;c++) { - s_vals[n*8+c] = tmp[c] * scale; - row_max = fmaxf(row_max, tmp[c] * scale); - } - } - row_max = wmax(row_max); - if (lane == 0) *sRowMax = row_max; - if (lane == 0) for (int j=0;j 0); - asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); - } - __syncthreads(); - } - } - - // Epilogue - if (wid == 0 && lane == 0) { - float rm = *sRowMax, rs = *sRowSum; - float inv = 1.0f / rs; - for (int n = 0; n < N_NSUB * 2; n++) { - float tmp[8]; - asm volatile("tcgen05.ld.sync.aligned.32x32b.x8.b32 {%0,%1,%2,%3,%4,%5,%6,%7},[%8];" - : "=f"(tmp[0]),"=f"(tmp[1]),"=f"(tmp[2]),"=f"(tmp[3]), - "=f"(tmp[4]),"=f"(tmp[5]),"=f"(tmp[6]),"=f"(tmp[7]) - : "r"(tb + n*8)); - asm volatile("tcgen05.wait::ld.sync.aligned;"); - for (int c = 0; c < 8; c++) { - int d = n*8 + c; - if (d < HD) o[d] = f32_to_bf16(tmp[c] * inv); - } - } - if (lse_out) *lse_out = logf(rs) + rm; - } - __syncthreads(); - - if (wid == 4) tmem_dealloc(tb, TMEM_N); +// Compute SMEM size matching the kernel layout exactly +static int compute_smem() { + // Mirror the kernel's SMEM layout: + // sTmemBase(8B) + sRowMax(128*4) + sRowSum(128*4) + align128 + // + sQ0(128*16*2) + sK0(128*16*2) + align128 + // + sPk(128*16*2) + align128 + sV(16*16*2) + slack + size_t off = 0; + off += 8; // sTmemBase + alignment + off += 128 * sizeof(float); // sRowMax + off += 128 * sizeof(float); // sRowSum + off = (off + 127) & ~(size_t)127; // align for sQ0 + off += 128 * MMA_K_BF16 * sizeof(bf16_t); // sQ0 + off += 128 * MMA_K_BF16 * sizeof(bf16_t); // sK0 + off = (off + 127) & ~(size_t)127; // align for sPk + off += 128 * MMA_K_BF16 * sizeof(bf16_t); // sPk + off = (off + 127) & ~(size_t)127; // align for sV + off += 16 * MMA_K_BF16 * sizeof(bf16_t); // sV + off += 256; // slack + return (int)off; } -int main() { - printf("Simple decode test (HD=%d, T=1)\n", HD); +static void reference_attention_multirow( + const bf16_t* q, const bf16_t* k, const bf16_t* v, + float* o_ref, float* lse_ref, + int hd, int T, int s_k, float scale +) { + for (int t = 0; t < T; t++) { + float s[512]; + for (int j = 0; j < s_k; j++) { + float dot = 0.0f; + for (int d = 0; d < hd; d++) + dot += bf16_to_f32_host(q[t * hd + d]) * bf16_to_f32_host(k[j * hd + d]); + s[j] = dot * scale; + } + float mx = -INFINITY; + for (int j = 0; j < s_k; j++) mx = fmaxf(mx, s[j]); + float sm = 0.0f; + for (int j = 0; j < s_k; j++) { s[j] = expf(s[j] - mx); sm += s[j]; } + for (int j = 0; j < s_k; j++) s[j] /= sm; + for (int d = 0; d < hd; d++) { + float ov = 0.0f; + for (int j = 0; j < s_k; j++) ov += s[j] * bf16_to_f32_host(v[d * s_k + j]); + o_ref[t * hd + d] = ov; + } + if (lse_ref) lse_ref[t] = logf(sm) + mx; + } +} + +static int test_single_T(int T, int n_h = 1, int batch = 1) { + printf("\n=== T=%d, n_h=%d, batch=%d, HD=%d, SK=%d ===\n", T, n_h, batch, HD, SK); const float SCALE = 1.0f / sqrtf((float)HD); + int total_heads = batch * n_h; - bf16_t *h_q = (bf16_t*)malloc(HD * sizeof(bf16_t)); - bf16_t *h_k = (bf16_t*)malloc(SK * HD * sizeof(bf16_t)); - bf16_t *h_v = (bf16_t*)malloc(HD * SK * sizeof(bf16_t)); - bf16_t *h_o = (bf16_t*)calloc(HD, sizeof(bf16_t)); + bf16_t* h_q = (bf16_t*)malloc(total_heads * T * HD * sizeof(bf16_t)); + bf16_t* h_k = (bf16_t*)malloc(total_heads * SK * HD * sizeof(bf16_t)); + bf16_t* h_v = (bf16_t*)malloc(total_heads * HD * SK * sizeof(bf16_t)); + bf16_t* h_o = (bf16_t*)calloc(total_heads * T * HD, sizeof(bf16_t)); + float* h_lse = (float*)calloc(total_heads * T, sizeof(float)); - srand(42); - for (int i=0;i>>(d_q, d_k, d_v, d_o, d_lse, SK, SCALE); + int smem = compute_smem(); + if (smem > 48 * 1024) + cudaFuncSetAttribute(fmha_6warp_multirow_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem); + + dim3 grid(1, n_h, batch); + fmha_6warp_multirow_kernel<<>>(params); cudaError_t err = cudaDeviceSynchronize(); if (err != cudaSuccess) { - printf("CUDA ERROR: %s\n", cudaGetErrorString(err)); - printf("FAILED\n"); - } else { - cudaMemcpy(h_o, d_o, HD*sizeof(bf16_t), cudaMemcpyDeviceToHost); - - // Reference - float o_ref[512]; - for (int d = 0; d < HD; d++) o_ref[d] = 0.0f; - float s[128]; - for (int j=0;j1e-4f) { cs+=a*o_ref[d]; na+=a*a; nb+=o_ref[d]*o_ref[d]; } - } - cs /= (sqrtf(na)*sqrtf(nb)+1e-10f); - printf("Cosine similarity: %.8f\n", cs); - printf("%s\n", cs > 0.999f ? "PASSED" : "FAILED"); + printf(" CUDA ERROR: %s\n", cudaGetErrorString(err)); + cudaFree(d_q); cudaFree(d_k); cudaFree(d_v); cudaFree(d_o); cudaFree(d_lse); + free(h_q); free(h_k); free(h_v); free(h_o); free(h_lse); + return 0; } + cudaMemcpy(h_o, d_o, total_heads * T * HD * sizeof(bf16_t), cudaMemcpyDeviceToHost); + cudaMemcpy(h_lse, d_lse, total_heads * T * sizeof(float), cudaMemcpyDeviceToHost); + + int failed = 0; float min_cos = 1.0f; + for (int b = 0; b < batch; b++) { + for (int h = 0; h < n_h; h++) { + int idx = b * n_h + h; + float o_ref[MAX_T * 512]; float lse_ref[MAX_T]; + reference_attention_multirow( + h_q + idx * T * HD, h_k + idx * SK * HD, h_v + idx * HD * SK, + o_ref, lse_ref, HD, T, SK, SCALE); + for (int t = 0; t < T; t++) { + float cs=0,na=0,nb=0; + for (int d=0;d1e-4f){cs+=a*b2;na+=a*a;nb+=b2*b2;} + } + cs /= (sqrtf(na)*sqrtf(nb)+1e-10f); + if(cs