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 1, 20:00 UTC
[Bugfix][Kernel] Fix mHC fused-RMSNorm big-fuse miscompile for hidden_size != 4096 (#44692)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>
Yongye Zhu committed
Jun 5, 2026 at 22:44 UTC
ec0a31d4aa019567395d9e7cce67d1e6ac56974a
1 file changed
+1
-1
vllm/model_executor/kernels/mhc/tilelang_kernels.py
+1
-1
@@ -309,7 +309,7 @@ def mhc_pre_big_fuse_with_norm_tilelang(
309
sumsq_per_pos = T.alloc_fragment(hidden_block, T.float32)
310
T.clear(sumsq_per_pos)
311
312
- for i0_h in T.Pipelined(hidden_size // hidden_block, num_stages=3):
312
+ for i0_h in T.Pipelined(hidden_size // hidden_block, num_stages=2):
313
xs = T.alloc_shared((hc_mult, hidden_block), T.bfloat16)
314
xl = T.alloc_fragment((hc_mult, hidden_block), T.float32)
315
T.copy(residual[i, 0, i0_h * hidden_block], xs)