Skip to content

Fix sgl-kernel per_token_quant fp8 kernel scale shared_memory bug #16886

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
Fridge003 merged 6 commits into from
Jan 13, 2026
Merged

Fix sgl-kernel per_token_quant fp8 kernel scale shared_memory bug #16886

Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
b20e54c
ud
BBuf Jan 11, 2026
db3e898
Merge branch 'main' into fix_sgl_kernel_per_token_quant_fp8
Fridge003 Jan 12, 2026
bb0434f
ud
BBuf Jan 12, 2026
6f8d618
Merge branch 'main' into fix_sgl_kernel_per_token_quant_fp8
BBuf Jan 12, 2026
44ab94d
Merge branch 'main' into fix_sgl_kernel_per_token_quant_fp8
BBuf Jan 13, 2026
b1e2975
revert changes in toml
Fridge003 Jan 13, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 4 additions & 3 deletions sgl-kernel/csrc/gemm/per_token_quant_fp8.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,13 +48,14 @@ __global__ void per_token_quant_fp8_kernel(

float warp_max = warpReduceMax(max_value);

__shared__ float scale;
scale = warp_max / FP8_E4M3_MAX;
// NOTE: one CTA has multiple warps (each warp handles one token), so `scale`
// must be per-warp/per-thread (register) instead of a single shared variable.
const float scale = warp_max / FP8_E4M3_MAX;
// Broadcast scale
if (lane_id == 0) {
token_scale[0] = scale;
}
float scale_inv = (scale == 0.f) ? 0.f : 1.0f / scale;
const float scale_inv = (scale == 0.f) ? 0.f : 1.0f / scale;

//
// Pass-2: quantize and write back
Expand Down
Loading