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

Match H3 NVFP4 row scale swizzle

Browse source
This commit is contained in:
Daniel Maddern 2026-08-15 02:40:09 +07:00
parent fbfb3453de
commit ef84a2786c
1 changed files with 3 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -130,8 +130,9 @@ __global__ void quantize_nvfp4_bf16_kernel(
const int64_t col_stride = scale_cols == 336 ? 16 : 128;
const int64_t phase_offset = scale_cols == 336 ? 176 : 512;
int64_t swizzled_col_base = pair * col_stride + phase * phase_offset;
scale_row = (row / 128) * 128 + pair * (scale_cols == 336 ? 3 : 1) + phase + swizzled_col_base / scale_cols + ((row_in_tile % 32) / 2);
scale_col = (swizzled_col_base % scale_cols) + (block_col % 4) + (row_in_tile / 32) * 4 + (row_in_tile % 2) * 16;
int64_t swizzled_col = (swizzled_col_base % scale_cols) + (block_col % 4) + (row_in_tile / 32) * 4 + (row_in_tile % 32) * 16;
scale_row = (row / 128) * 128 + pair * (scale_cols == 336 ? 3 : 1) + (scale_cols == 336 ? phase : 0) + swizzled_col_base / scale_cols + swizzled_col / scale_cols;
scale_col = swizzled_col % scale_cols;
if (scale_col >= scale_cols) {
scale_col -= scale_cols;
scale_row += 1;