[ROCm][Bugfix] Keep zero MiniMax MXFP8 activation blocks finite (#58089)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Andreas Karatzas committed
Sep 23, 2026 at 03:42 UTC
56b3acb55bb27c4c9df8c48e6effc343b986929c
1 file changed
+4
-2
vllm/models/minimax_m3/amd/ops/swiglu_oai.py
+4
-2
@@ -113,8 +113,10 @@ def _swiglu_oai_quant_kernel(
113
act = gate * tl.sigmoid(alpha * gate) * (up + beta) # [BLOCK_M, 32] fp32
114
amax = tl.maximum(tl.max(tl.abs(act), axis=1), TINY) # [BLOCK_M]
115
sb = tl.minimum(tl.maximum(tl.ceil(tl.log2(amax / FP8_MAX)) + 127.0, 0.0), 254.0)
116
- descale = tl.exp2(sb - 127.0)
117
- aq = (act / descale[:, None]).to(aq_ptr.dtype.element_ty)
116
+ # Computing the E8M0 byte-0 divisor (2^-127) can flush to zero on CDNA.
117
+ # Use its finite reciprocal so all-zero blocks do not produce 0/0.
118
+ rescale = tl.exp2(127.0 - sb)
119
+ aq = (act * rescale[:, None]).to(aq_ptr.dtype.element_ty)
120
tl.store(
121
aq_ptr + offs_m[:, None] * stride_qm + offs_c[None, :] * stride_qn,
122
aq,