Repository navigation
Measure and remove the draft attention layout-copy dispatch #754
Description
Each DFlash layer reduces attention into grouped query order and then copies it into projection order. Write the same normalized BF16 values directly into the existing projection-input buffer. This removes one reorder dispatch per layer and preserves the reduction order.
Inspired by LMSYS: Muse Spark Open Ultrafast.
What the follow-up established
- The original apparent 27B slowdown was not reproducible in the controlled comparison below. The unchanged baseline itself drifted 10.5–18.6% between its opening and closing processes; the candidate differences were smaller than that variation. Those original inconclusive measurements remain in the history section.
- Reusing one loaded model, allocator and pipeline cache, holding inputs/widths fixed, and alternating ABBA/BAAB gives aggregate spread at most 0.68% across the TTFT, prefill-GPU and decode metrics. Every aggregate comparison passes the repository's unchanged non-regression rule. This is a diagnostic result, not a bare-production release-gate certification or a general serving-speed claim.
- Every cold first-token request in this diagnostic executes zero changed draft-attention kernels, in both modes. This decode-stage change does not shorten that cold TTFT path.
- The replaced stage is 1.49–1.58× faster with byte-identical outputs. Its contribution to the complete model is small; the estimated budget below explains why this does not imply a comparable model tokens/sec gain.
The follow-up commit adds a permanent, model-free stage benchmark and tests the actual prepare → attention sequence while reusing consumed QKV storage. Tests verify dense row order, preservation of prepared inputs and an untouched allocation tail at every tested hidden width and B1–B4.
Related work: PR #702 already proposes in-place row-major reduction. This PR is an alternative with an explicit output buffer, reusing proposalQkv. The implementations must not be stacked. No performance superiority over #702 is claimed. It is independent of CommandGraph PR #757.
Current-commit validation
Candidate c4038a1a; original baseline d8af1054. Apple M5 Pro, 20 GPU cores, 48 GB; macOS 26.6, Apple clang 21, target macOS 26.4 / Metal 4.0.
- Production, native benchmark and stage-benchmark builds,
check-source, and all 11 build-identity tests pass. draft-attention,dflash-batch-control,draft-selectorand draft-preparation goldens pass with shader validation. The new QKV-reuse test includes every KV head and inactive NaN rings at empty history.- The stage benchmark's baseline/current BF16 outputs match exactly at B1–B4, both under shader validation and in its separate timing run.
- Both installed models pass the runtime oracle in FP32-state, BF16-state/two-spill-buffer, and 32K-spill/window-64 variants: six exact-head runs.
- The attention-only adapter passes shader-validated native Engine runs on both models, then all output/cache/acceptance checks in its timing run. Old-mode traces contain exactly one restored reorder per reduction; fused-mode traces contain none.
Hosted CI:
- Production, CPU runtime, native builds and the app: failure.
- Native sanitizers: success.
- Python compatibility 3.12: success.
- Python compatibility 3.13: success.
- Python compatibility 3.14: success.
The production/CPU/native-build/app job stopped before compilation: Apple's MetalToolchain asset catalog for RequestedBuild = 27A266a could not be fetched (xcodebuild exit 70). The original head passed all five hosted jobs. This new setup failure needs a maintainer rerun and is not reported as a code/test failure or a pass. Rerunning the failed job was denied: Must have admin rights to Repository.
Reproducible stage measurement
| Lanes | Reduce + reorder (µs) | Fused reduce (µs) | Paired stage speedup |
|---|---|---|---|
| B1 | 3.965 | 2.514 | 1.577× |
| B2 | 4.016 | 2.539 | 1.581× |
| B3 | 4.917 | 3.289 | 1.495× |
| B4 | 5.073 | 3.347 | 1.516× |
make benchmark-draft-attention BASELINE=../splash-baseline runs the committed benchmark against a built pre-fusion checkout. It reports all 20 ABBA rounds after four discarded warmups, with 256 repeated stages/command. The timing run has shader validation disabled; a separate validation run passes exact parity. Time columns are marginal medians and speedup is the paired median. It measures reduction plus layout transfer on synthetic partials, excluding the attention split pass, projections, MLPs and target verification.
A budget illustration using B1 stage savings and the original native baseline's GPU cycles:
- Qwen3.6-35B-A3B-4bit: 6 layers × the isolated-stage saving ≈ 8.7 µs, against a measured native GPU cycle of 18.3 ms (about 0.05%).
- Qwen3.8-27B-4bit: 5 layers × the isolated-stage saving ≈ 7.3 µs, against a measured native GPU cycle of 71.3 ms (about 0.01%).
These isolated timings are not in-situ full-graph savings or a hard ceiling: dependency/cache effects can change them. They show the scale of this component relative to the complete decode cycle.
Controlled whole-model diagnostic: TTFT and aggregate decode tokens/sec
Cold TTFT
| Model | Cold prompt | Original path (ms) | Fused path (ms) | Change | Aggregate spread | Verdict |
|---|---|---|---|---|---|---|
| Qwen3.6-35B-A3B-4bit | 512 | 165.381 | 165.017 | -0.22% | 0.68% | pass |
| Qwen3.6-35B-A3B-4bit | 2016 | 528.351 | 527.921 | -0.08% | 0.12% | pass |
| Qwen3.8-27B-4bit | 512 | 711.090 | 713.142 | +0.29% | 0.25% | pass |
| Qwen3.8-27B-4bit | 2016 | 2949.810 | 2945.895 | -0.13% | 0.66% | pass |
Aggregate decode output tokens/sec
| Model | Lanes | Original path (tok/s) | Fused path (tok/s) | Change | Aggregate spread | Verdict |
|---|---|---|---|---|---|---|
| Qwen3.6-35B-A3B-4bit | B1 | 308.41 | 308.55 | +0.05% | 0.16% | pass |
| Qwen3.6-35B-A3B-4bit | B2 | 553.11 | 553.11 | -0.00% | 0.14% | pass |
| Qwen3.6-35B-A3B-4bit | B3 | 628.56 | 630.48 | +0.31% | 0.35% | pass |
| Qwen3.6-35B-A3B-4bit | B4 | 788.93 | 789.73 | +0.10% | 0.09% | pass |
| Qwen3.8-27B-4bit | B1 | 96.01 | 96.16 | +0.16% | 0.13% | pass |
| Qwen3.8-27B-4bit | B2 | 182.42 | 182.75 | +0.18% | 0.03% | pass |
| Qwen3.8-27B-4bit | B3 | 212.14 | 212.33 | +0.09% | 0.03% | pass |
| Qwen3.8-27B-4bit | B4 | 281.53 | 280.79 | -0.26% | 0.45% | pass |
For each model, a single native Engine/backend/model remains loaded. Only the changed reduction and the old reorder use the original shader library in control samples; all other kernels, graph storage, weights, allocator and cached pipelines are shared. A command-copy adapter runs in both modes. Controls restore the original grouped reduction followed by reorder, while the candidate uses the fused output binding. The adapter refuses reduction commands with host-event positions, which these resident/GPU-only workloads do not use.
The native benchmark is restricted to two cold prompt bands and uses the identical input tokens for every repetition. Each decode width runs contiguously, with 64 decode output tokens per lane and every draft row verified. Modes alternate ABBA and BAAB. Forty samples per width/band are recorded; samples 0–7 are excluded only from timing, leaving eight balanced blocks. All samples remain in correctness checks. TTFT is the first emitted token's wall time; decode TPS is 1000 / ms per decode output token, using the raw native decode accounting.
For aggregation, each block is normalized to [old, fused, fused, old]; each of four positions contributes eight values to dev/benchmarks/abba.py. Its rule is unchanged: spread above 5% is inconclusive, otherwise the allowed regression is max(2%, 2 × spread). A pass does not establish speedup. Individual blocks are also retained: 179 pass, 12 inconclusive, 1 fail across 192 wall/GPU metric-block comparisons. One first-run 35B B4 GPU block flagged +2.08% against 2%; the rest of that scope aggregates to a pass. No block is dropped or relabeled.
An independent 35B decode-only repeat uses the same binary and eight fresh balanced blocks. B4 aggregate GPU change: -0.06%, spread 0.08%, pass; its block verdicts are {'pass': 8}. B4 wall change: -0.04%, spread 0.07%, pass. The first-run flagged block remains in the raw record below.
This diagnostic resolves the broad run-to-run confounder on this machine; it does not turn the original multi-process inconclusive scopes into production performance passes. Only this M5 Pro and the two installed MLX affine-4bit models were measured. No hardware-matrix or HTTP throughput qualification is claimed.
Raw stage rounds, interleaved block timings, parity certificates and build provenance
The normalized blocks preserve both arm values; their exact chronological order alternates ABBA/BAAB. The official aggregate comparison uses the median of each normalized position. Raw JSON hashes identify the retained complete files, and the values below allow all timing verdicts to be recalculated.
{"source_commit":"c4038a1a520bc490b5ecec7e6440c8252b3fd9ba","probe_binary_sha256":"0a500ebbf4314f45f0b7f94045161d25d82c61316aa51fe175391c6d5369624c","current_metallib_sha256":"a6403a17ac939193909f2e227abcce3f0e2eefdf9ac93040e0496d23fd9971f8","original_metallib_sha256":"7006f53dedd016809a16f4e5694c78e4f8959f31691e4473a35ad46be45b9b03","stage_benchmark_binary_sha256":"20c4e404ed3617d485b116915dd52e1c617c72684ebe623e58a2b6ce394325a5","stage":{"repeats_per_command":256,"warmup_rounds":4,"measured_rounds":20,"results":[{"lanes":1,"baseline_us":3.964925,"current_us":2.51416,"paired_speedup":1.576838,"exact_output_bytes":true,"rounds_abba_us":[[3.967773,2.50651,2.520345,3.947103],[3.941244,2.494141,2.510091,3.946289],[3.95752,2.494141,3.718425,3.975423],[3.967285,2.508463,2.517578,3.966471],[3.959798,2.515788,2.516113,3.969564],[3.974447,2.510254,2.522461,3.962402],[3.967448,2.509928,2.525391,3.962891],[3.973796,2.501465,2.519531,3.947266],[3.948568,2.582194,2.521321,3.964356],[3.9764,2.510742,2.520508,3.949707],[5.154297,2.514648,2.521973,3.962402],[3.966797,2.513184,2.520996,3.970215],[3.964356,2.508789,2.520345,3.958659],[3.967122,2.513184,2.499674,3.953776],[3.971517,2.497233,2.500977,3.960449],[3.961263,2.501953,2.508301,3.961263],[3.964844,2.509766,2.517741,3.964193],[3.965332,2.511719,3.697266,3.96582],[3.965983,2.508952,2.517578,3.971354],[3.971029,2.509603,2.516439,3.968262]]},{"lanes":2,"baseline_us":4.016073,"current_us":2.539469,"paired_speedup":1.581412,"exact_output_bytes":true,"rounds_abba_us":[[4.02181,2.542318,2.539388,4.023438],[4.01123,2.54362,2.544759,4.019043],[4.016113,2.542318,2.544271,4.019206],[4.01416,2.539713,2.540039,4.019531],[4.01709,2.544434,2.541829,4.025879],[4.018066,2.544434,2.54541,4.017578],[4.010742,2.519531,2.519531,3.985351],[5.200195,2.527344,2.516113,4.00651],[3.993164,2.537598,2.537272,4.018229],[4.02002,2.538574,2.545736,4.008952],[4.006348,2.529622,2.529785,4.002441],[4.005859,2.538411,2.528971,4.016601],[4.019368,2.536621,2.518555,3.985352],[4.000163,2.540039,2.537923,4.022949],[4.013672,2.518555,3.739746,4.03125],[4.026367,2.537598,2.537923,4.012207],[4.019531,2.539551,2.543132,4.017578],[4.014323,2.541016,2.537109,4.010742],[4.009765,2.538412,2.537435,4.020833],[4.01595,2.542155,2.540527,4.024414]]},{"lanes":3,"baseline_us":4.916504,"current_us":3.288981,"paired_speedup":1.494909,"exact_output_bytes":true,"rounds_abba_us":[[4.9139,3.294108,3.29362,4.921224],[4.911295,3.282715,3.287109,4.918294],[4.918457,3.279948,4.470703,4.933594],[4.900391,3.289063,3.288411,4.916667],[4.918783,3.262858,3.306641,4.900391],[4.915039,3.27832,3.28418,4.912109],[4.924805,3.289551,3.297363,4.918783],[4.907064,3.293945,3.287598,4.930501],[4.909342,3.288249,3.281087,6.136881],[4.919759,3.301921,3.283529,4.917969],[4.922363,3.288086,3.284831,4.924153],[4.930664,3.294922,3.287272,4.914551],[4.921875,3.294759,3.287923,4.910645],[4.917969,3.28418,3.295085,4.915527],[4.916341,3.301107,3.289876,6.121094],[4.901856,3.289063,3.289388,4.920736],[4.907064,3.274089,3.282389,4.895508],[4.924154,3.273112,3.27002,4.898926],[4.895508,3.285156,3.271484,4.912109],[4.901693,3.28776,3.27832,4.920899]]},{"lanes":4,"baseline_us":5.072998,"current_us":3.346598,"paired_speedup":1.51589,"exact_output_bytes":true,"rounds_abba_us":[[5.067871,3.344727,3.348795,5.056478],[5.082357,3.349284,3.333985,6.270996],[5.079264,3.355306,3.339518,5.067871],[5.076823,3.348796,3.344401,5.069824],[5.048665,3.329427,3.335286,5.046387],[5.044922,3.361328,3.343262,5.066895],[5.074382,3.349609,3.343587,5.07308],[5.053385,3.346191,3.356608,6.270996],[5.076335,3.361816,3.348307,5.068685],[5.078288,3.348144,3.34668,5.07373],[5.066895,3.339844,3.338867,5.056803],[5.076497,3.334961,3.330241,5.065918],[5.071777,3.345215,3.347656,5.069173],[5.074056,3.344401,3.348633,5.071289],[6.295085,3.354492,3.341471,5.069499],[5.073242,3.344238,3.353841,5.069824],[5.082357,3.350586,3.345703,5.082031],[5.080078,3.338379,3.34847,5.069987],[5.069824,3.346354,3.345215,5.081543],[5.067546,3.347656,3.343262,5.062174]]}]},"original_27b_same_binary_drift":[{"width":1,"same_baseline_initial_ms_per_token":9.47715625,"same_baseline_final_ms_per_token":10.679953125,"baseline_drift":0.12691537875615388,"candidate_change":-0.022675240494974824},{"width":2,"same_baseline_initial_ms_per_token":5.0469453125,"same_baseline_final_ms_per_token":5.5793671875,"baseline_drift":0.10549388630808543,"candidate_change":-0.01749864380995847},{"width":3,"same_baseline_initial_ms_per_token":4.250630208333333,"same_baseline_final_ms_per_token":5.0403125,"baseline_drift":0.18578004977203144,"candidate_change":-0.007936018009811252},{"width":4,"same_baseline_initial_ms_per_token":3.23077734375,"same_baseline_final_ms_per_token":3.81219140625,"baseline_drift":0.17996104362461152,"candidate_change":-0.01788719919731474}],"diagnostics":[{"model":"35b","loaded_layout_sha256":"964ce8afdb896d17505b4f1b322c0d81974ad8a38a16437c15dfd31fbc9bc377","raw_json_sha256":"efa961526debfa9790129f65fd582ea5ae5798e85bbf13bc8d316fe9a4014e4c","thermal_states":[1],"low_power_mode":[0],"all_output_cache_and_acceptance_parity":true,"prefill_changed_kernel_dispatches":0,"controls_restore_one_reorder_per_reduction":true,"ttft":[{"prompt_tokens":512,"ttft_ms":{"baseline":165.3805,"candidate":165.01725,"noise":0.006820634839052886,"regression":-0.002196450004686268,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[162.241,163.402,162.808,162.652],[162.413,164.117,165.218,163.966],[166.084,164.676,165.217,165.068],[166.013,164.958,165.844,165.352],[165.48,165.839,166.398,165.393],[165.876,165.325,166.028,164.281],[166.404,166.077,164.865,164.565],[166.157,162.538,165.071,166.506]]},"prefill_gpu_ms":{"baseline":160.025,"candidate":160.019,"noise":0.004655697135965133,"regression":-3.74941415404173e-05,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[157.365,157.742,157.399,158.051],[156.919,158.807,159.622,158.821],[159.717,159.426,159.848,160.128],[160.0,159.903,160.488,160.284],[160.057,160.436,160.454,160.602],[160.32,159.867,160.401,158.983],[161.056,160.605,160.382,159.915],[161.294,157.427,160.427,160.904]]}},{"prompt_tokens":2016,"ttft_ms":{"baseline":528.3505,"candidate":527.9214999999999,"noise":0.0011592632616780458,"regression":-0.0008119609993746391,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[528.888,526.851,526.894,526.225],[525.148,526.673,526.96,527.022],[526.65,527.078,523.213,528.916],[528.665,527.989,528.03,527.824],[528.484,528.466,528.077,526.808],[528.796,529.131,528.294,528.429],[529.251,529.588,527.201,529.249],[525.855,529.842,530.514,532.67]]},"prefill_gpu_ms":{"baseline":522.9005,"candidate":522.53125,"noise":0.0008663216042057229,"regression":-0.0007061572899623725,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[522.951,521.689,521.315,521.55],[520.922,520.84,521.35,522.635],[521.246,521.666,518.117,522.649],[523.616,522.853,522.88,522.699],[523.504,522.272,523.092,522.178],[523.716,523.816,523.587,523.383],[523.303,523.791,522.12,524.786],[520.591,524.247,524.629,525.847]]}}],"decode":[{"width":1,"wall_ms_per_token":{"baseline":3.2424023437500002,"candidate":3.24092578125,"noise":0.0015596430900218956,"regression":-0.0004553915102011308,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[3.2471875,3.241203125,3.23846875,3.240703125],[3.241625,3.245703125,3.235390625,3.2346875],[3.243328125,3.280859375,3.248234375,3.234984375],[3.24134375,3.238859375,3.241609375,3.246578125],[3.240375,3.240703125,3.238390625,3.245484375],[3.237703125,3.264859375,3.238203125,3.238828125],[3.24440625,3.238828125,3.23671875,3.24415625],[3.24584375,3.245796875,3.23840625,3.243953125]]},"gpu_ms_per_token":{"baseline":3.18625390625,"candidate":3.18485546875,"noise":0.0012338668233326982,"regression":-0.00043889706882971335,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[3.1911875,3.185546875,3.18325,3.185125],[3.18434375,3.19096875,3.179125,3.179890625],[3.186265625,3.22534375,3.191890625,3.179921875],[3.18540625,3.18396875,3.186125,3.19140625],[3.18525,3.184375,3.182640625,3.190671875],[3.182890625,3.210859375,3.181375,3.183875],[3.1891875,3.183640625,3.18228125,3.18821875],[3.189796875,3.18809375,3.183140625,3.189453125]]}},{"width":2,"wall_ms_per_token":{"baseline":1.8079433593749998,"candidate":1.807974609375,"noise":0.001449765412399907,"regression":1.7284833530917254e-05,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.8058984375,1.8097734375,1.8058203125,1.809484375],[1.8064609375,1.8061171875,1.8067421875,1.8066171875],[1.8065859375,1.807,1.8071484375,1.8066484375],[1.8072421875,1.8088671875,1.8072109375,1.8099140625],[1.810890625,1.808671875,1.8080546875,1.8070546875],[1.8066796875,1.806046875,1.8054609375,1.8090234375],[1.806375,1.8154375,1.84378125,1.8719453125],[1.923671875,1.8950859375,1.961796875,1.938140625]]},"gpu_ms_per_token":{"baseline":1.7794160156249998,"candidate":1.779556640625,"noise":0.0009900544529949627,"regression":7.902873682463785e-05,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.778390625,1.780640625,1.7770546875,1.7806640625],[1.77765625,1.77715625,1.7786796875,1.7787578125],[1.7774765625,1.778890625,1.7787421875,1.778984375],[1.77915625,1.7808671875,1.7792578125,1.781890625],[1.7816015625,1.7795859375,1.779375,1.779515625],[1.7786796875,1.7782734375,1.777578125,1.7799296875],[1.7781953125,1.786859375,1.8143671875,1.8405390625],[1.8936796875,1.8639921875,1.929671875,1.9075703125]]}},{"width":3,"wall_ms_per_token":{"baseline":1.5909401041666666,"candidate":1.5861002604166667,"noise":0.0034528476141609354,"regression":-0.003042128196607985,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.5904166666666668,1.5930156249999998,1.5901770833333335,1.5872083333333336],[1.5890833333333332,1.5745677083333334,1.5635052083333332,1.5900052083333334],[1.4600416666666665,1.49371875,1.5017604166666667,1.5191354166666666],[1.5760781250000002,1.5921041666666669,1.5818177083333333,1.5942552083333332],[1.5770364583333334,1.5991614583333333,1.554953125,1.5955312499999998],[1.5996197916666668,1.5853906249999998,1.5979114583333331,1.6033385416666668],[1.5972447916666666,1.586125,1.5898072916666666,1.5965364583333335],[1.5968125000000002,1.5915520833333332,1.5849062500000002,1.5555312499999998]]},"gpu_ms_per_token":{"baseline":1.5679921875,"candidate":1.5647747395833334,"noise":0.0027043961832233523,"regression":-0.0020519540481872545,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.568125,1.5706562499999999,1.5675729166666665,1.5653645833333334],[1.5660052083333333,1.5519687500000001,1.5412552083333333,1.5660989583333331],[1.4379791666666666,1.4705625,1.4795468749999998,1.4960572916666666],[1.5546770833333332,1.5698229166666666,1.5614635416666667,1.5731718749999999],[1.5536197916666667,1.5775625,1.5333645833333334,1.5717395833333334],[1.5775677083333333,1.56365625,1.5757239583333333,1.5816979166666665],[1.575234375,1.56496875,1.5673802083333335,1.5738489583333333],[1.5744895833333334,1.5688125,1.5638541666666665,1.5345833333333332]]}},{"width":4,"wall_ms_per_token":{"baseline":1.267544921875,"candidate":1.266255859375,"noise":0.0009399321707964826,"regression":-0.0010169757913535138,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.26846875,1.25312890625,1.2629296875,1.26028125],[1.2582265625,1.257921875,1.26632421875,1.26784765625],[1.2678125,1.266125,1.26621875,1.26240625],[1.269421875,1.26699609375,1.26640625,1.26605078125],[1.2659609375,1.25597265625,1.251375,1.27371484375],[1.2462578125,1.26635546875,1.2745625,1.24492578125],[1.27038671875,1.27004296875,1.26723828125,1.27226171875],[1.27571875,1.2701015625,1.2525859375,1.271703125]]},"gpu_ms_per_token":{"baseline":1.250873046875,"candidate":1.2497490234375,"noise":0.0008119329156042293,"regression":-0.000898591140250482,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":7,"fail":1},"blocks_normalized_abba":[[1.25154296875,1.23644921875,1.2455625,1.2438046875],[1.2409453125,1.2408671875,1.25015234375,1.251625],[1.25121875,1.24987890625,1.249671875,1.2458046875],[1.2523671875,1.2499609375,1.24994140625,1.24910546875],[1.24926953125,1.23966015625,1.2350625,1.25708203125],[1.2289296875,1.24950390625,1.2577890625,1.22737109375],[1.2533203125,1.2525,1.25079296875,1.2556796875],[1.258734375,1.25309375,1.23678125,1.2542578125]]}}],"request_cache_parity_certificate_sha256":"04c4d86212211cd0750a2d0efa781bc7dbe712acc2cebf798235faf6e3c329a4","decode_acceptance_certificate_sha256":"d6607261068ac3851129cba1381542fd723ad8b1f40b23d0e45199601d137411"},{"model":"27b","loaded_layout_sha256":"73ca2049777aee00d1c77cd21449b25380035a389d826cdc0492ab7959972a9e","raw_json_sha256":"4e60eda34e76cf6fc727ab9ee568544df2df010787c65fe18494c817b7d6b79a","thermal_states":[1],"low_power_mode":[0],"all_output_cache_and_acceptance_parity":true,"prefill_changed_kernel_dispatches":0,"controls_restore_one_reorder_per_reduction":true,"ttft":[{"prompt_tokens":512,"ttft_ms":{"baseline":711.08975,"candidate":713.1419999999999,"noise":0.002511421287765887,"regression":0.0028860632571345857,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":7,"inconclusive":1},"blocks_normalized_abba":[[703.954,699.689,703.733,706.354],[711.244,712.264,713.483,711.622],[711.353,717.674,715.178,724.734],[754.588,760.651,684.3,688.801],[692.152,732.397,714.592,712.323],[722.711,705.092,707.721,709.067],[718.551,712.229,718.103,709.934],[711.45,710.615,714.651,713.515]]},"prefill_gpu_ms":{"baseline":705.5160000000001,"candidate":707.283,"noise":0.0019418411488896132,"regression":0.002504549861378047,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":7,"inconclusive":1},"blocks_normalized_abba":[[698.394,694.158,696.671,699.744],[705.986,707.226,707.397,706.125],[704.705,711.003,710.126,719.508],[747.58,754.471,679.12,682.872],[687.074,727.548,707.66,704.986],[717.412,700.3,702.852,704.272],[712.443,706.849,712.076,704.676],[706.416,704.384,709.264,706.471]]}},{"prompt_tokens":2016,"ttft_ms":{"baseline":2949.8100000000004,"candidate":2945.895,"noise":0.006586864916723394,"regression":-0.0013272041250116384,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":6,"inconclusive":2},"blocks_normalized_abba":[[2865.11,2843.19,2887.6,2933.58],[2686.68,2867.22,2957.53,2992.68],[2970.38,2997.78,2929.45,2906.9],[3065.94,2982.69,2976.89,2918.65],[2985.0,2905.16,3026.87,2985.07],[2998.94,3073.82,2890.89,2856.43],[2938.08,2946.0,2939.68,2946.61],[2948.67,2956.91,2940.99,2967.19]]},"prefill_gpu_ms":{"baseline":2942.6875,"candidate":2938.605,"noise":0.006274876282310013,"regression":-0.0013873372554850105,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":6,"inconclusive":2},"blocks_normalized_abba":[[2858.4,2836.84,2880.11,2928.17],[2680.1,2860.33,2951.55,2985.97],[2962.89,2987.52,2923.8,2901.0],[3060.72,2976.86,2971.26,2911.57],[2975.23,2897.83,3020.63,2978.62],[2988.56,3067.22,2884.03,2848.28],[2931.31,2939.25,2931.67,2938.74],[2940.95,2950.58,2932.92,2959.25]]}}],"decode":[{"width":1,"wall_ms_per_token":{"baseline":10.41555859375,"candidate":10.3991875,"noise":0.0013253909884664377,"regression":-0.0015717921993951833,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[10.43534375,10.393546875,10.405078125,10.3991875],[10.4036875,10.390265625,10.39378125,10.41259375],[10.420828125,10.397296875,10.412359375,10.405921875],[10.409265625,10.40059375,10.363984375,10.411390625],[10.435265625,10.40321875,10.414515625,10.401796875],[10.4274375,10.436203125,10.350640625,10.3298125],[10.42409375,10.4015625,10.34065625,10.436171875],[10.411609375,10.392984375,10.4123125,10.47771875]]},"gpu_ms_per_token":{"baseline":10.36190234375,"candidate":10.34504296875,"noise":0.0012613815558570548,"regression":-0.001627054033197739,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[10.380484375,10.3391875,10.3521875,10.345015625],[10.349890625,10.3395625,10.341078125,10.357796875],[10.3668125,10.342640625,10.35884375,10.3529375],[10.35228125,10.3434375,10.310484375,10.359234375],[10.380703125,10.34921875,10.36046875,10.347734375],[10.37621875,10.384265625,10.2979375,10.273140625],[10.3700625,10.348265625,10.28728125,10.383828125],[10.35903125,10.34346875,10.357,10.425671875]]}},{"width":2,"wall_ms_per_token":{"baseline":5.481791015624999,"candidate":5.4719785156250005,"noise":0.0002572437084852327,"regression":-0.001790017162644486,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[5.46571875,5.4736484375,5.4698671875,5.4762890625],[5.467859375,5.4400703125,5.4750703125,5.478078125],[5.48125,5.4708125,5.4613125,5.486390625],[5.487125,5.4721640625,5.4650546875,5.4853828125],[5.4611953125,5.457390625,5.4751953125,5.4282578125],[5.4837421875,5.482765625,5.487375,5.435328125],[5.499328125,5.4620859375,5.478875,5.48409375],[5.4933125,5.47875,5.462109375,5.48778125]]},"gpu_ms_per_token":{"baseline":5.4547890625,"candidate":5.44401171875,"noise":0.00015038390863535654,"regression":-0.001975758113928028,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[5.4372421875,5.4452578125,5.4418984375,5.4485859375],[5.4395546875,5.410984375,5.4466328125,5.4510625],[5.4540078125,5.443046875,5.43409375,5.4589375],[5.459171875,5.44446875,5.4373125,5.4576953125],[5.43321875,5.431078125,5.446875,5.398703125],[5.456390625,5.4572421875,5.4595,5.4067109375],[5.4711796875,5.4330625,5.450734375,5.4578671875],[5.465828125,5.44934375,5.4343671875,5.461296875]]}},{"width":3,"wall_ms_per_token":{"baseline":4.713838541666666,"candidate":4.70965234375,"noise":0.0002961139464144569,"regression":-0.0008880656135468179,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":7,"inconclusive":1},"blocks_normalized_abba":[[4.713614583333333,4.707552083333334,4.706505208333334,4.708390625],[4.712666666666666,4.721796875,4.705369791666667,4.713666666666667],[4.699307291666667,4.712958333333334,4.716786458333334,4.712020833333333],[4.727979166666667,4.687265625,4.709755208333333,4.708171875000001],[4.695744791666667,4.715281249999999,4.70834375,4.71540625],[4.700286458333333,4.7002395833333335,4.7210624999999995,4.722218750000001],[4.729177083333333,4.718453125,5.071546875,5.203671875],[4.993994791666666,4.6886093749999995,4.6974375,4.814239583333333]]},"gpu_ms_per_token":{"baseline":4.695118489583333,"candidate":4.691145833333334,"noise":0.00042763830230358634,"regression":-0.0008461248121454101,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":7,"inconclusive":1},"blocks_normalized_abba":[[4.694427083333333,4.689953125,4.687994791666667,4.688583333333333],[4.693802083333334,4.7021927083333335,4.687588541666667,4.6947447916666665],[4.679942708333333,4.693796874999999,4.699348958333333,4.691859375],[4.708859375,4.666744791666667,4.690932291666667,4.689557291666667],[4.678661458333333,4.695817708333333,4.689901041666666,4.6975],[4.679729166666667,4.6820208333333335,4.702,4.702927083333333],[4.710994791666667,4.701328125,5.051276041666667,5.182729166666666],[4.973151041666667,4.670630208333333,4.676125,4.794239583333334]]}},{"width":4,"wall_ms_per_token":{"baseline":3.5519619140625,"candidate":3.5613906249999996,"noise":0.0044586252736597155,"regression":0.002654507893277369,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":6,"inconclusive":2},"blocks_normalized_abba":[[3.5604296875,3.55337109375,3.5427578125,3.5309921875],[3.54832421875,3.55319921875,3.54591796875,3.54375],[3.54189453125,3.55114453125,3.55202734375,3.5704921875],[3.55931640625,3.56041796875,3.554875,3.55216796875],[3.548234375,3.5782421875,3.57368359375,3.5931015625],[3.80625,3.764484375,3.71364453125,3.47859375],[3.55261328125,3.6835703125,3.86438671875,3.5102890625],[3.57276953125,3.92171875,3.41071484375,3.589453125]]},"gpu_ms_per_token":{"baseline":3.5368671875000004,"candidate":3.546384765625,"noise":0.004501723531170871,"regression":0.0026909628268305674,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":6,"inconclusive":2},"blocks_normalized_abba":[[3.54613671875,3.53844140625,3.527734375,3.51619921875],[3.53366015625,3.53758203125,3.53090234375,3.52924609375],[3.5268515625,3.53600390625,3.53748828125,3.55517578125],[3.54489453125,3.545828125,3.53931640625,3.53622265625],[3.53330078125,3.56290625,3.558203125,3.577875],[3.7901015625,3.7501640625,3.69846484375,3.46425],[3.53710546875,3.66886328125,3.84891015625,3.49442578125],[3.5584296875,3.905984375,3.3949375,3.5744296875]]}}],"request_cache_parity_certificate_sha256":"a213ea022e180fceb24980d1d4b13abbfc49faccb6dddaffa33cf34ad1ec51e9","decode_acceptance_certificate_sha256":"0a148e0fdc7c17f65b8d48a04483ece143debf9dbbce665baac39332594ccc0a"}],"focused_repeat":{"raw_json_sha256":"731c6de6ea840d5a1148129b3a1d0b44e9d6d5428f7dbf48360187664ce576c2","results":[{"width":1,"wall_ms_per_token":{"baseline":3.19404296875,"candidate":3.1929921875,"noise":0.0007509095912190452,"regression":-0.0003289815635796334,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[3.193203125,3.20753125,3.192828125,3.20096875],[3.189328125,3.19003125,3.186890625,3.1885],[3.190265625,3.195078125,3.1955,3.200625],[3.193296875,3.19175,3.188828125,3.18921875],[3.200140625,3.1925,3.201328125,3.20246875],[3.190359375,3.200984375,3.19121875,3.189609375],[3.198421875,3.190421875,3.199703125,3.197],[3.192484375,3.18840625,3.194890625,3.193484375]]},"gpu_ms_per_token":{"baseline":3.13758984375,"candidate":3.13877734375,"noise":0.0009984773842382762,"regression":0.00037847521796563477,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[3.138171875,3.147515625,3.136828125,3.1423125],[3.13484375,3.136078125,3.1313125,3.132609375],[3.135625,3.140328125,3.140828125,3.145890625],[3.135984375,3.137703125,3.13234375,3.1335],[3.14325,3.138875,3.147,3.147125],[3.1360625,3.148296875,3.13771875,3.1353125],[3.143234375,3.13590625,3.14453125,3.140578125],[3.135828125,3.13346875,3.1408125,3.137734375]]}},{"width":2,"wall_ms_per_token":{"baseline":1.782392578125,"candidate":1.7823535156249999,"noise":0.00108485422956174,"regression":-2.1915766750590926e-05,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.7822578125,1.78371875,1.7821640625,1.779203125],[1.7839921875,1.7779140625,1.77846875,1.78446875],[1.7798828125,1.7790546875,1.7815234375,1.782359375],[1.7879609375,1.782921875,1.778625,1.7823515625],[1.780203125,1.7813125,1.78125,1.784421875],[1.781703125,1.7863671875,1.7791171875,1.7801015625],[1.7830703125,1.7865546875,1.7829609375,1.78790625],[1.7816875,1.7840390625,1.784703125,1.78325]]},"gpu_ms_per_token":{"baseline":1.754390625,"candidate":1.75450390625,"noise":0.0013313948698994694,"regression":6.45701409855981e-05,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.7548046875,1.756203125,1.7531171875,1.7514140625],[1.757,1.7501484375,1.751421875,1.757078125],[1.751515625,1.750828125,1.7540859375,1.7543515625],[1.7599453125,1.755140625,1.75121875,1.7548046875],[1.752015625,1.7531640625,1.7535546875,1.7564140625],[1.7536015625,1.7584921875,1.7518828125,1.752671875],[1.7548125,1.7593984375,1.7545234375,1.760671875],[1.75290625,1.75653125,1.7561484375,1.7542890625]]}},{"width":3,"wall_ms_per_token":{"baseline":1.3894075520833333,"candidate":1.3881705729166667,"noise":0.0003654885128834573,"regression":-0.0008902925313828414,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.3897239583333334,1.3867083333333332,1.3858854166666665,1.3897395833333333],[1.3896145833333333,1.3888437500000002,1.38628125,1.388375],[1.3885260416666665,1.3911666666666667,1.3911875,1.3909583333333335],[1.3868385416666669,1.38803125,1.3895104166666667,1.3873541666666667],[1.3889947916666667,1.3903854166666667,1.386125,1.3895833333333334],[1.3926718749999998,1.3873072916666667,1.3899739583333333,1.3989062499999998],[1.389171875,1.3882135416666666,1.3869270833333334,1.392125],[1.3891354166666667,1.3877395833333335,1.3895937500000002,1.3876302083333334]]},"gpu_ms_per_token":{"baseline":1.3694388020833332,"candidate":1.3687213541666667,"noise":0.00044902005178949707,"regression":-0.0005238992173838275,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.3706458333333333,1.3670156249999998,1.3667604166666667,1.3699166666666667],[1.3702916666666667,1.3697447916666665,1.3663020833333333,1.3681875000000001],[1.3693697916666665,1.3717083333333333,1.3721041666666667,1.3715833333333334],[1.3675989583333334,1.3680312499999998,1.3700364583333335,1.3679166666666667],[1.369515625,1.3700885416666668,1.3664739583333334,1.368953125],[1.37353125,1.3682135416666668,1.370390625,1.378234375],[1.36865625,1.3683489583333335,1.3680520833333334,1.3719999999999999],[1.3689375000000001,1.3684791666666667,1.3700052083333334,1.3679739583333335]]}},{"width":4,"wall_ms_per_token":{"baseline":1.1117900390625,"candidate":1.1112988281249998,"noise":0.0006784010123291141,"regression":-0.00044181987627300856,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.11194921875,1.1090390625,1.1104453125,1.1105390625],[1.11229296875,1.113203125,1.10818359375,1.11106640625],[1.110265625,1.1125078125,1.113390625,1.11125390625],[1.1111328125,1.10818359375,1.1130234375,1.1119140625],[1.11204296875,1.111234375,1.10933984375,1.11116015625],[1.11580078125,1.1134609375,1.108953125,1.11267578125],[1.11006640625,1.1121171875,1.1113984375,1.11218359375],[1.132640625,1.11047265625,1.11365625,1.113234375]]},"gpu_ms_per_token":{"baseline":1.09712890625,"candidate":1.0964716796874998,"noise":0.0007748575642576182,"regression":-0.0005990422444949184,"bound":0.02,"verdict":"pass","paired_verdict_counts":{"pass":8},"blocks_normalized_abba":[[1.09783984375,1.09412109375,1.0951796875,1.09561328125],[1.0975546875,1.098484375,1.093796875,1.09646484375],[1.0955078125,1.09804296875,1.09873828125,1.09640625],[1.09647265625,1.09386328125,1.09876953125,1.09696484375],[1.0974453125,1.09688671875,1.0943046875,1.09655078125],[1.10095703125,1.0990703125,1.09416796875,1.09823046875],[1.0950234375,1.09690625,1.0969140625,1.098015625],[1.113859375,1.09596875,1.09881640625,1.09839453125]]}}]}}Portable diagnostic adapter and exact workload edits
This is experimental diagnostic code, excluded from the production source commits. Save the following as make_attention_diagnostic.py and run it from the pinned candidate checkout after building build/splash and build/engine-tests/backend-benchmark:
python3 make_attention_diagnostic.py ../splash-baseline/build/splash.metallib
build/attention-diagnostic/attention-native-probe build/splash.metallib MODEL_ROOT \
--samples 40 --scenario decode,short --ane-ffn-share 0Use MTL_SHADER_VALIDATION=1 with eight samples for the separate correctness run. The control uses the baseline library; output/cache/acceptance must match and each cold-request trace must report zero reductions/reorders. The stderr mode/counter traces and raw JSON distinguish the two arms. The primary benchmark requires a pre-fusion baseline ABI.
from pathlib import Path
import json,subprocess,time,sys
tree=Path.cwd().resolve()
out=tree/'build/attention-diagnostic';out.mkdir(parents=True,exist_ok=True)
baseline=Path(sys.argv[1]).resolve();assert baseline.is_file()
header=r'''#pragma once
#include <atomic>
#include <cstdint>
namespace splash::metal::attention_probe {
inline std::atomic<bool> legacy{false};
inline std::atomic<uint64_t> reductions{0}, reorders{0};
inline void select(uint32_t sample) {
const bool endpoints=sample%4==0||sample%4==3;
legacy.store((sample/4)%2==0?endpoints:!endpoints,std::memory_order_release);
}
}
'''
(out/'AttentionExperiment.hpp').write_text(header)
backend=(tree/'runtime/metal/MetalBackend.mm').read_text()
backend='#include "AttentionExperiment.hpp"\n#include "metal/abi/DraftAttention.h"\n'+backend
transform=r'''
namespace splash::metal::attention_probe {
struct Transform {
std::vector<ComputeDispatch> dispatches;
Command value;
explicit Transform(const Command &original):value(original) {
const bool old=legacy.load(std::memory_order_acquire);
uint64_t reduced=0,extra=0;
dispatches.reserve(original.dispatches.size()+8);
for(const auto &source:original.dispatches) {
auto d=source;
if(d.pipelineName=="draft_attention_bf16_reduce") {
++reduced;
if(old) {
if(!original.events.empty())throw MetalBackendError("attention probe does not rewrite host-event positions");
DraftAttentionBatchParams p{};
if(d.bytes.size()!=1||d.bytes[0].sizeBytes!=sizeof(p)||d.buffers.size()!=2)
throw MetalBackendError("attention probe reduction ABI mismatch");
std::memcpy(&p,d.bytes[0].data,sizeof(p));
const auto queries=d.buffers[0].buffer,output=d.buffers[1].buffer;
d.pipelineName="boundary_old_draft_attention_bf16_reduce";
d.buffers.resize(1);d.bytes[0].index=1;
dispatches.push_back(std::move(d));
ComputeDispatch reorder;reorder.pipelineName="boundary_old_draft_attention_reorder";
reorder.buffers={{0,queries},{1,output}};
reorder.threadgroups={(SPLASH_DRAFT_QUERY_ROWS*SPLASH_DRAFT_ATTENTION_WIDTH+255)/256,p.lanes,1};
reorder.threadsPerThreadgroup={256,1,1};dispatches.push_back(std::move(reorder));
++extra;
continue;
}
}
dispatches.push_back(std::move(d));
}
if(reduced)reductions.fetch_add(reduced,std::memory_order_relaxed);
if(extra)reorders.fetch_add(extra,std::memory_order_relaxed);
value.dispatches=dispatches;
}
};
}
'''
backend=backend.replace('namespace splash::metal {',transform+'\nnamespace splash::metal {',1)
needle=' NSString *key = checkedNSString(name, "pipeline name");'
assert backend.count(needle)==1
backend=backend.replace(needle,''' if(name.starts_with("boundary_old_")) {
static id<MTLLibrary> oldLibrary = [device newLibraryWithURL:[NSURL fileURLWithPath:@'''+json.dumps(str(baseline))+'''] error:nil];
source=oldLibrary;name.remove_prefix(13);
}
'''+needle)
assert backend.count('CommandTicket MetalBackend::submitCommandAsync(const Command &command,')==1
backend=backend.replace('CommandTicket MetalBackend::submitCommandAsync(const Command &command,','CommandTicket MetalBackend::submitCommandAsync(const Command &original,')
needle=' [[maybe_unused]] const NSUInteger buffers = checkEvents(command.events, command.dispatches.size());'
assert backend.count(needle)==1
backend=backend.replace(needle,' attention_probe::Transform transformed(original);\n const auto &command=transformed.value;\n'+needle)
(out/'MetalBackendAttention.mm').write_text(backend)
native=(tree/'dev/benchmarks/backend_benchmark.mm').read_text()
native='#include "AttentionExperiment.hpp"\n#import <Foundation/Foundation.h>\n'+native
def patch_function(source,signature,kind):
start=source.index(signature);end=source.index('\n}\n',start)+3
part=source[start:end]
brace=part.index(' {')+2
assert brace>1 and part.count(' return result;')==1
part=part[:brace]+'''
metal::attention_probe::select(sample);
const auto probeReductions=metal::attention_probe::reductions.load(std::memory_order_relaxed);
const auto probeReorders=metal::attention_probe::reorders.load(std::memory_order_relaxed);
'''+part[brace:]
trace='''
std::fprintf(stderr,"ATTN_'''+kind+''' sample=%u width=%u prompt=%u legacy=%d reductions=%llu reorders=%llu thermal=%ld low_power=%d\\n",
sample,'''+('width' if kind=='DECODE' else '0u')+''',static_cast<uint32_t>(tokens.size()),
metal::attention_probe::legacy.load(std::memory_order_relaxed),
static_cast<unsigned long long>(metal::attention_probe::reductions.load(std::memory_order_relaxed)-probeReductions),
static_cast<unsigned long long>(metal::attention_probe::reorders.load(std::memory_order_relaxed)-probeReorders),
static_cast<long>([NSProcessInfo processInfo].thermalState),[NSProcessInfo processInfo].lowPowerModeEnabled);
'''
if kind=='REQUEST':
# runRequest moves tokens into EngineRequest before the trace.
trace=trace.replace('static_cast<uint32_t>(tokens.size())','result.promptTokens')
trace='\n if(result.scenario=="short") {'+trace+' }\n'
part=part.replace(' return result;',trace+' return result;')
return source[:start]+part+source[end:]
native=patch_function(native,'Measurement runRequest(engine::Engine &engine','REQUEST')
native=patch_function(native,'runDecodeThroughput(engine::Engine &engine','DECODE')
old=''' for (uint32_t sample = 0; sample < samples; ++sample) {
for (uint32_t offset = 0;
offset < model::ExecutionLimits::maximumBatchWidth; ++offset) {
const uint32_t width =
1 + (sample + offset) % model::ExecutionLimits::maximumBatchWidth;'''
assert native.count(old)==1
native=native.replace(old,''' for (uint32_t width = 1; width <= model::ExecutionLimits::maximumBatchWidth; ++width) {
for (uint32_t sample = 0; sample < samples; ++sample) {''')
old=''' constexpr std::array<uint32_t, 8> shortRows{480, 512, 544, 640, 672, 1024, 1536, 2016};
for (uint32_t sample = 0; sample < samples; ++sample) {
for (uint32_t rows : shortRows) {'''
assert native.count(old)==1
native=native.replace(old,''' constexpr std::array<uint32_t, 2> shortRows{512, 2016};
for (uint32_t rows : shortRows) {
for (uint32_t sample = 0; sample < samples; ++sample) {''')
native=native.replace('prompt(rows, (uint64_t{rows} << 32 | sample) ^ 0x53484f5254ULL)','prompt(rows, (uint64_t{rows} << 32) ^ 0x53484f5254ULL)')
(out/'attention_native_probe.mm').write_text(native)
flags=['-std=c++20','-O3','-Wall','-Wextra','-Werror','-fobjc-arc','-mmacosx-version-min=26.4','-I'+str(out),'-Iruntime','-Iruntime/metal','-Idev/benchmarks']
with (out/'probe-build.log').open('w') as log:
subprocess.run(['xcrun','-sdk','macosx','clang++',*flags,'-c',str(out/'MetalBackendAttention.mm'),'-o',str(out/'attention-backend.o')],cwd=tree,check=True,stdout=log,stderr=subprocess.STDOUT)
subprocess.run(['xcrun','-sdk','macosx','clang++',*flags,'-include','build/engine/BuildIdentity.hpp',str(out/'attention_native_probe.mm'),str(out/'attention-backend.o'),'build/engine/libsplash.a','-framework','Foundation','-framework','Metal','-framework','IOKit','-framework','IOSurface','-o',str(out/'attention-native-probe')],cwd=tree,check=True,stdout=log,stderr=subprocess.STDOUT)
(out/'probe-build.status').write_text('0')
print('Attention-only interleaved native probe built',flush=True)Original standalone production-binary measurements at 744e642, preserved
These original multi-process timings and combined-build history predate the follow-up. Their failed/inconclusive records are retained; the controlled diagnostic above is separate evidence.
Measured component result
| Lanes | Without: reduce + reorder (µs) | With: fused reduce (µs) | Paired stage speedup |
|---|---|---|---|
| B1 | 4.225 | 2.742 | 1.575× |
| B2 | 4.017 | 2.537 | 1.583× |
| B3 | 4.913 | 3.290 | 1.493× |
| B4 | 5.073 | 3.343 | 1.516× |
The replaced stage was measured with the preserved old/new shader libraries, 20 ABBA rounds after four discarded rounds, 256 repeated stages/command. All BF16 output bytes match at B1–B4. Timing columns are marginal medians and the speedup column is the paired median; their quotient need not equal the paired statistic. This is one reduce/layout stage, not a complete draft, target, or serving speedup; it is not directly comparable to #702’s five-layer graph benchmark.
TTFT: standalone change versus no change
| Model | Cold prompt tokens | Without (ms) | With (ms) | Change | ABBA spread | Verdict |
|---|---|---|---|---|---|---|
| Qwen3.6-35B-A3B-4bit | 480 | 141.555 | 140.780 | -0.55% | 2.07% | pass |
| Qwen3.6-35B-A3B-4bit | 512 | 146.523 | 144.976 | -1.06% | 2.68% | pass |
| Qwen3.6-35B-A3B-4bit | 544 | 153.611 | 153.136 | -0.31% | 3.05% | pass |
| Qwen3.6-35B-A3B-4bit | 640 | 173.755 | 170.697 | -1.76% | 3.51% | pass |
| Qwen3.6-35B-A3B-4bit | 672 | 180.443 | 178.244 | -1.22% | 4.92% | pass |
| Qwen3.6-35B-A3B-4bit | 1024 | 255.663 | 249.505 | -2.41% | 6.10% | inconclusive |
| Qwen3.6-35B-A3B-4bit | 1536 | 364.599 | 356.317 | -2.27% | 7.11% | inconclusive |
| Qwen3.6-35B-A3B-4bit | 2016 | 469.512 | 457.382 | -2.58% | 7.68% | inconclusive |
| Qwen3.8-27B-4bit | 480 | 691.687 | 707.115 | +2.23% | 12.37% | inconclusive |
| Qwen3.8-27B-4bit | 512 | 713.147 | 731.004 | +2.50% | 14.47% | inconclusive |
| Qwen3.8-27B-4bit | 544 | 756.564 | 787.989 | +4.15% | 13.57% | inconclusive |
| Qwen3.8-27B-4bit | 640 | 879.004 | 910.374 | +3.57% | 13.16% | inconclusive |
| Qwen3.8-27B-4bit | 672 | 928.962 | 969.433 | +4.36% | 12.63% | inconclusive |
| Qwen3.8-27B-4bit | 1024 | 1396.360 | 1442.590 | +3.31% | 12.98% | inconclusive |
| Qwen3.8-27B-4bit | 1536 | 2087.915 | 2162.020 | +3.55% | 12.84% | inconclusive |
| Qwen3.8-27B-4bit | 2016 | 2761.935 | 2870.020 | +3.91% | 12.81% | inconclusive |
Decode tokens/sec: standalone change versus no change
| Model | Concurrent lanes | Without (tok/s) | With (tok/s) | Change | ABBA spread | Verdict |
|---|---|---|---|---|---|---|
| Qwen3.6-35B-A3B-4bit | B1 | 312.47 | 312.60 | +0.04% | 0.21% | pass |
| Qwen3.6-35B-A3B-4bit | B2 | 561.62 | 561.86 | +0.04% | 0.33% | pass |
| Qwen3.6-35B-A3B-4bit | B3 | 719.21 | 719.78 | +0.08% | 0.46% | pass |
| Qwen3.6-35B-A3B-4bit | B4 | 899.23 | 902.01 | +0.31% | 0.22% | pass |
| Qwen3.8-27B-4bit | B1 | 99.22 | 96.97 | -2.27% | 11.93% | inconclusive |
| Qwen3.8-27B-4bit | B2 | 188.21 | 184.92 | -1.75% | 10.02% | inconclusive |
| Qwen3.8-27B-4bit | B3 | 215.26 | 213.56 | -0.79% | 17.00% | inconclusive |
| Qwen3.8-27B-4bit | B4 | 283.97 | 278.89 | -1.79% | 16.51% | inconclusive |
Method and limits
Apple M5 Pro, 20 GPU cores, 48 GB; macOS 26.6 (25G72), Apple clang 21.0.0, Metal 32023.922. Target macOS 26.4 / Metal 4.0.
Baseline: d8af1054. Candidate: 744e6426. Both split branches start at the baseline and contain only their own change. Generated reports and prototypes are excluded from source commits, per CONTRIBUTING.
These are standalone native Engine measurements for this change alone, not the earlier two-change build and not HTTP throughput. Both installed MLX affine-4bit models use unchanged weights/layouts, GPU-only prefill (--ane-ffn-share 0) and the benchmark's fixed verification of every draft row.
Four fresh processes in baseline / candidate / candidate / baseline (ABBA) order on one machine, serialized GPU work. Each process records four samples; sample 0 is retained for correctness and excluded from timing, leaving the median of samples 1–3. Every output, cache work counter, loaded layout, decode output hash and drafted/accepted count is checked across builds.
TTFT is wall time to the first emitted token of a cold, single-output-token request. Decode rate is aggregate wall-clock decode output tokens/sec across the stated number of lanes, with 64 decode output tokens/lane. Each arm averages its two round-median ms/output-token values; TPS is 1000 / that average. It is not an arithmetic average of TPS samples.
The repository's dev/benchmarks/abba.py rule treats >5% ABBA spread as inconclusive; otherwise the allowed regression is max(2%, 2 × spread). pass is a non-regression verdict, not proof of a speedup. Both negative and inconclusive results are included; no general model/HTTP throughput improvement is claimed.
Reproduction: build the pinned baseline and this branch using their own production backend-benchmark and splash.metallib, then run each round with the same installed model root:
build/engine-tests/backend-benchmark build/splash.metallib MODEL_ROOT \
--samples 4 --scenario decode,short --ane-ffn-share 0Use ABBA order, discard sample 0 only for timing, and compare with the same abba.py rule. All native benchmark settings and output hashes are retained in the raw round summary below.
Broader fusion experiments rejected: expert gate/up+down fusion raised the tested 35B B1 GPU cycle from about 19 to 43 ms; combine/norm and convolution+RMS/Table16 fusions preserved tested values but gave no useful model-level gain. Only the small reduction/layout fusion is proposed here.
Earlier broad combined-change runs were mostly neutral/noisy. A 28-token 27B TTFT comparison flagged +8%, but it did not reproduce in a validated, one-loaded-model interleaved restoration of old/new behavior (TTFT −0.32%, GPU prefill −0.38%, spread below 0.8%, exact tokens). Those results do not prove either standalone PR faster and do not erase the original failed/inconclusive measurements. The tables above are newly measured one-change-at-a-time comparisons.
Validation
On this standalone commit: production/benchmark build and check-source pass; draft-attention, dflash-batch-control, draft-selector, and the draft-preparation golden checks pass with MTL_SHADER_VALIDATION=1. The native A/B below checks both models at B1–B4, all 32 cold-request samples and all 16 decode samples/round.
The earlier combined candidate additionally passed native CPU/Metal and Metal TSAN, CPU ASAN/UBSAN/TSAN, Python (1,641 tests, 32 expected skips), Swift (218 tests), and both models’ three runtime oracles, vision parity and HTTP smoke (including BF16 state and 32K spill). Those are combined-build results, not a claim that the entire suite was rerun on this split head. Python used an isolated standalone 3.12 environment because the existing pyenv 3.14 interpreter aborts under the unchanged PDF sandbox; Swift used per-process en_US for two locale-sensitive number-format expectations. No frontend, sandbox, global locale or existing .venv change is included.
Only this M5 Pro and the two installed MLX models were measured. This is not a hardware-matrix release qualification. CUDA TMA/NVLink/cluster/code-address cache operations were unavailable; local analogues were separate experiments. No full attention+MoE megakernel, native shared-prefix tree verifier or GPU VM is claimed.
Recalculable raw ABBA medians, parity certificates and build provenance
Round arrays are in baseline/candidate/candidate/baseline order. Request/cache and decode/acceptance hashes must match across all four phases; raw JSON SHA-256 values identify the retained complete local files. No model weights or private paths are included here.
{
"provenance": {
"baseline_commit": "d8af1054aab0f97de21d043e01aa98af3754f10e",
"candidate_commit": "744e64260178ad94192e2a2fda54a1a8dc1a139e",
"candidate_build_identity": {
"build_id": "src-16fcd427f275474f8214e4a6163c7c11b219b7a85e0dfaa2b577cbca1ed803e3",
"schema_version": 1,
"source_sha256": "16fcd427f275474f8214e4a6163c7c11b219b7a85e0dfaa2b577cbca1ed803e3"
},
"binary_sha256": "9bbac83f63b8bfcef8d8238af1e5e4b953c2af38f73a3705096acb56f73f3e42",
"metallib_sha256": "a6403a17ac939193909f2e227abcce3f0e2eefdf9ac93040e0496d23fd9971f8",
"baseline_binary_sha256": "b61a0ccd08c59d98c2dcc42764ee0729ff2768791e72541727aebb45cf4f5bd2",
"baseline_metallib_sha256": "7006f53dedd016809a16f4e5694c78e4f8959f31691e4473a35ad46be45b9b03"
},
"measurements": [
{
"model": "Qwen3.6-35B-A3B-4bit",
"source_commit": "744e64260178ad94192e2a2fda54a1a8dc1a139e",
"loaded_layout_sha256": "964ce8afdb896d17505b4f1b322c0d81974ad8a38a16437c15dfd31fbc9bc377",
"output_parity": true,
"cache_parity": true,
"ttft_round_medians_ms": {
"480": [
140.089,
139.858,
141.702,
143.02
],
"512": [
144.562,
145.08,
144.871,
148.484
],
"544": [
151.265,
152.35,
153.922,
155.956
],
"640": [
170.703,
170.311,
171.084,
176.806
],
"672": [
176.006,
178.172,
178.316,
184.88
],
"1024": [
247.861,
248.213,
250.797,
263.466
],
"1536": [
351.632,
353.848,
358.786,
377.565
],
"2016": [
451.483,
451.544,
463.22,
487.542
]
},
"decode_round_medians_ms_per_token": {
"1": [
3.1969375,
3.20009375,
3.197859375,
3.203625
],
"2": [
1.77759375,
1.7785234375,
1.7810625,
1.783515625
],
"3": [
1.3871874999999998,
1.3886614583333332,
1.3899739583333333,
1.393640625
],
"4": [
1.110859375,
1.10867578125,
1.10859375,
1.11326171875
]
},
"phases": [
{
"phase": 0,
"arm": "baseline",
"raw_json_sha256": "d7dea43d91712d6f30fab9816303a60be6f89594991022530927e3d22fb0289e",
"request_outputs_cache_sha256": "1b91f64ee4280cc16b6d272e341cb85e6d6e7bf9611684c5d58f1cf5ba37cffd",
"decode_outputs_acceptance_sha256": "b4a27441886c8e9322925675567c5285a6815fee92709c9cd8e99ab9be5c185e"
},
{
"phase": 1,
"arm": "candidate",
"raw_json_sha256": "496dc71079cddd2782dee435b9f75937e25ea46f9d127c0e96b2e4e8a377a2b9",
"request_outputs_cache_sha256": "1b91f64ee4280cc16b6d272e341cb85e6d6e7bf9611684c5d58f1cf5ba37cffd",
"decode_outputs_acceptance_sha256": "b4a27441886c8e9322925675567c5285a6815fee92709c9cd8e99ab9be5c185e"
},
{
"phase": 2,
"arm": "candidate",
"raw_json_sha256": "9d27c0e74efb6c49999c3f32891e04d89d8c8f02c5a9a6020a4cd72ff4a952c2",
"request_outputs_cache_sha256": "1b91f64ee4280cc16b6d272e341cb85e6d6e7bf9611684c5d58f1cf5ba37cffd",
"decode_outputs_acceptance_sha256": "b4a27441886c8e9322925675567c5285a6815fee92709c9cd8e99ab9be5c185e"
},
{
"phase": 3,
"arm": "baseline",
"raw_json_sha256": "d0b9d8a2e43c8e9439f065ac63482fe129e05c3e0b4d4041224ea217b56a4695",
"request_outputs_cache_sha256": "1b91f64ee4280cc16b6d272e341cb85e6d6e7bf9611684c5d58f1cf5ba37cffd",
"decode_outputs_acceptance_sha256": "b4a27441886c8e9322925675567c5285a6815fee92709c9cd8e99ab9be5c185e"
}
]
},
{
"model": "Qwen3.8-27B-4bit",
"source_commit": "744e64260178ad94192e2a2fda54a1a8dc1a139e",
"loaded_layout_sha256": "73ca2049777aee00d1c77cd21449b25380035a389d826cdc0492ab7959972a9e",
"output_parity": true,
"cache_parity": true,
"ttft_round_medians_ms": {
"480": [
648.899,
681.525,
732.705,
734.474
],
"512": [
661.542,
698.838,
763.17,
764.751
],
"544": [
705.246,
746.499,
829.479,
807.882
],
"640": [
821.16,
867.987,
952.761,
936.847
],
"672": [
870.319,
922.087,
1016.78,
987.605
],
"1024": [
1305.71,
1404.12,
1481.06,
1487.01
],
"1536": [
1953.82,
2052.55,
2271.49,
2222.01
],
"2016": [
2584.98,
2743.28,
2996.76,
2938.89
]
},
"decode_round_medians_ms_per_token": {
"1": [
9.47715625,
10.107703125,
10.517078125,
10.679953125
],
"2": [
5.0469453125,
5.3130390625,
5.50253125,
5.5793671875
],
"3": [
4.250630208333333,
4.483229166666667,
4.882036458333333,
5.0403125
],
"4": [
3.23077734375,
3.44691796875,
3.72432421875,
3.81219140625
]
},
"phases": [
{
"phase": 0,
"arm": "baseline",
"raw_json_sha256": "4d964a06a96f52a35538ef9695476d07e2ec883e6360e08260051fee9958e375",
"request_outputs_cache_sha256": "61aa5c339936495f81b2b2657c72f6203b03f02f797e967dab7b5f563281db7d",
"decode_outputs_acceptance_sha256": "11436014ce1fecd54d5a349765e9ba435c372bee953984a828b48b83d0ff749f"
},
{
"phase": 1,
"arm": "candidate",
"raw_json_sha256": "24f13c77fbe10758cb397dd4fffccc36a18d1200ed2aeecdc2e4d26f8480cf00",
"request_outputs_cache_sha256": "61aa5c339936495f81b2b2657c72f6203b03f02f797e967dab7b5f563281db7d",
"decode_outputs_acceptance_sha256": "11436014ce1fecd54d5a349765e9ba435c372bee953984a828b48b83d0ff749f"
},
{
"phase": 2,
"arm": "candidate",
"raw_json_sha256": "2bc41e2f73f97d020ec2524bf1e1587800fd2edc3aba86cd82130ba3b6f43270",
"request_outputs_cache_sha256": "61aa5c339936495f81b2b2657c72f6203b03f02f797e967dab7b5f563281db7d",
"decode_outputs_acceptance_sha256": "11436014ce1fecd54d5a349765e9ba435c372bee953984a828b48b83d0ff749f"
},
{
"phase": 3,
"arm": "baseline",
"raw_json_sha256": "d16148bc2da4f0f2fb46779806c675a996165c2b429a268b37a2db6e1a1dd8b6",
"request_outputs_cache_sha256": "61aa5c339936495f81b2b2657c72f6203b03f02f797e967dab7b5f563281db7d",
"decode_outputs_acceptance_sha256": "11436014ce1fecd54d5a349765e9ba435c372bee953984a828b48b83d0ff749f"
}
]
}
]
}
Activity