From 2d3bbbd4450175cd5305a523311513ca1a34b179 Mon Sep 17 00:00:00 2001 From: LZ-QWQ Date: Fri, 16 Oct 2026 15:44:02 +0800 Subject: [PATCH 1/2] fix(kernel): preserve fp4 scale rounding at subnormal ties --- python/freetoken/kernel/triton/dsv41/pack.py | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/python/freetoken/kernel/triton/dsv41/pack.py b/python/freetoken/kernel/triton/dsv41/pack.py index 47a9164e6e..fb1eea3176 100644 --- a/python/freetoken/kernel/triton/dsv41/pack.py +++ b/python/freetoken/kernel/triton/dsv41/pack.py @@ -87,7 +87,8 @@ def _pack_rows_kernel( amax = tl.maximum(tl.max(tl.abs(xe), axis=1), tl.max(tl.abs(xo), axis=1)) if FMT == 2: amax = tl.maximum(amax, 6.0 * (2.0 ** -9)) - s = _e4m3_round(amax * (1.0 / 6.0)) + # Round the division before e4m3's magic-add; fusing it changes subnormal ties. + s = _e4m3_round(tl.math.div_rn(amax, 6.0)) tl.store(dst + D // 2 + offs_g, _e4m3_codes(s)) else: amax = tl.maximum(amax, 6.0 * (2.0 ** -126)) From 0dc42918f59b67335be5f2c8d0bcc931b1754554 Mon Sep 17 00:00:00 2001 From: LZ-QWQ Date: Sat, 17 Oct 2026 02:18:46 +0800 Subject: [PATCH 2/2] fix(kernel): preserve native e4m3 scale calculation --- python/freetoken/kernel/triton/dsv41/pack.py | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/python/freetoken/kernel/triton/dsv41/pack.py b/python/freetoken/kernel/triton/dsv41/pack.py index fb1eea3176..c0fca83162 100644 --- a/python/freetoken/kernel/triton/dsv41/pack.py +++ b/python/freetoken/kernel/triton/dsv41/pack.py @@ -87,8 +87,12 @@ def _pack_rows_kernel( amax = tl.maximum(tl.max(tl.abs(xe), axis=1), tl.max(tl.abs(xo), axis=1)) if FMT == 2: amax = tl.maximum(amax, 6.0 * (2.0 ** -9)) - # Round the division before e4m3's magic-add; fusing it changes subnormal ties. - s = _e4m3_round(tl.math.div_rn(amax, 6.0)) + if e4m3_native_cx(): + s = amax * (1.0 / 6.0) + else: + # Round before the emulated magic-add; fusing it changes subnormal ties. + s = tl.math.div_rn(amax, 6.0) + s = _e4m3_round(s) tl.store(dst + D // 2 + offs_g, _e4m3_codes(s)) else: amax = tl.maximum(amax, 6.0 * (2.0 ** -126))