diff --git a/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu b/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu index fd4e29e..18c3a56 100644 --- a/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu +++ b/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu @@ -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;