Skip to content

npu: quantise the activation row on an AVX2 path, byte-identical to the shipped loop - #908

Merged
bong-water-water-bong merged 1 commit into
mainfrom
1bit/quant-i8-row
Nov 3, 2026
Merged

bong-water-water-bong merged 1 commit into
mainfrom
1bit/quant-i8-row

Conversation

@bong-water-water-bong

Copy link
Copy Markdown
Collaborator

What

Every NPU GEMM quantises one activation row to int8 on the host before its launch, in the scalar loop in I8Ctx::quantize_async / quantize_async_rows. The engine builds at the baseline x86-64 targets, where roundf() is a libm call, so that loop is a large part of the host time inside npu_gemm.

This moves the row loop into npu/forward/include/quant_i8.h as a scalar reference plus an AVX2 path, dispatched at run time on __builtin_cpu_supports("avx2") with the reference used for the tail. Nothing else changes: the callers keep their amax/scale handling and pass the same 1.0f / scale.

This is the second half of the audit row #873 started: #873 removed 99.2 % of the activation pad memset (measured +2.60 ms/token) and left this loop alone.

Measured

Workload is the shipping container's own plan — bonsai2-27b-t2s2-serve/plan.json: 288 projections, 2,310,144 A elements per decode token (48 gdn layers × K = 5120/5120/6144/5120/17408 + 16 attention layers × 5120/5120/17408). Host replay of exactly that mix, best of 7, five runs:

arm ms/token ns/element
scalar (shipped loop) 3.96 – 4.05 1.72 – 1.76
AVX2 dispatch 0.55 0.238

7.2 – 7.4×, bytes identical on every projection of the mix in all five runs.

In situ, npu-hold, 3 prompts × 24 greedy tokens on two packs, one round each:

pack npu_gemm total delta
bonsai2-27b-t2s2-serve 267.35 → 264.01 ms −2.46 ms/token
bonsai2-27b-pb16i 376.66 → 345.54 ms −8.69 ms/token

Both in the direction and order of the host measurement. Caveat, stated rather than buried: this is a single round per arm, and the derived [lh] gemm host row is not used here because the two pb16i arms ran at 53 °C and 63 °C and that row moved 31 ms while the top-level total moved 8.7. rounds=2 alternates the arms so drift cancels; not run, because the identity gate is the gate and the array is shared.

Gate: byte-identical output, at three strengths, on two packs

Tokens, every logprob float, the top-5 floats per token, and the sha256 of the final hidden state — all identical to the same-commit base:

pack prompt 0 prompt 1 prompt 2
t2s2-serve (what :18097 serves) ✔ 5 tokens ✔ 2 tokens ✔ 24 tokens
pb16i ✔ 5 tokens ✔ 2 tokens ✔ 24 tokens

0 amdxdna errors, 0 TDR, no canary short reads. Both arms are the same commit (a2879378), one with the header, one without — symbol check: 0 vs 1 q4nx_quant_i8::row_avx2.

Why it is byte-identical by construction, not by luck

roundf is round-half-away-from-zero, and

t = trunc(x)   exact
f = x - t      exact, |f| < 1
x rounds away from zero iff |f| >= 0.5   (compared exactly)
t ± 1          exact

so trunc + an exact half adjustment is the same integer roundf produces. The vector path reproduces the shipped loop's three sharp edges, each checked against the shipped code rather than assumed:

  • NaN and ±inf are zeroed on the input (isfinite), so a large finite input whose product overflows is not zeroed;
  • such an input saturates through the (int) cast, which x86 makes INT_MIN, so the clamp yields −127, not 127 — the vector path must not "fix" this, and the inrange mask exists precisely because t + adj would wrap to INT_MAX for negative overflow;
  • |v·ais| just inside 2³¹ takes the adjustment normally.

On non-x86 hosts, or without AVX2, the scalar loop runs unchanged.

Tests

tests/quant_i8_row_test.cpp — no device, no model, 0.11 s in CI: the .5 table (including 126.5 → 127), both clamps, NaN, ±inf, FLT_MAX, ±2³¹ and just inside, denormals, -0.0, six real layer scales, tail lengths, and 6.5 M elements of random bit patterns per run (env-tunable via QUANT_I8_FUZZ_ROUNDS), where NaN payloads and denormals appear at their real density instead of only where a case list remembers to put them.

The test is load-bearing, verified by mutation rather than asserted: flipping _CMP_GE_OQ → _CMP_GT_OQ on the half comparison makes it fail with round-half-to-even answers (in=67.5 … got 67 want 68). A 97.5 M-element run of the same fuzz against the shipped loop was clean.

tools/copyright.py --check clean; 1bit builds under the thermal guard.

…he shipped loop (target 2)

Every NPU GEMM quantises one activation row to int8 on the host before its launch, in a scalar loop
built at the baseline x86-64 targets where roundf() is a libm call. On the container's own plan
(bonsai2-27b-t2s2-serve/plan.json: 288 projections, 2,310,144 A elements per decode token) that loop
costs 3.96-4.05 ms/token measured; the AVX2 path costs 0.55 ms/token (0.238 ns/element), 7.2-7.4x,
bytes identical on every projection of the mix.

The loop moves to npu/forward/include/quant_i8.h as a scalar reference plus an AVX2 path dispatched on
__builtin_cpu_supports("avx2"), with the reference used for the tail. Identity is by construction, not
by luck: trunc + an exactly-compared half adjustment is roundf's round-half-away-from-zero, and the
vector path reproduces the shipped loop's three sharp edges -- NaN/+-inf zeroed on the *input*, a finite
input whose product overflows left to saturate through the (int) cast (which x86 makes INT_MIN, so the
clamp yields -127, not 127), and |v*ais| just inside 2^31 taking the adjustment normally.

tests/quant_i8_row_test.cpp proves it with no device and no model: the .5 table, both clamps, NaN, +-inf,
fmax, +-2^31 and just inside it, denormals, -0.0, six real layer scales, lengths that leave a tail, and
6.5 M elements of random *bit patterns* per run (env-tunable rounds), where NaN payloads and denormals
turn up at their real density. It is load-bearing, not decorative: GE -> GT on the half comparison makes
it fail with round-half-to-even answers (67.5 -> 67 vs 68); 97.5 M elements of the same fuzz ran clean
against the shipped loop.

This is the other half of audit row 1: #873 removed 99.2% of the pad memset (+2.60 ms/token) and left
this loop alone.
@github-actions

github-actions Bot commented Nov 3, 2026

Copy link
Copy Markdown

PR Reviewer Guide 🔍

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶

873 - Partially compliant

Compliant requirements:

  • The PR removes the activation BO's pad rows zeroing when a previous launch wrote them
  • The performance improvement is measured and reported (+2.60 ms/token = 0.63 %)
  • The implementation is bit-exact with the previous scalar loop
  • Tests are added to verify quantization correctness

Non-compliant requirements:

  • Documentation updates are not mentioned in the PR description or diff

Requires further human verification:

  • Verification that the documentation (docs/, README, wiki) has been updated to reflect the changes
⏱️ Estimated effort to review: 4 🔵🔵🔵🔵⚪
🧪 PR contains tests
🔒 No security concerns identified
⚡ Recommended focus areas for review

AVX2 Implementation Edge Case Handling

The AVX2 implementation in row_avx2 handles edge cases like NaN, infinities, and values that overflow to 2^31, but the handling may differ from the scalar implementation in subtle ways. Specifically, the check for finite values and the handling of overflow cases should be carefully reviewed to ensure bit-exactness. The comment on line 95 mentions that outside +-2^31 the cast saturates, but it's important to verify that this behavior matches exactly with the scalar path.

    for (; k + 8 <= n; k += 8) {
        // isfinite() is applied to the input, so only NaN and +-inf are zeroed -- a large finite input
        // whose product overflows to inf is NOT zeroed, and saturates through the cast like the scalar.
        __m256 in = _mm256_loadu_ps(a + k);
        in = _mm256_and_ps(in, _mm256_cmp_ps(in, in, _CMP_ORD_Q));
        __m256 isinf = _mm256_or_ps(_mm256_cmp_ps(in, pinf, _CMP_EQ_OQ),
                                    _mm256_cmp_ps(in, ninf, _CMP_EQ_OQ));
        const __m256 v = _mm256_mul_ps(_mm256_andnot_ps(isinf, in), vais);
        const __m256i t = _mm256_cvttps_epi32(v);
        const __m256 f = _mm256_sub_ps(v, _mm256_cvtepi32_ps(t));
        __m256i adj = _mm256_and_si256(
            _mm256_castps_si256(_mm256_cmp_ps(f, half, _CMP_GE_OQ)), one);
        adj = _mm256_or_si256(
            adj, _mm256_and_si256(_mm256_castps_si256(_mm256_cmp_ps(f, nhalf, _CMP_LE_OQ)),
                                  minusone));
        // outside +-2^31 the cast has saturated: leave INT_MIN alone so the clamp yields -127, as scalar
        const __m256i inrange =
            _mm256_castps_si256(_mm256_cmp_ps(_mm256_and_ps(v, absmask), lim, _CMP_LT_OQ));
        __m256i q = _mm256_add_epi32(t, _mm256_and_si256(adj, inrange));
        q = _mm256_min_epi32(_mm256_max_epi32(q, vmin), vmax);
        // 8 int32 -> 8 int8 in lane order (packs saturate, but every lane is already inside +-127)
        const __m128i p16 =
            _mm_packs_epi32(_mm256_castsi256_si128(q), _mm256_extracti128_si256(q, 1));
        _mm_storel_epi64(reinterpret_cast<__m128i*>(out + k), _mm_packs_epi16(p16, p16));
    }
    if (k < n) row_scalar(a + k, n - k, ais, out + k);   // tail: the reference, unchanged
}
Memory Zeroing Optimization

The change optimizes memory zeroing by only zeroing rows [am, a_written_rows_) instead of the entire activation BO. This optimization is correct as long as the a_written_rows_ is properly managed and updated. The change correctly sets a_written_rows_ = am after zeroing, ensuring that the next launch will not re-zero already written rows. However, it's important to ensure that this logic is consistent across all code paths and that there are no race conditions or incorrect assumptions about the state of the activation BO.

memset(Am, 0, (size_t)MD * KD);
for (int m = 0; m < am; m++)
    q4nx_quant_i8::row(A + (size_t)m * ak, ak, ais, Am + (size_t)m * KD);
return Am;

@bong-water-water-bong bong-water-water-bong left a comment

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

VERDICT: READY

At 87dbe494f. The claim is "byte-identical", so I fuzzed the two arms against each other myself, in the engine's configuration (no -mavx2: row_avx2 carries target("avx2") and row() dispatches at run time, so the scalar arm is what a baseline x86-64 build compiles).

82,241 comparisons, 0 mismatches      (-O2 and -O3, and again with no -mavx2)

over, in one run: named edges — NaN, ±inf, ±0.0, ±0.5 and nextafter either side of it, denormals, ±2^31 exact and nextafter(2^31, 0), ±3e38, ±2^24 — crossed with nine scales (1, 0.5, 127, 63.5, 1e-30, 1e30, +inf, -1, 2) and every length 1..33 so the tail after the last 8-lane block is covered; then every edge value placed alone in each of the eight lanes; then 40,000 random raw bit patterns (every exponent, every NaN payload) and 40,000 normal activations at a realistic scale. Zero mismatches anywhere.

The three cases the comment says are modelled rather than assumed, checked in the same run:

  • NaN in, 0 out — and with the ordered mask zeroes the lane, matching if (!isfinite(v)) v = 0.
  • A finite input whose product overflows — scalar roundf(±inf) casts to INT_MIN and clamps to -127, and the vector path reproduces that (the inrange mask suppresses the rounding adjustment so the saturated cvttps value survives to the clamp). Both signs.
  • |v| ≥ 2^31 — same suppression, so it saturates rather than being "fixed" to 127. The boundary is exact because at that magnitude the float spacing is 256, so there is no representable value between 2^31 - 256 and 2^31.

Their own test passes here too: quant_i8_row: PASS (scalar + avx2 arms), 200,000 rounds / 6,535,086 elements / 0 failures. Worth saying that the test re-spells the scalar reference instead of calling the header's copy — that is what stops an edit to the reference from moving both sides together, and it is the right shape.

The speed claim, measured here (rows of 5120 f32, the shipping element count of 2,310,144/token, best of 7, pinned core, under thermal-run):

scalar 1.539 ns/element -> 3.56 ms/token
avx2   0.235 ns/element -> 0.54 ms/token
speedup 6.56x

So the AVX2 side matches the body (0.235 vs 0.238 ns/element, 0.54 vs 0.55 ms/token) but I get 6.6x, not 7.2-7.4x — the scalar figure in the comment (1.72-1.76 ns/element) is about 12 % slower than this box measures today. Not a gate; just make the number in the comment one that a re-run reproduces, so the next person does not go looking for a regression.

Checked by reading, for the wiring: both call sites (quantize_async, quantize_async_rows) now call q4nx_quant_i8::row(...) with the memset(Am, 0, MD*KD) still in front and only the first ak columns written, so the padded columns keep the zero behaviour they had; the CMake target is bound to a target ci.yml already builds (the #635 "Not Run" trap), and a non-AVX2 host still compiles and takes the scalar loop.

No device work: this path is host-side, and the device effect (3.4 ms of a ~415 ms token) is the lane's to measure on the array.

@bong-water-water-bong bong-water-water-bong added the tprs-ready TPRS: review verdict READY posted by the standing reviewer; merge gate satisfied label Nov 3, 2026
@bong-water-water-bong
bong-water-water-bong merged commit 39121fe into main Nov 3, 2026
11 checks passed
@bong-water-water-bong
bong-water-water-bong deleted the 1bit/quant-i8-row branch November 3, 2026 17:17
bong-water-water-bong added a commit that referenced this pull request Nov 4, 2026
…duces (follow-up to #908) (#919)

* npu: make the quantiser header's speed figure one that a re-run reproduces

#908's reviewer measured the same mix on this box and got 6.56x (1.539 ns/element
scalar, 0.235 AVX2) where the comment claimed 7.2-7.4x (1.72-1.76, 0.238). The
AVX2 side agrees; the scalar figure was ~12 % optimistic, so a later re-run would
look like a regression. The comment now carries the range both runs measured,
labelled as two independent measurements on the same box and mix.

Comment-only change: no code, no behaviour, no re-gate.

* npu: the quantiser comment's saving range must include both runs' endpoints

Copilot's review of #919: the comment said 6.6-7.4x is 'about 3.4-3.5 ms' saved, but
the slower scalar run (3.56) minus the faster AVX2 run (0.54) is 3.02 ms, so the
range excluded the second measurement. Now stated as the subtraction it is:
3.0-3.5 ms/token (3.56-0.54 to 4.05-0.55), i.e. 0.73-0.85 % of a ~415 ms token.

Comment-only; no code, no behaviour, no re-gate.

---------

Co-authored-by: agent <agent@local>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Review effort 4/5 tprs-ready TPRS: review verdict READY posted by the standing reviewer; merge gate satisfied

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant