daniel/h3-blackwell-runtime
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Match wide NVFP4 block scale swizzle

Browse source
This commit is contained in:
Daniel Maddern 2026-08-15 02:30:41 +07:00
parent b3a2544936
commit 1200152d79
1 changed files with 8 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

10
src/h3_blackwell_runtime/csrc/nvfp4_scale.cu
View file

@ -120,8 +120,14 @@ __global__ void quantize_nvfp4_bf16_kernel(
block_scale_value = decode_fp8_e4m3(scale_byte); block_scale_value = decode_fp8_e4m3(scale_byte);
} }
if (row < scale_rows) { if (row < scale_rows) {
const int64_t tile_row = (row / 128) * 128 + (row % 32) * 4 + ((row % 128) / 32); int64_t scale_row = row;
block_scale[tile_row * scale_cols + block_col] = scale_byte; int64_t scale_col = block_col;
if (scale_cols >= 32) {
const int64_t row_in_tile = row % 128;
scale_row = (row / 128) * 128 + (block_col / 4) * 16 + ((row_in_tile % 32) / 2);
scale_col = (block_col % 4) + (row_in_tile / 32) * 4 + (row_in_tile % 2) * 16;
}
block_scale[scale_row * scale_cols + scale_col] = scale_byte;
} }
if (row < q_rows && block_col * 8 < q_cols) { if (row < q_rows && block_col * 8 < q_cols) {
#pragma unroll #pragma unroll