01 · LOAD
10-byte block
Eight packed code bytes and two UE4M3 scale bytes arrive from memory.
q0q1…e0/e1
Current answer
Yes—FPX-IFP2 has a place, but as a speed-and-capacity tier on the bulk expert path, not as a blanket every-tensor recipe.
On HY3, FPX-IFP2 Speed‑v1—a ROCmFP2-dominant artifact—is 2.55% smaller than IQ2_M and cuts end-to-end no-MTP latency by about 2.4% on a short code prompt and 6.2% on a 19,654-token code prompt. The unweighted control scored 81 on HermesAgent‑20 and 86 on full Tool‑Eval; the unchanged imatrix edition preserved targeted quality at 425/425 and raised full Tool‑Eval to 88. With the completed disk-only target+draft cache, the imatrix model finished HA‑20 at 80—one point from the control and within the run-to-run variance observed here—while holding 14.05 GB of reusable state on SSD instead of reserving it in scarce UMA. A larger FP3-heavy candidate was slower and scored worse on the discriminating tests.
ROCmFP2 block
10 B
32 weights · 2.50 BPW
Artifact
90.76 GiB
2.55% below IQ2_M
Control HA‑20
81
1,457.29 s wall
Imatrix Tool‑Eval
88
122/138 · control 86
SSD-cache HA‑20
80
38:16.652 · tie-level result
Disk / max RSS
14.05 / 2.00
GB cache payload / GiB process RSS
Naming clarification: FPX-IFP2 is the public model and recipe lane. ROCmFP2 is the compact two-bit storage and codebook format underneath it. The hardware does not execute native FP2 floating-point math; the runtime turns stored selectors into exact signed int8 values and uses AMD integer dot products.
01 · Storage
The whole ROCmFP2 block fits in ten bytes
A block covers 32 original weights. It is split into two 16-weight halves so each half can choose its own scale. Every group of four weights becomes one byte: four 2-bit selectors packed little-endian.
Serialized layout
q0weights 0–3
q1weights 4–7
q2weights 8–11
q3weights 12–15
q4weights 16–19
q5weights 20–23
q6weights 24–27
q7weights 28–31
e0scale 0–15
e1scale 16–31
8 code bytes × 8 bits = 64 bits
2 scale bytes × 8 bits = 16 bits
total = 80 bits / 32 weights = 2.50 BPW
Why metadata matters at two bits
Two raw bits can name only four states. ROCmFP2 spends the remaining 0.50 bit/weight on two shared scales. That lets the same four integers represent tiny or large weights without storing a floating-point value per weight.
Exact serialized block sizes. FP3 is 12 code bytes + 2 scale bytes; BF16 is 64 data bytes.
02 · Quantization
The codebook numbers are intentionally simple
The frozen S40/MORD order is not a 256-entry table and it has no zero. It is exactly four signed integer reconstruction levels. Sign chooses the side; normalized magnitude chooses inner or outer at the midpoint 2.5.
00₂
−4
negative outer01₂
−1
negative inner10₂
+1
positive inner11₂
+4
positive outernormalized magnitude = |x / scale|
inner = 1 when magnitude ≤ 2.5
outer = 4 when magnitude > 2.5
reconstructed weight = codebook[code] × scale
Concrete byte: scale metadata 0x38 decodes to s = 0.5. Source values [−1.8, −0.4, +0.7, +2.1] normalize to [−3.6, −0.8, +1.4, +4.2], select [−4,−1,+1,+4], and pack as 0xE4. Reconstruction is [−2.0,−0.5,+0.5,+2.0].
One scale byte, 127 legal values
The scale metadata is unsigned UE4M3: four exponent bits and three mantissa bits. Bytes 0x00…0x7E are valid; 0x7F is rejected. The positive range runs from 2⁻¹⁰ = 0.0009765625 to 224.
E = byte >> 3; M = byte & 7
E = 0 : scale = M × 2⁻¹⁰
E > 0 : scale = (8 + M) × 2^(E − 11)
How a block is chosen: start near max(|x|)/4, test legal UE4M3 scales, map all 16 values to the four codebook levels, and keep the scale with the lowest squared reconstruction error. With an importance matrix, that same error is weighted so mistakes on frequently activated channels cost more.
03 · Runtime
From scale metadata to int8—and into AMD’s dot unit
The fast path never expands FP2 weights to FP16 one by one. It expands each stored byte directly into four signed int8 lanes, performs four multiply-adds together against Q8_1 activations, accumulates in int32, then applies the two scales.
02 · SPREAD
2-bit selectors
Bit-plane arithmetic spreads four adjacent 2-bit codes into four byte selectors.
00011011
03 · PERMUTE
Four int8 lanes
v_perm_b32 selects bytes from the packed constant table {−4,−1,+1,+4}.
−4−1+1+4
04 · DOT4
Integer MAC
RDNA3/RDNA4 uses __builtin_amdgcn_sudot4: four signed 8-bit products plus an int32 accumulator.
w₀×a₀w₁×a₁w₂×a₂w₃×a₃
05 · RESCALE
Float result
Each half’s integer sum is multiplied by its decoded weight scale and the Q8_1 activation scale.
dₐ × (d₀·Σ₀ + d₁·Σ₁)
Why this shape fits AMD
The codebook is a four-byte constant, not a large irregular lookup grid. One GPU byte-permute produces exactly the packed signed-byte shape consumed by AMD’s native four-lane integer dot instruction. That keeps the arithmetic dense while the model traffic stays at 2.5 BPW.
The key is separation of concerns: the codebook carries shape, UE4M3 carries range, int8 dot4 carries compute, and scaling happens after accumulation.
Why “AMD-friendly” does not mean “always faster”
At only two bits, unpack cost is large relative to payload. The first kernel expanded FP2 separately for every MTP verification column, leaving the dot unit underfed. The optimized multi-column path expands a target block once and reuses it across all Q8_1 columns.
That change is FP2-only and compile-time opt-in; other quants keep their existing kernel paths.
Measured expert-kernel recovery
Clean paired average, exact HY3 expert shapes. Lower is better.
What the optimization actually removed
before: for each MTP column
expand FP2 byte → int8×4
decode FP2 scale
dot against one activation column
after: expand FP2 byte → int8×4 once
decode FP2 scale once
reuse against every activation column
Paired average
−23.8%
66.52 → 50.72 μs
Correctness
8 / 8
1, 4, 32, 128 tokens
04 · Measured evidence
What we have found on HY3 so far
FPX-IFP2 Speed‑v1 puts ROCmFP2 on all 237 trunk routed-expert tensors, keeps attention/shared tensors at higher precision, and leaves the three MTP expert tensors at FP3. That is deliberately a ROCmFP2-dominant recipe—not a claim that every tensor tolerates two bits.
Artifact footprint
Five-shard realized GGUF sizes. Smaller is better.
No-MTP total latency
128 generated tokens. Log axis keeps both prompt lengths legible.
Quality gates
Control HA‑20 raw score 81.25; the imatrix full Tool‑Eval result is 88; disk-only cache-on HA‑20 scored 80. The one-point HA spread is treated as run variance, not a quality failure.
Prefill wins; raw decode is almost tied
The end-to-end win comes from faster prefill and lower model traffic. Raw no-MTP TG trails IQ2_M by about 2.1% short and 1.5% long.
MTP must be tuned on real workloads
All five bars scored 100/100; the chart isolates wall time. FPX-IFP2 n=2 was selected.
Prompt sensitivity is real
On the short code prompt, n=2 MTP lowers total time to 8.979 s. On the 19,654-token prompt it raises total time to 166.333 s even though decode itself jumps to 7.865 TG, because draft-prefill overhead dominates. This is why tiny prompts cannot select an MTP config alone.
Quality and wall-time results
| Gate | Score | Wall time | Read |
|---|---|---|---|
| HermesAgent‑08 · control | 100 / 100 | 62.446 s | n=2 MTP selector |
| Tool‑Eval core15 · control | 100 / 100 | 175.3 s | 21.2% faster than IQ2_M |
| HermesAgent‑20 · control | 81 / 100 | 1,457.29 s | 81.25 raw; strong promotion result |
| Tool‑Eval full69 · control | 86 / 100 | 1,697.287 s | 119/138 · 57 pass · 5 partial · 7 fail |
| Tool‑Eval full69 · imatrix | 88 / 100 | 1,885.736 s | 122/138 · full-suite winner; earlier 8 GiB RAM-cache runner |
| HermesAgent‑20 · imatrix + SSD | 80 / 100 | 2,296.652 s | 38:16.652 · completed cache-on · tie-level vs 81 control |
FP3 did not buy a quality tier: Quality‑v1 was 13.27% larger, HA‑08 slowed to 75.306 s, targeted Hermes scenarios fell from 425 to 405 raw points, and Tool core15 fell from 100 to 97 while slowing to 205.3 s. The failed 100+ GiB artifact was deleted rather than archived.
05 · Completed experiment
Can importance weighting improve FP2 without changing its size?
Coverage
237 / 237
target trunk FP2 tensors
No-MTP total
4.3–4.8%
faster on both measured prompts
Full Tool‑Eval
88 > 86
122/138 vs 119/138 control
What the matrix changes
It does not alter the ten-byte block, the codebook, or the AMD kernel. It changes only which UE4M3 scale wins the offline error search: important channels contribute more to the objective.
plain objective: Σ (xᵢ − x̂ᵢ)²
weighted objective: Σ importanceᵢ · √(σ² + xᵢ²) · (xᵢ − x̂ᵢ)²
Constructed implementation test: weighted error 0.510835 → 9.13573×10⁻⁵. This proves the weighting path works; it does not predict model-level quality by itself.
Gate and operational status
| Loader / coherence / tool | PASS · 24.56 s · 323 · native tool JSON |
|---|---|
| Tensor payload/layout | PASS · identical; +224-byte provenance header |
| No-MTP parity | PASS · no repeatable regression |
| MTP long prompt | 6.521 TG · 20/28 accepted · total +0.75% |
| Quality screen | HA‑08 100 · core15 100 · targeted 425/425 |
| Full Tool‑Eval69 | 88/100 · 122/138 · beats control 86/100 |
| Full69 cache · 8 GiB RAM | completed · 11 updates · 2 better-cache hits · largest update 173.03 ms · earlier pre-final-boundary runner |
| HA20 cache · 8 GiB RAM | NO SCORE · 13 saves / 0 recovered entries · aborted at HA‑13 · 7,657.434 MiB + next 348.502 MiB state triggered host OOM |
| HA20 cache · 4 GiB | CANCELED · eval not run; 4 GiB UMA cost is unjustified at the measured hit rate |
| HA20 · disk-only SSD | PASS · 80/100 · 38:16.652 wall · 185.42 PP · 17.07 TG · no OOM |
| SSD target+draft restore | PASS · exact and extension hits · MTP acceptance and TG preserved · suffix-only prefill |
| Forced context shift | PASS · 300/300 generated · 122/129 drafts accepted · 16.83 TG |
| Disposition | PASS · promoted quality/no-MTP FPX-IFP2 lane with operational SSD cache |
The imatrix applies to every covered quantized tensor in the unchanged recipe, including protected higher-precision tensors—not only FP2. The MTP block lacks matrix entries and uses the existing unweighted path. The +224 bytes are four provenance fields and alignment; all 1,298 tensor types, payload lengths, and capacities match the control.
The targeted tie was preservation—not a stop signal. HA‑08 remained 100, Tool‑Eval core15 remained 100, and the eight targeted Hermes scenarios reproduced the control’s exact 425 raw points. The full 69-case operational run then exposed the broader gain: 88/100 and 122/138, ahead of the unweighted control’s 86/100 and 119/138. The imatrix build therefore preserves the FPX-IFP2 profile’s existing advantages and adds a measured full-suite quality win.
The RAM-cache OOM identified the wrong storage tier; the SSD implementation resolves it. The old HA20 attempt placed target+draft state in UMA, reached 7,657.434 MiB across 12 states, then OOM-killed the server on the next 348.502 MiB allocation. The completed disk-only run stored 14.05 GB across 21 entries while process RSS peaked at 2.00 GiB, finished all 20 scenarios, and scored 80 versus the 81 control. That one-point spread is ordinary eval variance in this evidence set, not a cache failure. The earlier Tool‑Eval 88 result used the 8 GiB RAM cache and predates the final MTP boundary patch; it remains quality evidence, while the final HA20 and focused restore smoke are the operational SSD evidence.
The imatrix build was faster in four of five comparable timing views
Positive means lower wall time. It improved both no-MTP prompts, short-prompt MTP, and the targeted Hermes scenario sum. The one exception was long-prompt MTP: total latency was 0.75% slower and decode acceptance fell, so that lane remains the control build’s strength. These are observed A/B results; repeat them before treating the exact percentages as universal.
06 · Completed runtime proof
SSD-backed target+draft cache makes the operational profile viable
The patched ROCmFPX runner streams the target KV state, MTP draft state, and exact speculative boundary metadata to a bounded directory on fast SSD. It keeps no automatic prompt states in RAM, validates a complete pair before reuse, and cold-falls back if the target, draft, or MTP boundary cannot be accepted together.
SSD-cache HA‑20
80 / 100
control 81 · within observed variance
Full eval wall
38:16.652
2,296.652 seconds · all 20 complete
Engine throughput
185.42 / 17.07
PP / TG tok·s⁻¹ over the full run
SSD payload
14.05 GB
21 entries · 79,675 cached tokens
Process RSS ceiling
2.00 GiB
0.895 GiB minimum · no host OOM
Cache integrity
21 / 0 / 0
saves / evictions / save failures
Cache-hit time to first token
Two natural cold repeats are shown rather than selecting one baseline. Exact reuse processed one replay token; extension reuse prefetched only the 15-token suffix.
Focused target+draft restore checks
| Exact hit | 0.264 s TTFP · 1 prompt token · 1,362 cached · 13/13 MTP accepted · 13.87 TG |
|---|---|
| Extension hit | 0.478 s TTFP · 15 prompt tokens · 1,426 cached · 18/19 MTP accepted · 15.81 TG |
| Disk load | 59.8 ms exact · 88.3 ms extension |
| First token / boundary | matched cold · exact one-token replay · extension suffix-only prefill |
| Forced context shift | 1,536 ctx · 767 discarded · 300/300 generated · 122/129 accepted · 16.83 TG |
Cold HY3 repeats themselves diverged at token 27 for the exact prompt and token 1 for the extension prompt. The cache runs matched the first cold token and did not diverge earlier than natural cold-repeat variance, so full-stream bit equality was not used as a false failure gate.
What HA20 proves—and what the smoke proves
HA20 contains independent scenarios, so its 21 saved states had no eligible reuse and the run correctly recorded zero loads. That makes it a conservative cache-on stability, memory, throughput, and quality test. The exact and extension smoke deliberately revisited matching prefixes; those runs prove the speed benefit and target+draft restoration path.
Both values are plotted in GiB. SSD payload is reusable prompt state; RSS is the server process working set, not an extra cache reservation.
Model-card configuration
Users already need the pinned ROCmFPX runner for the FPX-IFP2 codebook. The SSD tier is therefore a normal profile setting, not another model conversion step:
MODEL_CONTEXT=65536
NO_MMAP=1
NO_CONTEXT_SHIFT=1
SPEC_TYPE=draft-mtp
SPEC_DRAFT_N_MAX=2
SPEC_DRAFT_P_MIN=0.75
CACHE_TYPE_K=q8_0
CACHE_TYPE_V=q8_0
CACHE_TYPE_K_DRAFT=q8_0
CACHE_TYPE_V_DRAFT=q8_0
CACHE_RAM=0
CACHE_DISK=/path/on/fast/ssd
CACHE_DISK_LIMIT=16384
Equivalent server flags are --cache-ram 0 --cache-disk PATH --cache-disk-limit 16384. The promoted 65k profile keeps context shift disabled; the shift-aware MTP metadata hook was separately forced at 1,536 tokens and passed. Cache entries are private to the live server run and are cleaned up at shutdown.
07 · Recipe decision
Where FPX-IFP2 belongs today
Use it for
Very large MoE expert tensors where memory traffic and prefill dominate, the codebook kernel has a tuned launch policy, and higher precision protects attention, embeddings, shared experts, and other fragile tensors.
Do not assume
That two bits automatically beat IQ2_M in raw token generation, that more FP3 automatically improves quality, or that one MTP n_max wins on both short and long prompts.
Best next move
Promote the imatrix edition as the quality/no-MTP and short-workload FPX-IFP2 lane: it scored 88 on full Tool‑Eval and was faster in four of five timing views. Keep the original build for long-generation MTP. Ship the validated disk-only target+draft cache in the pinned runner profile instead of reserving 4–8 GiB of UMA.
Evidence ledger
Local sources behind the report
| Claim family | Primary evidence | Status |
|---|---|---|
| Block layout, UE4M3, S40 codebook | src/ROCmFPX-hy3/ggml/rocmfpx/rocmfpx.hsrc/ROCmFPX-hy3/ggml/rocmfpx/rocmfpx.c | source-audited |
| AMD byte permute and dot4 path | src/ROCmFPX-hy3/ggml/src/ggml-cuda/vecdotq.cuhsrc/ROCmFPX-hy3/ggml/src/ggml-cuda/common.cuhsrc/ROCmFPX-hy3/ggml/src/ggml-cuda/mmvq.cu | source-audited |
| Kernel sweep and exact correctness | results/hy3-adaptive-rpb.tsvresults/hy3-adaptive-q2-final-correctness.log | measured |
| No-MTP / MTP API speed | results/*-api-short/results.jsonlresults/*-api-long/results.jsonl | measured |
| HA20 and Tool-Eval | /srv/ssd/p3700ba/data/llm-benchmarking-lab/.benchlocal-home/.benchlocal/runs/hermesagent-20/20260713T121558Z-hermesagent20-hy3-chadrock-speed-v1-mtp-n2-clean-eval/summary.json/srv/ssd/p3700ba/data/llm-benchmarking-lab/runs/2026/07/2026-07-13T13-56-30.458189Z_713cc7e5.md | measured |
| Imatrix H4 validation | experiments/h4-imatrix/results/artifact-gate.mdexperiments/h4-imatrix/results/{no-mtp,mtp-n2}-{short,long}/results.jsonl20260713T174811Z-…-targeted8/summary.json2026-07-13T18-19-21.426469Z_db0e22bb.md | validated; full Tool‑Eval winner on earlier 8 GiB RAM-cache runner |
| RAM-cache boundary | experiments/h4-imatrix/results/imatrix-mtp-n2-ha20-cache8g-final-server.log20260713T190117Z-…-ha20-cache8g-final/summary.jsonexperiments/h4-imatrix/results/imatrix-mtp-n2-ha20-cache4g-final-server.log | 8 GiB HA20 aborted; 4 GiB eval canceled |
| SSD exact / extension restore | experiments/h5-ssd-cache/results/live-mtp-cache-smoke-v3.jsonexperiments/h5-ssd-cache/results/server-live-smoke-v3.log | all focused checks passed |
| SSD cache-on HA20 | 20260713T210935Z-…-imatrix-mtp-n2-ssd-cache-final/summary.jsonexperiments/h5-ssd-cache/results/ha20-ssd-final/live-metrics.tsvexperiments/h5-ssd-cache/results/server-ha20-ssd-final.logexperiments/h5-ssd-cache/results/ha20-ssd-final/verify-store.json | 80/100; 38:16.652; store verified |
| MTP context shift | experiments/h5-ssd-cache/results/mtp-context-shift-enabled-1536.jsonexperiments/h5-ssd-cache/results/mtp-context-shift-reference-65536.jsonexperiments/h5-ssd-cache/results/server-mtp-context-shift-enabled.log | forced shift completed |
Benchmark host: AMD Ryzen AI Max+ 395 / Radeon 8060S unified-memory system, ROCm llama.cpp fork. Timings are local single-host observations, not claims about all AMD architectures. MTP quality tests use the selected n=2 / Q8 target+draft cache profile. The imatrix full Tool‑Eval result of 88 ran on the earlier 8 GiB RAM-cache revision and produced two better-cache hits across 11 updates. The final SSD implementation adds paired target+draft streaming, exact MTP boundary metadata, atomic size-validated state, fail-closed fallback, and context-shift rebasing. Its HA20 run completed at 80 with zero eligible loads because the scenarios are independent; exact and extension reuse were validated separately by the focused smoke. The 80 versus 81 HA20 spread is treated as ordinary run variance, and no additional reruns were used to fish for a preferred score.