biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

auto: pre-test commit

Browse Source
This commit is contained in:
biondizzle
2026-05-28 16:30:55 +00:00
parent 2a6d72912a
commit e71c957561
1 changed files with 1 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
dsv4/kernels/attention/fmha_6warp.cuh
View File

@@ -211,7 +211,6 @@ fmha_6warp_kernel(
// Epilogue: TMEM → regs → normalize → BF16 → GMEM (warp 0)
// ================================================================
if (wid == 0) {
float inv_sum = 1.0f / *sRowSum;
float o_vals[HD];
for (int n = 0; n < HD / 8; n++) {
float tmp[8];
@@ -220,7 +219,7 @@ fmha_6warp_kernel(
"=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] * inv_sum;
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<HD;d++) o[d] = f32_to_bf16(o_vals[d]);
}