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

Match NVFP4 block scale swizzle

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

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

@ -120,7 +120,8 @@ __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) {
block_scale[row * scale_cols + block_col] = scale_byte; const int64_t tile_row = (row / 128) * 128 + (row % 32) * 4 + ((row % 128) / 32);
block_scale[tile_row * scale_cols + block_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