diff --git a/vllm/patches/staging_kernel.py b/vllm/patches/staging_kernel.py index 679f3a97..769a4b5f 100644 --- a/vllm/patches/staging_kernel.py +++ b/vllm/patches/staging_kernel.py @@ -150,9 +150,8 @@ def _deepseek_v4_stage_mega_moe_inputs_kernel( # Pack E2M1 pairs into single bytes (2 per byte, low nibble first) # mxf4nvf4 reads FP4 packed from SMEM — must match kernel's TMA layout - e2m1_flat = tl.reshape(e2m1_4bit, [BLOCK_K]) - e2m1_lo = e2m1_flat[0::2] # even indices → low nibble - e2m1_hi = e2m1_flat[1::2] # odd indices → high nibble + e2m1_lo = e2m1_4bit[0::2] # even indices → low nibble + e2m1_hi = e2m1_4bit[1::2] # odd indices → high nibble e2m1_packed = (e2m1_hi << 4 | e2m1_lo).to(tl.uint8) # [BLOCK_K // 2] k_offsets_out = k_block_id * (BLOCK_K // 2) + tl.arange(0, BLOCK_K // 2)