From 01319d724777689839decd1f6580ac8e66a2096d Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 15:59:22 +0000 Subject: [PATCH] auto: pre-test commit --- tests/unit/test_fmha_gen.cu | 237 ++++++++++++++++++++++++++++ tests/unit/test_fmha_gen_kernel.cuh | 191 ++++++++++++++++++++++ tests/unit/test_fmha_hd128_gen.cu | 3 + tests/unit/test_fmha_hd16_gen.cu | 3 + tests/unit/test_fmha_hd16_v2.cu | 195 +++++++++++++++++++++++ tests/unit/test_fmha_hd256_gen.cu | 3 + tests/unit/test_fmha_hd64_gen.cu | 3 + 7 files changed, 635 insertions(+) create mode 100644 tests/unit/test_fmha_gen.cu create mode 100644 tests/unit/test_fmha_gen_kernel.cuh create mode 100644 tests/unit/test_fmha_hd128_gen.cu create mode 100644 tests/unit/test_fmha_hd16_gen.cu create mode 100644 tests/unit/test_fmha_hd16_v2.cu create mode 100644 tests/unit/test_fmha_hd256_gen.cu create mode 100644 tests/unit/test_fmha_hd64_gen.cu diff --git a/tests/unit/test_fmha_gen.cu b/tests/unit/test_fmha_gen.cu new file mode 100644 index 00000000..717b2fbf --- /dev/null +++ b/tests/unit/test_fmha_gen.cu @@ -0,0 +1,237 @@ +/** + * Generalized FMHA for HD=16/64/128/256 using N=16 PV sub-tiles. + * Compile with -DHD_VAL=64 etc. + * Default HD_VAL=64. + */ + +#include +#include +#include +#include +#include + +#ifndef HD_VAL +#define HD_VAL 64 +#endif + +#include "dsv4/kernels/attention/fmha_common.cuh" +#include "dsv4/kernels/attention/fmha_umma_desc.cuh" + +using namespace dsv4::kernels::attention; + +static bf16_t f32_to_bf16_host(float f) { uint32_t u; memcpy(&u,&f,4); return (uint16_t)(u>>16); } +static float bf16_to_f32_host(bf16_t h) { uint32_t u=(uint32_t)h<<16; float f; memcpy(&f,&u,4); return f; } + +constexpr int HD = HD_VAL; +constexpr int SK = 128, BLOCK_MN = 128; +constexpr int NKT_QK = HD / MMA_K_BF16; +constexpr int NKT_PV = SK / MMA_K_BF16; // 8 +constexpr int TILE_SZ = BLOCK_MN * MMA_K_BF16; // 2048 BF16 +constexpr int N_NSUB = HD / 16; +constexpr int V_SUB_SZ = 256; // (16,16) canonical BF16 +constexpr int TMEM_N = (HD <= 128) ? 128 : 256; + +__global__ void __launch_bounds__(128) +fmha_kernel(const bf16_t* q, const bf16_t* k, const bf16_t* v, + bf16_t* o_out, float* o_scalar, float scale) +{ + const int tid = threadIdx.x, wid = tid / 32, lane = tid % 32; + + extern __shared__ char sbuf[]; + uint32_t* sTmemBase = (uint32_t*)sbuf; + bf16_t* sQ0 = (bf16_t*)(((uintptr_t)(sbuf + 4) + 15) & ~(uintptr_t)15); + 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); + float* s_p_vals = (float*)(sV + V_SUB_SZ); + + // TMEM alloc + if (wid == 1) tmem_alloc(__cvta_generic_to_shared(sTmemBase), TMEM_N); + __syncthreads(); + uint32_t tb = *sTmemBase; + + // ===== QK GEMM (one K-tile at a time) ===== + { + uint32_t idesc = make_idesc(BLOCK_MN, BLOCK_MN); + for (int kt = 0; kt < NKT_QK; kt++) { + for (int i = tid; i < TILE_SZ; i += 128) sQ0[i] = 0; + for (int d = tid; d < MMA_K_BF16; d += 128) { + int ck = d / 8, lc = d % 8; + sQ0[ck * 16 * 64 + lc] = q[kt * MMA_K_BF16 + d]; + } + for (int i = tid; i < TILE_SZ; i += 128) sK0[i] = 0; + for (int r = 0; r < SK; r++) { + for (int d = tid; d < MMA_K_BF16; d += 128) { + int ck = d / 8, lc = d % 8; + int tmn = r / 8, lr = r % 8; + sK0[ck * 16 * 64 + tmn * 64 + lr * 8 + lc] = k[r * HD + kt * MMA_K_BF16 + d]; + } + } + __syncthreads(); + + uint64_t dq = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sQ0), BLOCK_MN); + uint64_t dk = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sK0), BLOCK_MN); + if (tid == 0) umma_ss_f16(tb, dq, dk, idesc, kt > 0); + asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); + __syncthreads(); + } + } + + // ===== Softmax (row 0 only) ===== + if (wid == 0) { + float s_vals[SK], row_max = -INFINITY; + 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); + float row_sum = 0.0f; + if (lane == 0) for (int j=0;j 0); + asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); + __syncthreads(); + } + } + } + + // ===== Epilogue ===== + if (wid == 0) { + float o_vals[HD]; + for (int n = 0; n < HD / 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++) o_vals[n*8+c] = tmp[c]; + } + if (lane == 0) for (int d=0;d 48 * 1024) { + cudaFuncSetAttribute(fmha_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem); + } + fmha_kernel<<<1, 128, smem>>>(d_q, d_k, d_v, d_o, d_o_scalar, SCALE); + + cudaError_t launch_err = cudaGetLastError(); + if (launch_err != cudaSuccess) { printf("LAUNCH ERROR: %s\n", cudaGetErrorString(launch_err)); return 1; } + + cudaError_t err = cudaDeviceSynchronize(); + if (err != cudaSuccess) { printf("CUDA ERROR: %s\n", cudaGetErrorString(err)); return 1; } + + cudaMemcpy(h_o, d_o, HD*sizeof(bf16_t), cudaMemcpyDeviceToHost); + cudaMemcpy(h_o_scalar, d_o_scalar, HD*sizeof(float), cudaMemcpyDeviceToHost); + + printf("O[0..7] MMA: "); for(int d=0;d1e-4f) { cs+=a*b; na+=a*a; nb+=b*b; } + } + cs /= (sqrtf(na)*sqrtf(nb)+1e-10f); + printf("Filtered cosine: %.8f\n", cs); + printf("Test %s\n", cs > 0.999f ? "PASSED" : "FAILED"); + + cudaFree(d_q); cudaFree(d_k); cudaFree(d_v); cudaFree(d_o); cudaFree(d_o_scalar); + free(h_q); free(h_k); free(h_v); free(h_o); free(h_o_scalar); + return cs > 0.999f ? 0 : 1; +} diff --git a/tests/unit/test_fmha_gen_kernel.cuh b/tests/unit/test_fmha_gen_kernel.cuh new file mode 100644 index 00000000..6d40649e --- /dev/null +++ b/tests/unit/test_fmha_gen_kernel.cuh @@ -0,0 +1,191 @@ +/** + * Generalized FMHA for HD=16/64/128/256 using N=16 PV sub-tiles. + * + * Key design: PV MMA uses N=16 sub-tiles to avoid the Layout D N≠16,128 bug. + * For HD values: n_n_subtiles = HD/16 PV calls per K-tile. + * Each sub-tile writes 16 TMEM columns starting at offset n*16. + * + * Pipeline: QK(SS, N=128) → softmax → PV(SS, N=16, sub-tiled) → epilogue + */ + +#include +#include +#include +#include +#include + +#include "dsv4/kernels/attention/fmha_common.cuh" +#include "dsv4/kernels/attention/fmha_umma_desc.cuh" + +using namespace dsv4::kernels::attention; + +static bf16_t f32_to_bf16_host(float f) { uint32_t u; memcpy(&u,&f,4); return (uint16_t)(u>>16); } +static float bf16_to_f32_host(bf16_t h) { uint32_t u=(uint32_t)h<<16; float f; memcpy(&f,&u,4); return f; } + +constexpr int SK = 128, BLOCK_MN = 128; +constexpr int NKT_QK = HD_VAL / MMA_K_BF16; +constexpr int NKT_PV = SK / MMA_K_BF16; // 8 +constexpr int TILE_SZ = BLOCK_MN * MMA_K_BF16; // 2048 BF16 +constexpr int N_NSUB = HD_VAL / 16; // Number of N=16 sub-tiles for PV +constexpr int V_SUB_SZ = 256; // (16,16) canonical BF16 + +__global__ void __launch_bounds__(128) +fmha_kernel(const bf16_t* q, const bf16_t* k, const bf16_t* v, + bf16_t* o_out, float* o_scalar, float scale) +{ + const int tid = threadIdx.x, wid = tid / 32, lane = tid % 32; + + extern __shared__ char sbuf[]; + uint32_t* sTmemBase = (uint32_t*)sbuf; + bf16_t* sQ0 = (bf16_t*)(((uintptr_t)(sbuf + 4) + 15) & ~(uintptr_t)15); + bf16_t* sK0 = sQ0 + NKT_QK * TILE_SZ; + bf16_t* sPk = (bf16_t*)(((uintptr_t)(sK0 + NKT_QK * TILE_SZ) + 127) & ~(uintptr_t)127); + bf16_t* sV = (bf16_t*)(((uintptr_t)(sPk + TILE_SZ) + 127) & ~(uintptr_t)127); + float* s_p_vals = (float*)(sV + V_SUB_SZ); + + // Load Q K-tiles + for (int kt = 0; kt < NKT_QK; kt++) { + bf16_t* sq = sQ0 + kt * TILE_SZ; + for (int i = tid; i < TILE_SZ; i += 128) sq[i] = 0; + for (int d = tid; d < MMA_K_BF16; d += 128) { + int ck = d / 8, lc = d % 8; + sq[ck * 16 * 64 + lc] = q[kt * MMA_K_BF16 + d]; + } + } + + // Load K K-tiles + for (int kt = 0; kt < NKT_QK; kt++) { + bf16_t* sk = sK0 + kt * TILE_SZ; + for (int i = tid; i < TILE_SZ; i += 128) sk[i] = 0; + for (int r = 0; r < SK; r++) { + for (int d = tid; d < MMA_K_BF16; d += 128) { + int ck = d / 8, lc = d % 8; + int tmn = r / 8, lr = r % 8; + sk[ck * 16 * 64 + tmn * 64 + lr * 8 + lc] = k[r * HD_VAL + kt * MMA_K_BF16 + d]; + } + } + } + + __syncthreads(); + + // TMEM alloc: need max(128, HD) columns for QK + PV output + constexpr int TMEM_N = 128; // Always 128 (enough for QK N=128 and PV up to HD=128) + if (wid == 1) tmem_alloc(__cvta_generic_to_shared(sTmemBase), TMEM_N); + __syncthreads(); + uint32_t tb = *sTmemBase; + + // ===== QK GEMM (N=128, proven working) ===== + { + uint32_t idesc = make_idesc(BLOCK_MN, BLOCK_MN); + for (int kt = 0; kt < NKT_QK; kt++) { + bf16_t* sq = sQ0 + kt * TILE_SZ; + bf16_t* sk = sK0 + kt * TILE_SZ; + uint64_t dq = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sq), BLOCK_MN); + uint64_t dk = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sk), BLOCK_MN); + if (tid == 0) umma_ss_f16(tb, dq, dk, idesc, kt > 0); + asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); + __syncthreads(); + } + } + + // ===== Softmax (row 0 only for T=1 decode) ===== + if (wid == 0) { + float s_vals[SK], row_max = -INFINITY; + 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); + float row_sum = 0.0f; + if (lane == 0) for (int j=0;j 0); + asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); + __syncthreads(); + } + } + } + + // ===== Epilogue ===== + if (wid == 0) { + float o_vals[HD_VAL]; + for (int n = 0; n < HD_VAL / 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++) o_vals[n*8+c] = tmp[c]; + } + if (lane == 0) for (int d=0;d +#include +#include +#include +#include + +#include "dsv4/kernels/attention/fmha_common.cuh" +#include "dsv4/kernels/attention/fmha_umma_desc.cuh" + +using namespace dsv4::kernels::attention; + +static bf16_t f32_to_bf16_host(float f) { uint32_t u; memcpy(&u,&f,4); return (uint16_t)(u>>16); } +static float bf16_to_f32_host(bf16_t h) { uint32_t u=(uint32_t)h<<16; float f; memcpy(&f,&u,4); return f; } + +constexpr int HD = HD_VAL; +constexpr int SK = 128, BLOCK_MN = 128; +constexpr int NKT_QK = HD / MMA_K_BF16; +constexpr int NKT_PV = SK / MMA_K_BF16; // 8 +constexpr int TILE_SZ = BLOCK_MN * MMA_K_BF16; // 2048 BF16 +constexpr int N_NSUB = HD / 16; // Number of N=16 sub-tiles for PV +constexpr int V_SUB_SZ = 256; // (16,16) canonical BF16 +// TMEM columns: need at least 128 for QK, and HD for PV output +constexpr int TMEM_N = (HD <= 128) ? 128 : ((HD <= 256) ? 256 : 512); +// Ensure power of 2, min 32 +// (All our values 128/256/512 are already powers of 2 and >= 32) + +__global__ void __launch_bounds__(128) +fmha_kernel(const bf16_t* q, const bf16_t* k, const bf16_t* v, + bf16_t* o_out, float* o_scalar, float scale) +{ + const int tid = threadIdx.x, wid = tid / 32, lane = tid % 32; + + extern __shared__ char sbuf[]; + uint32_t* sTmemBase = (uint32_t*)sbuf; + bf16_t* sQ0 = (bf16_t*)(((uintptr_t)(sbuf + 4) + 15) & ~(uintptr_t)15); + bf16_t* sK0 = sQ0 + TILE_SZ; // Only 1 K-tile at a time + 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); + float* s_p_vals = (float*)(sV + V_SUB_SZ); + + // TMEM alloc + if (wid == 1) tmem_alloc(__cvta_generic_to_shared(sTmemBase), TMEM_N); + __syncthreads(); + uint32_t tb = *sTmemBase; + + // ===== QK GEMM (one K-tile at a time) ===== + { + uint32_t idesc = make_idesc(BLOCK_MN, BLOCK_MN); + for (int kt = 0; kt < NKT_QK; kt++) { + // Load Q K-tile kt + for (int i = tid; i < TILE_SZ; i += 128) sQ0[i] = 0; + for (int d = tid; d < MMA_K_BF16; d += 128) { + int ck = d / 8, lc = d % 8; + sQ0[ck * 16 * 64 + lc] = q[kt * MMA_K_BF16 + d]; + } + // Load K K-tile kt + for (int i = tid; i < TILE_SZ; i += 128) sK0[i] = 0; + for (int r = 0; r < SK; r++) { + for (int d = tid; d < MMA_K_BF16; d += 128) { + int ck = d / 8, lc = d % 8; + int tmn = r / 8, lr = r % 8; + sK0[ck * 16 * 64 + tmn * 64 + lr * 8 + lc] = k[r * HD + kt * MMA_K_BF16 + d]; + } + } + __syncthreads(); + + uint64_t dq = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sQ0), BLOCK_MN); + uint64_t dk = make_umma_desc_kmajor_none(__cvta_generic_to_shared(sK0), BLOCK_MN); + if (tid == 0) umma_ss_f16(tb, dq, dk, idesc, kt > 0); + asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); + __syncthreads(); + } + } + + // ===== Softmax (row 0 only for T=1 decode) ===== + if (wid == 0) { + float s_vals[SK], row_max = -INFINITY; + 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); + float row_sum = 0.0f; + if (lane == 0) for (int j=0;j 0); + asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); + __syncthreads(); + } + } + } + + // ===== Epilogue ===== + if (wid == 0) { + float o_vals[HD]; + for (int n = 0; n < HD / 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++) o_vals[n*8+c] = tmp[c]; + } + if (lane == 0) for (int d=0;d