@tduncan / Nvidia-Optimization-Stack-VLLM Mirror · read-only / commits / 6fdb49392e
Fork 0

A high-throughput and memory-efficient inference and serving engine for LLMs

Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git

Synced Nov 3, 22:45 UTC

[Bugfix] Fix int32 overflow in DeepGEMM SiLU/mul FP8 Triton kernel (#42201)

Signed-off-by: vensen <vensenmu@gmail.com> Signed-off-by: Vensen <vensenmu@gmail.com> Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>

Vensen committed May 12, 2026 at 01:52 UTC 6fdb49392e5183c310c4703b883e7edbf25f6a96
1 file changed +3 -3
vllm/model_executor/layers/quantization/utils/fp8_utils.py
+3 -3
@@ -171,7 +171,7 @@ def _silu_mul_quant_fp8_packed_kernel(
171
172 pid_pack = tl.program_id(0)
173 pid_m = tl.program_id(1)
174 - m_offset = pid_m * BLOCK_M
174 + m_offset = pid_m.to(tl.int64) * BLOCK_M
175
176 if m_offset >= M:
177 return
@@ -321,8 +321,8 @@ def _silu_mul_per_token_group_quant_fp8_colmajor(
321 pid_n = tl.program_id(1)
322 N_2 = N // 2
323
324 - m_offset = pid_m * BLOCK_M
325 - n_offset = pid_n * BLOCK_N
324 + m_offset = pid_m.to(tl.int64) * BLOCK_M
325 + n_offset = pid_n.to(tl.int64) * BLOCK_N
326 if m_offset >= M:
327 return
328