Data types on SM86

The result in five numbers

KL(BF16 ‖ model) over the positions generated by the assistant, in 48 windows of real coding-agent traffic. On prompt text KL means nothing: the chat model predicts <|im_end|> at 99% in user turns and any perturbation flips it. Prefill with pp.py (cold prompt, max_tokens 1). On Sep 28 the Sep 25 checkpoint was compared against the current one at the same time: 2570 / 2180 vs 2590 / 2180 tok/s, the same. On Oct 10, with the same checkpoint and today's stack (dynamic prefill chunk, 8 GiB L2, PN174), it gave 2368 / 2014: −9% / −8%. The weights did not change, so it is the stack: still to be bisected, together with the −5% from Sep 25 to Sep 28. The decode step is the Sep 28 profiler measurement, not repeated.

Against BF16: the Qwen3.8-27B model card benchmarks

Every benchmark in the official card that can be run with an endpoint and a public dataset, against the served model (W4A8, int8 KV, DFlash2 tree, two 3090s). Official prompts and graders, the card's sampling (thinking, temperature 1.0, top_p 0.95, top_k 20), reasoning_effort=xhigh, one sample per item. Finished on Oct 10; HLE and LiveCodeBench were re-run with a 200k-token cap (the first run used 81,920).

This is not an even comparison. Reference model: the BF16 column is the number Qwen publishes for the official model. idiotSavant comes from orcarouter's abliterated BF16, so the difference mixes abliteration, quantization and the harness. Judge: HLE, CharXiv and BabyVision are graded by the same model, not GPT-4o as in the card. Samples: HLE is a fixed sample of 500 and MathVision is testmini. Suite: benchmarks/qwen38_card in genesis-vllm-patches.

What happened to each proposal from Sep 23

The silicon rule

On GA102 the accumulator type matters as much as the input type. An fp16 mma with an fp32 accumulator, the default in FlashInfer, cuBLAS and Triton, runs at half the rate of the same with an fp16 accumulator and at a quarter of int8. That gap matters where compute dominates (prefill). In decode memory dominates, and bytes are what count.

int8 activations are the condition for using the 4× row. The problem that came up is that with int8 per token the scale is set by each token's maximum, and in this model a single channel (3994, the massive activation) is 70 times larger than the rest. Rotation solves that without touching the kernel.

SM86 is not SM80: how to take advantage of it

Almost all attention and GEMM software (FlashAttention, FlashInfer, cuBLAS, the Triton and Marlin templates) is designed for the A100 (SM80). The 3090 (SM86, GA102) shares the instruction set and runs the same binary, but has different ratios between units. What is free on SM80 can cost double here, and vice versa. This is the record of those differences and how best to use them; rows are added as they are measured.

What changes, per SM and per cycle

SM86 measured with tests/proto/sm86_throughput.cu (Sep 27): results per cycle and SM with clock64. Tensor: the 3090 spec (71 / 142 / 284 / 568 dense TFLOPS or TOPS at 1.695 GHz over 82 SMs); int4 at 2× s8 is measured. A100: the board spec.

Rules for writing kernels here

1. Every fp16 mma with an fp32 accumulator is a candidate. Moving it to int8 gives 4×, and to an fp16 accumulator gives 2× if the sum is short (within a tile) and is then folded into fp32. Where it pays: prefill attention, the GDN chunk, the lm_head with many rows.
2. No integer ALU in the hot loop. It shares a path with float and runs at half width. Softmax and epilogues in float, with MUFU exp; integers only where they replace bytes (weights, KV).
3. Registers decide occupancy, and occupancy hides latency. Without wgmma, overlap between mma, softmax and loads only comes from having ≥8 warps per SM. With head_dim 256: fp16 accumulators per tile, split dimensions across warps and recompute what is cheap (QKᵀ int8) rather than synchronizing through shared memory.
4. Tiles for 100 KB and 1536 threads, not SM80's. Fewer cp.async stages and smaller tiles, and the grid tied to the hardware (warps per SM), not to the problem.
5. Decode is bytes. With 6 MB of L2 and ~640 GB/s, a decode kernel is measured in % of the DRAM ceiling. Marlin is at 98–103%; SK-18h at ~44% (35% before moving to 8 warps per block).
6. Few communications. Over PCIe each all-reduce costs ~25 µs even if it moves 92 KB: fusing or eliminating all-reduces is worth more than compressing them.
7. Measure on SM86 and with a fixed clock. Figures from papers (H100, A100) do not transfer: the exp that weighs 50% on H100 is ~7% of the pair here; the fp32 that is free on A100 costs double here.

Where it is already applied and where it is missing

Why rotation was needed: the A8 error per linear

Relative error of each linear's output with the activation in per-token int8, on real samples held out from calibration. It is the Marlin W4A8 computation, with the TP=2 partitioning simulated.

Whole layer: the A8 error was 13.6% of what the layer adds (median), the same size as the int4 weight error (~12%). In the whole model, W4A8 doubled the KL of W4A16.

Per-layer sensitivity, with and without rotation

KL on responses when a single layer is quantized and the other 63 stay in BF16. Each layer was measured twice: with int4 weights and fp16 activations (W4A16) and with int8 activations (W4A8). So each bar can be split in two: what the int4 weights cost (W4A16) and what int8 activations add (W4A8 − W4A16). That split is what explains why the model works well.

Why it works so well, in four readings

1. Unrotated, what broke things was not the weights but int8. The orange on top (0.0232) is larger than the blue (0.0149) and piles up in early attention: layers 3, 7, 11 and 15 account for 19% of that whole penalty. It is channel 3994, which sets each token's int8 scale and crushes the other 5119.
2. Rotated, the orange almost disappears. With a crest factor of 2.7 instead of 70, per-token int8 costs the same as fp16 (median 1.06×). In other words: idiotSavant computes in int8, at the 4× tensor-core rate, with the quality of a W4A16. The bars below are almost all blue.
3. Rotation also improves the weights. The blue drops 19%: GPTQ quantizes a matrix better without outlier columns, because the Hadamard spreads them in the weights too.
4. What remains is even and small. The 8 worst layers account for 21% of the total: there is no single layer to fix, which is why per-layer mixed precision pays little. In total, the served model reaches KL 0.0178 per token and the same top-1 in 95.6% of positions. That is on the order of the difference between two identical server starts (~0.019 measured in probes on Sep 14, with another model). A task decided with margin (multiple choice, a format constraint, a short answer) does not flip because of that, and that is why it ties with BF16 on IFBench, GPQA and ERQA.

Where it falls short, and what the chart says about it

There are two things the chart measures and two it does not. What it does measure:

• The visible exception is layer 0. Rotated, its orange (5.1 ×1e-4) is 35% of all the int8 error that remains. The massive activation is born in its MLP and its attention input has a crest factor of 3.8, not 2.7. Serving it in W4A16 is the cheapest improvement left: no recalibration needed and it removes ~4% of the KL. The cost is ~1% of prefill, because that layer loses the int8 rate; in decode, which reads weights, it is barely noticeable.
• Layer 63 is the most sensitive with and without rotation. It writes what the lm_head reads.

What it does not measure:

• Length. These are 1024-token windows. In the benchmarks, what fails is long reasoning: in HLE 16% of responses end without closing </think>, with a rate that goes from 0% below 8k tokens to 25% between 64k and 82k. With thousands of reasoning tokens, a small per-token error has more chances of derailing the trajectory.
• Calibration. GPTQ was calibrated with 4096-token windows and never saw a 50k GDN state.

The next step is to repeat this same chart on 32–80k-token trajectories. It is done by teacher-forcing the BF16, layer by layer from the SSD, with the tokens idiotSavant generated on HLE, without generating anything. That tells whether the profile stays even at long context or whether culprit layers appear. And, in the responses where idiotSavant did not close the reasoning, whether the BF16 did emit </think>.

Measured on Sep 25, before rotating o_proj/out_proj online (today's served model also has that per-head Hadamard, which lowered the served KL from 0.0193 to 0.0178). 4 windows of 1024 response tokens of real agent traffic, exact KL over the full vocabulary, torch. The sum of the single-layer KLs (0.0136) is close to the KL of the whole model in torch (0.0116): errors from different layers add up with almost no compounding.

What the build says, layer by layer

The script that rebuilds the model leaves a report for each layer (informe_idiotsavant.json, which ships with the model). These are its three readings. The crest factor and the local error are from the Sep 25 build; the Sep 28 update redid only o_proj, out_proj and in_proj_a/b (idiotsavant.py realinear), and the weights table already includes them.

The crest factor of residual-stream inputs

Max over rms per token, at the median of calibration tokens, for the attention input and the MLP input of each layer. Unrotated it starts at ~71 and decreases with depth; rotated it stays between 2.6 and 4.1 in all of them. Layer 0 is the exception at the attention input (6.7): the massive activation is only created by its MLP.

The local error of each served layer

Error of the layer's output already in its servable form (rotated int4 weights, folded norms, Hadamard in down_proj) on the samples held out from calibration, relative to what the layer adds to the residual stream. It is the weight error; the A8 error is measured separately. Median 10.3% in the GatedDeltaNet layers and 9.2% in the attention layers.

The local error guides but does not decide: its correlation with the KL each layer causes is 0.6. Layer 0 has the lowest error (0.9%) because what it adds to the residual stream is large compared to its error; layer 63 has 5.3% error and is the most sensitive in the model.

The int4 weight error, by linear type

Relative Frobenius error between the rotated weight and its quantization (symmetric int4 g128, GPTQ). It is even, because rotation spreads outliers in the weights too. The ones that write the residual stream are slightly higher, and layer 0 is the worst in qkv and down_proj. A single dead channel in the whole model: layer 7, MLP input, channel 3994. With the Hadamard at the input, the Frobenius error of o_proj/out_proj rises slightly (13.6/13.3 → 13.8/13.9%): GPTQ minimizes the Hessian-weighted error, not this one, and the model's KL goes down.

Whole-model quality

KL(BF16 ‖ model), top-20, response positions only. Torch: 4 windows. vLLM: 48 windows, with DFlash2 and the tree enabled. The rebuilt model shares 98.6% of its int4 values with the research build. The Sep 28 update with the Hadamard in torch gave 0.0176; served (SK-25 fused) gives 0.0178. Drafter acceptance under real agent load does not change: 5.546 before and 5.553 after (40 requests, 4 in parallel, 2 server starts).

The pipeline, operation by operation

The data types are today's (Sep 28). The percentages come from the Sep 20 and 21 traces (with noon) and remain as a reference for relative weight; the current breakdown, with idiotSavant at 62k context, is further below. The green stripe marks the rows that changed since Sep 23.

* Estimated: SK-23 removes 1 of the 4 quantization sites in each layer, so ~¾ of what was measured on Sep 21 remains (2.7% and 1.3%). Marlin types according to the vLLM 0.29.0 ScalarType: Marlin<int8, uint4b8, float16, float16> for the projections and <float16, uint4b8, float16, float16> for the lm_head.

The measured step, layer by layer and kernel by kernel

Two snapshots of the same stack (idiotSavant + DFlash2 + 8-node tree, production compose in an isolated instance, no outside traffic), with the torch profiler on both ranks: the Sep 25 baseline (traces is25_*) and the Sep 28 state (is27_v2u_*), after steps 1–9 of Improvements. Three loads: one request with 62k context (9 rows per step: the root and 8 nodes), four requests of ~15k (36 rows) and a cold prefill of 15.4k in chunks of 1760. Each step is split into regions using SK-23 as the anchor: it runs once per target layer and each layer ends at the all-reduce that follows its down_proj. Both ranks give the same result within a 0.1% difference, so TP imbalance does not show up here.

Where the time goes

The target GEMMs are the floor and are already at the DRAM ceiling (gate_up reads 44.6 MB in 60.5 µs, 737 GB/s): on Sep 25 they were 30% of the step with one request and today 41%, because everything else shrank around them. What remains far from its ceiling is decode attention (20%, ~44% of bandwidth) and the drafter (14%). The Sep 28 bars come from summing kernels by category (the Sep 25 breakdown came from the region analyzer), so the comparison is approximate.

One layer, kernel by kernel (1 request, 9 rows), Sep 28

Decode attention reads 63 MB of int8 KV per layer and rank (62k × 2 heads × 256 × K and V). On Sep 25 it did so in 284 µs (35% of what the power-capped card delivers) with two blocks of 4 warps that loaded the same page; with one block of 8 warps, 224 µs (~44%). What holds it back now: 254 registers (one block per SM) and 880-token pages, which at 62k give 142 blocks for 82 SMs (1.7 waves). With 4 requests, sk18h_salida went from 84 to 10 µs per layer: it was doing an emulated 64-bit integer division for every element.

Outside the layers

Prefill, one chunk of 1760 (Sep 25)

A chunk takes 575–650 ms. The GEMMs are half and run at 90% of int8 peak at the 810 MHz of the 220 W cap (gate_up: 314 GOP in 2.55 ms, 123 TOPS). The int8 all-reduce is 22%: two per layer of ~0.95 ms each, 9 MB over PCIe at ~10.5 GB/s, not overlapped with anything. Prefill attention (FlashInfer fp16 with f32 accumulator, half rate on GA102) is 11.6% at 15k, plus 0.5 ms per layer of idle GPU in its preparation (sort and unique on the host).

In mixed steps (one request in prefill while others decode), sk18h_decuant dequantizes the KV to fp16 before FlashInfer: 0.6–1.3 ms per attention layer, 10–20 ms per step. In a prefill alone it is 0.14 ms. This is the opencode pattern: subagents that come in while the main thread generates. Since Sep 27 the requests that decode in a mixed step go through SK-18h and only the prefill dequantizes its pages: −63% dequantization.

The drafter: qwen3.8_27b_idiotSavant_sm_86_dflash2

It is Inco's DFlash2 (5 layers, block of 8, top-16 selector for the tree) fine-tuned directly against the served idiotSavant, with no intermediate models, and quantized to W4A16 (1.2 GB). It cannot be rotated entirely: its k_proj/v_proj serves two streams with different norms, the target's features and the block's tokens. So it stays in its original basis and is converted only where it touches the target's basis.

Where margin remains, and how much

The table was built from the Sep 25 profile. Each row is a change that does not touch the model's numbers; the saving is the measured time minus the kernel's ceiling, so it is an upper bound. The estimate was that, if everything worked out, the one-request step would go from 29.7 to ~19.6 ms and the four-request step from 41.2 to ~25.6 ms, and that in practice half of that would materialize. By Sep 28 levers 1 (partially), 2, 3, 4 and 6 had landed: 22.0 and 30.0 ms, more than the realistic expectation. The status column says what is left of each one.

Possible improvements, in order

A single ordering that mixes two kinds of change. Code changes give the same output, bit for bit. Layer changes alter how a linear is rotated or quantized, and they are now possible because the rotated residual stream has a crest factor of ~2.7 instead of ~56. Everything measured with noon is measured again on idiotSavant. The criterion: cheap and safe first, then group the changes that touch the same kernel, and the layer changes in batches, because each batch pays for a quality validation.

The KL values shown in each step are those from when that step was measured, almost all with the Sep 25 checkpoint (0.0193). The current one, since batch B on Sep 28, is 0.0178.

Prefill attention: why softmax in float and not integer

SK-18's all-integer path does, for each (query, key) pair, the exp via shift and quadratic, the weight split into two bytes and two mma passes for P·V. That is ~15 operations on the INT32 pipe, which on GA102 is half as wide as the FP32 one. The hardware exp (MUFU) is a separate unit that runs in parallel with the tensor cores. With head_dim 256 the exp is not the bottleneck, as it is on H100 with 128: what dominates is the mma type.

Expected 1.6–2× in attention (Sage measures ~2× over FlashAttention2 on 3090/4090), and without dequantizing. Measured throughput and rules in SM86 is not SM80. The same applies to decode: SK-18h uses the integer softmax and sits at 254 registers and 4 warps per SM. Moving it to float is a variant of step 5. Literature: SageAttention (arXiv 2410.02367), SageAttention2 (2411.10958), FlashAttention-3 (2407.08608), FlashAttention-4 (2603.05451).

The server: one crash, three chained problems and the empty responses

During the benchmark run (Oct 10), the engine crashed once and could not recover on its own. It was three different things, one after another. The fourth row is from Oct 11 and is not a crash: responses that arrived empty.

PN174 does the unlink from the scheduler, which is the last process to open the region (8 s after the workers). If it were done earlier, the scheduler would create another empty file and work on different memory. Verified: all three processes map the (deleted) file and speed did not change (~130 tok/s of prose with one request). The container went down from 21 to 18 GB.