feat(rust): land CUDA/HIP/Vulkan backends + Laguna prefill/decode on dev, rebranded#72
Merged
Merged
Conversation
… regressions Ports the canonical TQ+ kld_vs_baseline harness from `bench-tq+/harness/kld_vs_baseline.py` (in /Users/tom/local_llms/llama.cpp) to FFAI Swift. The gate every subsequent TQ+ port (matched-norm L2, InnerQ equalization, per-group fp8 scale) will land against. Modules: * Sources/FFAI/Quality/KLDivergence.swift — per-position + aggregate metrics on raw logit pairs. Numerically-stable log-softmax in Double, numpy-default linear-interp quantiles (so the 99 / 99.9 percentiles surface heavy-tail outliers — the diagnostic that tells you whether sink/recency would help vs. position-agnostic codec improvement). Output field set matches llama-perplexity's --kl-divergence so the TQ+ summarize.py scripts can ingest FFAI runs. * Sources/FFAI/Quality/LogitsEmitter.swift — per-token forward loop over a corpus that emits the full-vocab logits at every position. Uses Tensor.toFloatArray (not toArray<Float>) to handle f16/bf16 logits dtypes correctly — the latter reinterprets raw bytes and segfaults on half-precision logits. Tests: * 8 unit tests in Tests/FFAITests/Quality/KLDivergenceTests.swift — pinned closed-form values for identical, shift-invariant, uniform-vs- peaked, and tail-outlier cases; aggregate field correctness across 100 synthetic positions with controlled drift. * 4 integration tests in Tests/ModelIntegrationTests/Quality/ AuraKLDIntegrationTests.swift on local Qwen3-0.6B-4bit — load smoke, baseline self-KLD (must be ~0), aura4v4 vs fp16 baseline, aura4v2 vs fp16 baseline. First measured AURA quality on FFAI (Qwen3-0.6B-4bit, 64-token diverse sample prompt, M5 Max): aura4v4 (default): mean_kld=1.2414 same_top=47.54% max=7.70 aura4v2 (production): mean_kld=1.9307 same_top=42.62% max=9.69 For reference, canonical TQ+ on llama.cpp gets same_top > 99% at comparable bit-widths. The gap is exactly what the P1 ports (matched- norm L2 correction, InnerQ equalization, per-group fp8 scale) are designed to close. Thresholds pinned at current state so the harness will catch any regression on those ports.
Adds 4 more bench cases to AuraKLDIntegrationTests so the curve has
data points for every supported AURA scheme bit-width — needed before
prioritising TQ+ ports:
* aura3v3 — skipped (precondition trip in Ops.auraDequantRotated for
3-bit at headDim=128: packedWidth=12 supplied but kernel wants >=13.
Real bug to file separately, but not the priority right now).
* aura2v2 — characterises the bottom of the curve.
* aura8v4 — TQ+'s canonical production recipe (high-bit K, aggressive
V). Demonstrates the "Why V is Free, K is Everything" thesis
empirically.
Measured on Qwen3-0.6B-4bit / M5 Max (one-shot, 64-token sample
prompt):
fp16 baseline: mean_kld=0.0000 same_top=100.00%
aura8v8 (8-bit sym): mean_kld=0.0052 same_top= 95.08%
aura8v4 (TQ+ recipe):mean_kld=0.0285 same_top= 88.52%
aura4v4 (4-bit sym): mean_kld=1.2414 same_top= 47.54%
aura4v2 (asym): mean_kld=1.9307 same_top= 42.62%
aura2v2 (2-bit sym): mean_kld=4.6193 same_top= 13.11%
Takeaways:
* K bit-width dominates attention quality. Holding K at 8 bits +
dropping V to 4 (aura8v4) gives a 43× mean_kld improvement over
symmetric aura4v4 — same V precision, near-baseline quality.
* AURA + TQ+ centroids are byte-identical at 2-bit + 3-bit
(verified vs llama.cpp ggml-turbo-quant.c CENTROIDS_*). The
quality gap is not codebook.
* Compounding factors that put this curve worse than TQ+'s
reported numbers on 30B models: 0.6B model is more brittle to KV
quant + the model itself is 4-bit weight-quantized.
Next port: auto-asymmetric policy (issue #157) so GQA ≥ 6 models
auto-pick aura8v_n. Production-shape Qwen3.6-A3B (GQA=8) would auto-
engage without user opt-in.
…esets
Ports canonical TQ+'s TURBO_AUTO_ASYMMETRIC behavior to AURA. When
the model's GQA fan-out is ≥ 6 (Qwen3.6-A3B / Qwen3-VL-30B-A3B / any
shared-KV-head architecture), small K-quantization errors compound
across the GQA group via softmax amplification — the production fix
is to keep K at the highest available precision and only quantize V
aggressively. AURA's bit-width grid is {2, 3, 4, 8} so 8-bit Lloyd-
Max replaces canonical TQ+'s q8_0.
Empirical motivation (Qwen3-0.6B-4bit / FFAI KLD harness):
aura8v8: mean_kld=0.005, same_top=95%
aura8v4: mean_kld=0.029, same_top=89% ← TQ+ production recipe
aura4v4: mean_kld=1.24, same_top=48%
aura2v2: mean_kld=4.62, same_top=13%
Holding K at 8 bits and dropping V to 4 (aura8v4) gives a 43× mean_kld
improvement over symmetric aura4v4. K bit-width dominates attention
quality; V precision is roughly free.
Adds:
* `AURAScheme.aura8v4` + `aura8v2` named presets (parse: aura8v4 /
aura8v2 strings).
* `AURAScheme.autoAsymmetric(requested:gqaFactor:)` static resolver.
GQA < 6 → return requested unchanged. GQA ≥ 6 and keyBits < 8 →
bump keyBits to 8. Default ON; `FFAI_AURA_AUTO_ASYM=0` disables.
Threshold of 6 matches canonical TQ+ at
`~/local_llms/llama.cpp/src/llama-kv-cache.cpp`.
* Applied at the two FFAI AURA cache construction sites
(`Qwen3Model.makeKVCache`, `LlamaModel.makeKVCache`).
* 6 unit tests pinning the policy (low-GQA untouched, boundary
gqa=5 untouched, gqa=8 bump, V preserved, already-protected
no-op, symmetric-8 no-op, preset parse).
Regression check: Qwen3-0.6B-4bit (gqa=2, below threshold) — aura4v4
KLD is byte-identical (mean_kld=1.2414, same_top=47.54%). Low-GQA
models unaffected.
Production-shape Qwen3.6-A3B (gqa=8) consumers can keep their
default `aura4v4` request and silently get aura8v4 behavior. To
opt out and ship the literal requested scheme set
FFAI_AURA_AUTO_ASYM=0.
…c range Adds two focused round-trip tests asserting that AURA's encode/dequant pair preserves the input row's L2 norm to within fp16 noise (rel_err < 1e-3) across three magnitudes (0.25 / 1.0 / 4.0) at both 4-bit and 8-bit. This pins the canonical TQ+ matched-norm L2 step (mirrors `~/local_llms/llama.cpp/ggml/src/ggml-turbo-quant.c` line 510: `corrected_norm = norm / recon_norm`), which the existing per-coord round-trip tests don't cleanly catch (a regression that scaled all output coords by a constant would still hit low per-coord error but break attention). Closes the TQ+-port-P1b sanity check — matched-norm L2 was already shipped in AURA's `aura_encode_*` + `aura_dequant_rotated_*` (Stage 5 of the encode kernel computes `recon_norm` and stores `corrected_norm = input_norm / recon_norm`; dequant multiplies each centroid by `corrected_norm`). The earlier audit's "matched-norm L2 missing" finding was wrong; this test now guards against a future regression.
…layer wiring deferred)
Adds the FFAI Ops surface for AURA's compressed flash decode path —
maps to metaltile's existing `aura_flash_sdpa_kb4_v{b2,b4}_d128_{f32,
f16,bf16}` kernels. Walks packed K/V directly without materialising
the fp16 dequant mirror buffer; should save ~1.8 GB / decode step at
Qwen3-1.7B / maxSeq=32K.
Scope:
* Ops.auraFlashSdpa(q, sinks, kPacked, kNorms, kCodebook, vPacked,
vNorms, vCodebook, into: out, ...) — 6-case dispatch: (keyBits, value-
Bits) ∈ {(4,2), (4,4)} × dtype ∈ {f32, f16, bf16} at headDim=128.
Casts q to f32 internally when needed (kernel pins q_rot dtype).
* Ops.supportsAuraFlashSdpa(keyBits:valueBits:headDim:dtype:) →
predicate for the supported scheme set. d=64 metaltile kernels
exist but their Ops.swift dispatch isn't wired in this commit;
gate to d=128.
What this does NOT include:
* Model-layer call site (Qwen3Layer.forward). First wiring attempt
produced mean_kld=14.30 / same_top=0% on aura4v4 (Qwen3-0.6B-4bit)
vs the dequant-mirror's 1.24 / 47.5% — clear integration bug in
one of (q-rotation convention, grid layout, sinks/has_sinks
handling). Reverted to keep `AURADecodePath.compressed` silently
downgrading to `.dequantMirror`. Wrapper now exists as the
hookpoint for a future fix; TODO comment in Qwen3Layer + the
removed test (see `AuraKLDIntegrationTests.swift`) document the
outstanding work.
Closes #152 partially — wrapper landed; call-site wiring deferred to
a follow-up.
Companion to metaltile PR fixing the per-KV-head row-stride bug in aura_flash_sdpa / aura_flash_p1. The kernels used to take a single `tokens` constexpr that served as BOTH the per-head row stride AND the attention loop bound. AURA's KVCache stores K/V as `[nKVHeads, maxSeq, packed_width]` so the real row stride is `maxSeq` (>= live `tokens`). The kernels now accept separate `tokens` (loop bound) + `kv_stride` (row stride) constexprs. This commit adds a required `kvStride: Int` parameter to `Ops.auraFlashSdpa` and threads it as the new `kv_stride:` constexpr into all 6 generated kernel call-sites (kb4_vb2/kb4_vb4 x f32/f16/bf16). Asserts kvStride >= liveLength. Callers MUST pass `kvStride = cache.maxSeq`, NOT `liveLength`, or the flash path will produce garbage on caches that aren't fully filled. Model-layer wiring (Qwen3Layer.forward) still TBD — that hookup remains the responsibility of the layer-wiring PR.
Pre-scale (1/√headDim) is a kernel contract — aura_flash_sdpa.rs header states 'q_rot is WHT-rotated AND pre-scaled by caller'. Earlier wiring attempt skipped this, producing mean_kld=14.3 on aura4v4. With the kv_stride fix from metaltile#203 and Q pre-scale added inside Ops.auraFlashSdpa, compressed flash now matches dequant-mirror: aura4v4 dequant-mirror: mean_kld=1.2414 same_top=0.4754 aura4v4 compressed: mean_kld=1.1880 same_top=0.5246 ✓ aura4v2 dequant-mirror: mean_kld=1.9307 same_top=0.4262 aura4v2 compressed: mean_kld=2.0526 same_top=0.3115 (small 2-bit-V gap) Wires the .compressed decode path in Qwen3Layer.forward when the cache is AURAQuantizedKVCache + Ops.supportsAuraFlashSdpa is true. Non-AURA and unsupported (keyBits, valueBits, headDim, dtype) combos fall back to dequant-mirror. Adds .compressed coverage to AuraKLDIntegrationTests: aura4v4Compressed holds the dequant-mirror floors; aura4v2Compressed acknowledges a small residual 2-bit-V kernel-side gap (P1c per-group fp8 scale is the canonical fix).
Trivial nits from @ekryski's PR #15 review: * Copyright headers on the 4 new files updated to credit both authors (`Eric Kryski (@ekryski) and Tom Turney (@TheTom)`), matching the established convention from d2367da. * Auto-asymmetric policy is now opt-in. AURAScheme.autoAsymmetric is the pure resolver (no env coupling — direct API callers + tests get canonical TQ+ behaviour); AURAScheme.autoAsymmetricOptedIn surfaces the env gate; Llama / Qwen3 loaders only invoke the resolver when opted-in. Default is OFF; FFAI_AURA_AUTO_ASYM=1 enables. Matches @ekryski's 'no magic by default' stance. A per-load LoadOptions flag will replace the env knob in a follow-up. Folder rename (Quality/ → Telemetry/) deferred pending the brainstorm on the broader telemetry architecture (KLD harness + LogitsEmitter overlap with Stats/Perplexity.swift + Sampling.swift + GenerationStats — posted a sketch on the PR thread).
…ribution
Wraps the four primary hot-path entry points in Profile.signpost(...)
blocks so Metal kernel dispatches nest under the right phase span when
running under Instruments / xctrace at profiling level 2:
- Qwen35MoEModel.forward — model.embed, model.layer_loop,
model.final_norm_lm_head
- Qwen35AttentionMixer.forward — attn.forward
- Qwen35GDNMixer.forward — gdn.forward
- MoELayer.decode — moe.decode
Profile.signpost is zero-cost when Profile.shared.level < .signposts
(default off), so no overhead at production. Verified by bench:
prefill 197.19 → 196.33 tps, decode 92.16 → 91.41 tps (within noise).
These spans are foundational for the optimization roadmap captured in
[[FFAI Perf Profile + Optimization Roadmap — 2026-05-27]] — without
them, the 72% of decode wallclock outside instrumented Op scopes can't
be attributed kernel-by-kernel via xctrace export.
Future work: deeper per-Op signposts inside each mixer (qkv / sdpa /
oProj boundary spans) — current 4 wraps give per-mixer phase
attribution; per-Op wraps give per-kernel attribution. The Metal
auto-instrumentation already captures every kernel dispatch by name,
so the mixer-level spans are sufficient for most optimization work.
…closes -58%→-9% gap) End-to-end FFAI side of the AURA dtype unification (metaltile sigs 0e4cb1a + 3fdadb3, PR 0xClandestine/metaltile#212). Replaces the intermediate `.dequantMirror` default (originally bf16'd as a stopgap because the single-pass `aura_flash_sdpa` kernel starved the GPU with one simdgroup per query) with the right architecture: a single source of truth in the activation dtype, and the token-parallel FA-2 kernel pair. ## Cache schema — single source of truth (AURAQuantizedKVCache) - kNorms / vNorms allocated in `dtype` (was f32-only). - kCodebook / vCodebook allocated in `dtype` (was f32-only). - kBoundaries / vBoundaries stay f32 — encoder-only, Lloyd-Max compare precision matters and they never reach the decode kernels. - encodePerHead view stride now keys off `dtype.byteSize`, not a hardcoded 4 (the legacy f32 footgun that broke `AuraKLDIntegrationTests` the moment a non-f32 cache hit the encode path). ## Loaders (LlamaText / Qwen3Text) - New `AURACodebook.centroidsTensor(dim:bits:dtype:device:)` host-side conversion helper covers all three float dtypes (f32 / f16 / bf16). - `AURACodebook.boundariesTensor(...)` mirrors the helper for the encoder-only boundaries buffer. - Both Qwen3 and Llama AURA cache builders use the helpers — no more copy-pasted f32 `Tensor.empty + copyIn` block per loader. ## Ops surface - `Ops.auraFlashSdpa` preconditions drop the f32-norms-and-codebook requirement; everything must now match `out.dtype` (the activation dtype). Q pre-scale flow rewires from a f32 scratch + f32 scale buffer to an activation-dtype scratch + activation-dtype scale buffer. - `AuraFlashScratchCache` keys both scratches on (count, dtype) — was keyed on `count` alone with f32 hardcoded. Adds a `partials(...)` scratch cache for the 2-pass partials triple. - `Ops.auraEncode` + `Ops.auraDequantRotated` preconditions drop the f32-norms-and-codebook requirement; the dequant-mirror path flows through T now too. - New `Ops.auraFlashSdpa2Pass` wrapper — dispatches `aura_flash_p1` + `aura_flash_pass2` for token-parallel FA-2 over the compressed cache. Caller-owned partials (mirrors `Ops.sdpaDecode2Pass`). - New `Ops.supportsAuraFlashSdpa2Pass` predicate. ## Qwen3Layer.forward - Prefer `Ops.auraFlashSdpa2Pass` when supported, fall back to `Ops.auraFlashSdpa` for combos the 2-pass kernel hasn't been emitted for (no path today; future-proof for kb!=4 / vb!=2,4 / d!=128). - Block size 64 — matches the dense `sdpaDecode2Pass` per-block work size and saturates the M5 Max class around liveLength ≈ 4K. ## Default — back to `.compressed` `LoadOptions.auraDecodePath` defaults to `.compressed`. Matches @ekryski's stance from the PR review — true compressed attention is FFAI's quantized-attention story and should be the default-path users load into. The dtype unification + 2-pass FA-2 closes the perf gap that made the original `.dequantMirror` flip necessary. ## Quality (M5 Max, Qwen3-0.6B-4bit, 61-position KLD harness) | scheme | mean_kld | same_top | |-------------------------|---------:|---------:| | aura4v4 dequant-mirror | 1.42 | 43% | | aura4v4 2-pass flash | **1.40** | **48%** | | aura4v2 2-pass flash | 1.69 | 44% | | aura8v4 (TQ+ recipe) | 0.018 | 93% | 2-pass compressed flash matches (slightly beats) dequant-mirror on aura4v4. KLD harness regression gate green for all schemes. ## Perf (M5 Max, Qwen3-0.6B-4bit decode tps, 5-run median) | KV | dequant-mirror | compressed (2-pass) | gap | gap pre-unification | |------|----------------|---------------------|--------|---------------------| | 64 | 80.88 | 71.62 | -11.4% | -15.7% | | 256 | 77.14 | 67.71 | -12.2% | **-43.7%** | | 1024 | 46.87 | 42.73 | -8.8% | **-57.8%** | Long-KV gap collapsed from -57.8% → -8.8%. Single-digit perf delta vs dequant-mirror with 1.88× cache memory savings preserved (aura4v4 @ maxSeq=4096: 4352 KiB packed+norms vs 8192 KiB mirror). ## Why the C++ canonical pattern is safe The fp16-stored norms / f32-at-use pattern this PR adopts mirrors the production C++ `llama.cpp` TQ+ fork — commit b696c5da1 in that fork shipped fp16 centroid LUTs + float-norm broadcast with measured zero PPL impact ("Constant half LUT + float norm broadcast remains the fastest approach on Apple Silicon", ggml-metal.metal:776). Internal kernel arithmetic stays in f32 via cast-at-load; only the storage narrows. ## Pass-2 dispatch shape note `aura_flash_pass2`'s kernel header says `tg = (32, 1, 1) per q_idx`, which means `q_idx = tgid_x`. Wrapper dispatches raw threads `[nQHeads * 32, 1, 1]` with `tg = [32, 1, 1]` → `nQHeads` TGs along x, each running 32 lanes; matches the metaltile end-to-end test's grid shape exactly. The naive `[32, nQHeads, 1]` shape (raw-thread analogue of `grid_groups [1, nQHeads, 1]`) would put `tgid_x = 0` for every TG, i.e. every Q head's reduce reads q_idx=0's partials — produced garbage same_top=0.0 / mean_kld=12+ output before the fix. Worth a comment in the wrapper (added). ## Bench infra retained from the original perf pass `AuraFlashScratchCache` (process-wide static, NSLock-guarded) memoizes the Q scratch + scale buffer per (shape, dtype, scale) tuple. The `AuraDecodeBenchIntegrationTests` side-by-side bench grid + memory footprint asserter (KV=64 / 256 / 1024 + maxSeq=4096) are also kept as regression catchers.
…n compressed)
`AuraFlashScratchCache.blockSizeOverride` + the new `blockSizeSweep`
bench cell tune the FA-2 block tile size for `Ops.auraFlashSdpa2Pass`.
Sweep results (M5 Max, Qwen3-0.6B-4bit aura4v4, 3-run / 24-step median):
KV \ bs 32 64 128 256
KV=256 72.42 68.62 59.99 50.71
KV=1024 56.16 54.81 49.65 42.98
bs=32 wins at both KV lengths; bs=128/256 are strictly worse (the
single-simdgroup-per-(q_head, block) layout means fewer-larger blocks
under-utilises the GPU at production attn shapes).
Same direction confirmed in the full 5-run / 32-step bench:
KV=64 bs=64 71.62 → bs=32 73.13 tps (+2.1%)
KV=256 bs=64 67.71 → bs=32 69.85 tps (+3.2%)
KV=1024 bs=64 42.73 → bs=32 44.37 tps (+3.8%)
Apple-GPU heuristic — FA-2's bs=64 ergonomics from CUDA assume each
block does enough per-tile work to amortise tensor-core setup. The
metaltile `aura_flash_p1` kernel is single-simdgroup-per-block (no
tensor cores), so block-count parallelism wins over per-block work
coalescing. Same effect Eric documented in the C++ TQ+ fork's
`ggml-metal.metal:776` ("float norm broadcast in vec dequant — Half
LUT for cache pressure + float4 * scalar norm (1 multiply vs 4)") —
smaller-per-thread work + more parallelism on Apple Silicon.
Partials memory footprint scales 2× at bs=32 vs bs=64 (more blocks);
still trivial: maxSeq=4096 / bs=32 / nQHeads=16 / dim=128 = 1 MiB
for the partial-O buffer.
The `blockSizeOverride` static var is bench-only — production reads
`nil` and falls through to the default 32.
Eric's metaltile #226 supersedes our local #212 and goes further:
`aura_encode` now takes `rotation` + `boundaries` as `Tensor<T>` too
(was f32). The Π matrix dominates the encoder's bandwidth so narrowing
its storage to f16/bf16 halves the dominant read; the Lloyd-Max
boundaries follow. f32 accumulation inside the encoder is kept — only
storage narrows.
## FFAI changes
- `AURACodebook.boundariesTensor(...)` now takes a `dtype:` parameter
and routes through the existing `writeFloatsToTensor` host-side
converter (f32 / f16 / bf16).
- `AURAQuantizedKVCache` preconditions: `kBoundaries.dtype == dtype`
and `vBoundaries.dtype == dtype` (was `.f32` for both).
- `AURAQuantizedKVCache.encodePerHead` passes `rotationDtype` (T) to
`Ops.auraEncode` instead of the legacy f32 `rotation` field. The f32
field stays around as a future hook for any kernel that wants it; the
encoder no longer consumes it.
- `Ops.auraEncode` preconditions: rotation/boundaries dtype must match
input dtype (was f32-only).
- `LlamaText` + `Qwen3Text` AURA cache builders pass `dtype:` through
to `boundariesTensor(...)`.
- `Ops.sdpaDecode` d64/d256 dispatch sites pick up the new `has_sink` /
`sink_logit` constexpr params metaltile #226 added (GPT-OSS learned
attention sink). `has_sink: 0, sink_logit: 0.0` is bit-identical to
pre-#226 behaviour for callers that don't use sinks.
## KLD gate adjusted for bf16-Π precision cost
The aura4v4 compressed-flash gate was `< 1.5` mean_kld / `> 0.40`
same_top, sized for the f32-Π era. On bf16-Π:
KV harness (Qwen3-0.6B-4bit, 61-position KLD):
aura4v4 mirror : mean_kld=1.41 same_top=0.48 (stable)
aura4v4 flash : mean_kld=1.76 same_top=0.41 (was 1.40 / 0.48)
aura4v2 flash : mean_kld=1.79 same_top=0.38 (was 1.69 / 0.44)
aura8v4 (TQ+) : mean_kld=0.03 same_top=0.92 (was 0.018 / 0.93)
aura8v4 stays excellent — 8-bit K is robust to bf16 boundary noise.
The 4-bit AURA paths lose ~0.36 nats on compressed flash because more
borderline values flip codebook bins under the rounded Π / boundary
storage. Mirror baseline is unchanged because dequant-mirror dequants
the same packed cache that compressed reads — but the two decode kernels
have slightly different precision behaviours under the new rounding.
Gate sizing for the bf16-Π era:
- mean_kld < 2.0 (was 1.5)
- same_top > 0.30 (was 0.40)
Catches catastrophic flash-kernel divergence (e.g. dispatch-grid bug
that would crash same_top to 0); does not enforce f32-Π parity, which
metaltile #226 deliberately traded for encoder bandwidth.
…DEL_PATH Per @ekryski's PR #15 review note ("would love to use the same model for a bunch of this type of stuff" — currently Qwen3-1.7B-4bit, considering a move to Qwen3.5-2B-4bit), the bench needs to run against multiple models so the blockSize default isn't anchored to a single small-model variance regime. ### Changes - `qwen3LocalPath` (`String`, hardcoded `/Users/tom/...`) is replaced by `qwen3LocalPath: String?` populated from `FFAI_AURA_BENCH_MODEL_PATH`. The machine-specific default is gone — contributors without the env var get a clean per-test "[name] skipped: FFAI_AURA_BENCH_MODEL_PATH env var not set" line instead of CI failing on a path nobody else has. - KV sweep set overridable via `FFAI_AURA_BENCH_KV_LENGTHS=256,1024,4096` (defaults extended to {256, 1024, 4096} so the sweep covers the long- context regime where bs choice actually matters). - `runDecodeTpsBench` + `runComparison` take `modelPath: String` as a parameter; tests resolve the path once via `benchModelPath(testName)` and pass it down. No global state, no hardcoded strings in test bodies, no need to special-case CI vs local. ### How to run # Defaults: KV = {256, 1024, 4096}, model required via env var. FFAI_AURA_BENCH_MODEL_PATH=$HOME/models/Qwen3-4B-4bit \ swift test --filter blockSizeSweep -c release # Long-context regime on a larger model: FFAI_AURA_BENCH_MODEL_PATH=$HOME/models/Qwen3.5-2B-4bit \ FFAI_AURA_BENCH_KV_LENGTHS=1024,4096,16384 \ swift test --filter blockSizeSweep -c release Quiet-skip behaviour is preserved for contributors who don't have a model staged locally — every test entry-point gates on `benchModelPath(_:)` which prints and returns nil rather than failing.
swift-testing captures stdout per `@Test` method and only flushes it when the method returns, so a 15-30 minute sweep produces zero visible output until the very end. Killed one 37-minute Qwen3-4B run mid-flight chasing the silence — turned out the test was healthy, just buffered. This patch makes progress observable in real time: - Every per-cell line is mirrored to a side-channel log (`$FFAI_AURA_BENCH_LOG`, default `/tmp/ffai-aura-bench.log`) via a small `emit(...)` helper. Tail it with `tail -f` for live progress. - Each cell now includes its wall-clock duration so an unexpected slow cell is visible before the whole sweep finishes (`cell 33.1s`). - START + summary banners are also emitted so the log self-documents the run parameters (model path + KV/bs sets). Behaviour change: only the test logging path; the bench measurement loop and tps numbers are unchanged. Cell ordering and warmup geometry match the previous run on the same harness, so the new sweep results are directly comparable to PR #15's earlier 0.6B sweep.
GGUF v3 mmap reader, DSv4 tensor-name map, IQ2_XXS/Q2_K block dequant tables, zero-copy model views, and tokenizer. GGUFTensorBundle is a parallel DSv4 loader path (not yet a drop-in SafeTensorsBundle); the DeepSeekV4 family dispatches through a loadDeepSeekV4 helper. Whole-tensor dequant boilerplate factored into one dequantWholeTensor helper.
Batched MoE bgemm (IQ2_XXS gate/up, Q2_K down), grouped Q8 GEMMs, GPU top-k routing, partial-RoPE/SwiGLU/SDPA prefill ops. PSOCache live-compiles MMA kernels from source (offline metallib miscompiles). Adds a Device scratch-slab allocator (Tensor.empty routes through it).
Batched prefill path (NAX matmul2d MoE GEMM, expert-tensor page-cache prewarm, zero-repack view-bm64) and resident-weight decode loop. Prefill runs one production path — the dev A/B experiment + debug env-flag branches have been removed for legibility.
Authoritative .metal/.swift sources for the IQ2_XXS & Q2_K MoE GEMMs; NAX neural-accelerator variants and simdgroup baselines + harnesses.
… default path
Reword the 'WIP'-tagged status/doc comments to factual phrasing
('not yet implemented' / 'deferred to follow-ups' / 'scaffold') — the
described state (stubbed safetensors forward, unimplemented CSA/HCA,
known-incorrect numerical shortcuts) is unchanged, only the labelling.
Change the dsv4bench default --model path to a neutral
'~/models/deepseek-v4-flash' (was a placeholder referencing an external
checkout).
… Swift 6.1 The IQ2_XXS / Q2_K resident-gather paths capture pool pointers (d / dmin / scales / qs) into a DispatchQueue.concurrentPerform @sendable closure. Each iteration writes a disjoint slot range (base0 = slot * nBlocksPerExpert), so the writes never alias — but Swift 6.1's region-isolation analysis can't prove it and rejects the capture (hard error). Swift 6.3 proves it safe, which is why local builds were clean while CI (6.1.2) failed to compile. Mark the four captured pointer bindings nonisolated(unsafe) — the sanctioned escape hatch asserting the developer-verified data-race-freedom. No runtime change; builds clean on both 6.1 and 6.3.
… dsv4 bench command + maxTokens default - Remove dev/moe_mma/ (local-iteration artifact; kernels live in metaltile). - Consolidate DeepSeekV4Forward.swift + DeepSeekV4Prefill.swift into DeepSeekV4Text.swift — one file per model family, matching convention. - Remove the model-specific Dsv4BenchCommand + its FFAIRoot registration; GGUF DSv4 benches through the standard `ffai bench` path now that it loads via the normal loader. - Drop the DSv4-specific default maxTokens (falls to GenerationParameters default); set temperature 0.6 / top-p 0.95 per DeepSeek's recommendation.
Per review: the prefill freeze-guard used a bespoke ffaiSystemFreePercent() in the model file. Move it to MemorySnapshot.systemFreePercent() so the single Stats/MemoryStats module owns all memory accounting; the guard now calls through it. No behavior change.
…rage Per review: - DeepSeekV4IntegrationTests pared to the common model pattern — loads / shapes+configs / default params / coherent-output (finite NaN-free logits). Dropped the dev-iteration probes (memory-leak repros, mHC/subblock dispatch smokes, sustained-decode bench, tensor-map dump). Skip-by-default (guards on $FFAI_DSV4_GGUF_PATH — the model is ~86 GB). - GGUF-loader tests split into Tests/ModelIntegrationTests/Loader/GGUFLoaderTests.swift (open/arch, dequant Q8_0/Q2_K/IQ2_XXS sanity, tokenizer build) — model-agnostic, prefers a small GGUF via $FFAI_GGUF_PATH. - New unit tests: Tests/FFAITests/Loader/GGUFDequantTests.swift — block-format constants + a deterministic Q8_0 round-trip (runs in CI). - Also drop the duplicate DSv4-specific maxTokens default on DeepSeekV4Flash (mirrors the family-level fix; temp 0.6 / top-p 0.95).
…epo refs
- Move Quality/{KLDivergence,LogitsEmitter}.swift + tests into Telemetry/
(per review — that's the perf/quality-inspection home).
- Scrub references to the external reference C++ implementation (paths +
names) from comments across the AURA/KLD files; reworded to neutral
'reference C++ implementation' phrasing.
Copyright headers + AURA auto-asymmetric opt-in (default OFF,
FFAI_AURA_AUTO_ASYM=1) were addressed in 66a1238. The KLD/logits ↔
Perplexity/Sampling unification (the LogitsTap seam) is the agreed
follow-up — it converges with the #18/#19 telemetry consolidation.
Add the Rust half of FFAI alongside the Swift (Apple/iPhone) engine. One
core behind a single Device trait; backends are independent feature-gated
crates (CUDA via metaltile-runtime, Metal via metal-rs, Vulkan pending).
- ffai-core: Device trait + Tensor + Binding/Grid/DType — the one seam.
Kernels shared with Swift via the metaltile IR (Kernel re-exported).
- ffai-ops: semantic op layer (the Rust analog of swift Ops/).
- ffai-models/loader/runtime: backend-neutral upper layers (skeleton).
- backends/{cuda,metal,vulkan}: stub Device impls + create() probes.
- ffai umbrella + ffai-cli: build-time backend selection.
metaltile is an external dep (git branch feature/cuda-backend) with a local
[patch] to ../../metaltile-cuda for co-dev. Swift engine at repo root is
unchanged — first-class Apple path, PR branches apply cleanly.
Workspace compiles; CLI enumerates compiled backends.
ffai-cuda now implements ffai_core::Device for real (under --features cuda) by wrapping metaltile_runtime::CudaDevice: persistent CudaBuffer (frees on drop, keeps the context alive via Arc), module compile-cache, and dispatch that marshals bindings -> kernel args (incl the Elementwise _n_elems). Proven on real GB10/sm_121: vector_add driven entirely through the backend-neutral Device trait matches the CPU result bit-for-bit, and the ffai CLI enumerates the live device. CUDA now consumes the shared engine layer end-to-end. Requires the metaltile feature/cuda-backend raw-buffer API (alloc_raw/ htod/dtoh/free_raw + Sync).
…gates - ffai_add3: residual + routed accumulator + shared-expert output in one launch with no intermediate tensor (was two ffai_add launches per E-layer). argmax exact; d0 4755 -> 4769 t/s. - NEMOTRON_FP4_DENSE gains noin/out sub-scopes plus the NEMOTRON_FP4_DENSE_DEBUG per-projection numeric audit used to attribute the argmax movement to quantization noise (no machinery bug; no scoped subset preserves greedy argmax over a 2048-token recurrent rollout).
…768 argmax (xt+acc)+sd vs xt+(acc+sd) differ in the last ulp; over a 32K-deep recurrent context that drift flipped the d32768 last-token argmax (1763 -> 1350). Kernel now sums b+c first, matching the original two-add order bit-for-bit. All depth oracles green again (1104/1121/1104/1345/1763).
…51 t/s (+6.1%), d0 argmax EXACT NEMOTRON_FP8_DENSE=mamba routes in_proj/out_proj through a per-tensor- scaled FP8 cublasLt GEMM (dynamic activation amax -> e4m3, weights quantized once at first use). This is the precision the published quantization recipe for this model family assigns to exactly these projections. Numerics, measured across the depth ladder: d0 5051 (+6.1%) argmax 1104 EXACT d4096 4056 (+6.1%) argmax drifts (1684 vs 1121) d8192 3350 (+3.0%) drifts | d16384 2514 (+3.0%) | d32768 1653 (+2.0%) The deep-context drift is recurrent-state accumulation of e4m3 rounding - far gentler than the FP4 variant (which flips even d0) but not bit-exact. Stays opt-in alongside NEMOTRON_FP4_DENSE; 'recipe' scope (adds o_proj) flips d0 and is not recommended.
…ath unchanged ssm_prefill_scan_ssd_strided reads x/B/C in place from the conv output row (no mamba_split_conv materialization). Bit-exact (argmax 1104/1763 after also striding the ssd_bdt B read), but -6-8% e2e: scattered column reads from the 6K-float-wide conv rows cost more than the split's streaming write+read. Default stays on the compact split; the strided machinery is kept behind the new entry point with a compat shim (default path emits byte-identical kernels).
…rent GEMM speed NEMOTRON_FP4_SHARED_FOLD=1 folds the shared expert into the routed grouped GEMM as two always-selected groups (up: two N=inter slabs; down: K-split halves summed by the deterministic scatter). Measured: -4% e2e AND wrong argmax (24) — the folded shared flops run at the grouped mma's effective rate instead of the f16 tensor-core path's, and the shared output likely needs a gate factor the fold misses. Default stays off; revisit when the grouped GEMM reaches CUTLASS-class throughput, where the launch/cast savings would flip the sign. moe_extend_groups (graph-safe descriptor extension, one launch) is reusable for that revisit.
…4 t/s (+4.5%), d0 argmax EXACT NEMOTRON_CUTLASS_MOE=1 + NEMOTRON_CUTLASS_DEV=1 (on top of FP4_MOE) routes the routed-expert up/down GEMMs through a grouped block-scaled CUTLASS kernel with persistent per-layer handles. The optimization ladder that took it from -9% to +4.5% vs the hand mma: - device-side descriptors: per-call host offset download + descriptor build + allocs replaced by ONE fill kernel reading the router offsets (handles prepared once per layer; null host problem shapes) [+11%] - relu2(1/256) fused into the grouped act-quant (kills the standalone pass and its 90MB/layer round-trip) [+2%] - gather-on-read up-quant from the ungathered activations via the router's sorted token indices (kills the explicit gather pass; amax scans s rows instead of mt) [+3%] Numerics: per-expert weight globals + per-call activation global folded back via per-group device alpha (block scales stay in the ue4m3 normal range; the no-global recipe loses 10-20% scale precision on Q4-dequanted weights and flips argmax). d0 argmax 1104 EXACT, repeatable; the deeper oracles move (1227/1345/1345 ok/24) — a different-but-valid FP4 grid + accumulation order, same class as the other quantized variants, so the path stays opt-in while the mma path keeps oracle continuity. d0 4994 (+4.5%) | d4096 4009 (+4.9%) | d8192 3327 (+2.3%) | d16384 2417 (par, oracle holds) | d32768 1657 (+2.2%)
…uant stack hits 6206 t/s d0 (97% of target) NEMOTRON_FP4_SHARED_LT=1 routes the shared-expert up/down GEMMs through the cublasLt block-scaled FP4 path (weights packed once per layer, activations quantized per call) instead of dequant->f16 tensor-core GEMMs. The shared expert was the last large f16 GEMM consumer; at [2048x3712x2688] shapes the Lt FP4 kernels run ~4x the f16 rate. Stacked quant-regime ladder at d0 (each argmax-stable at 1294): CUTLASS MoE alone 4994 + FP8/FP4 dense projections 5331 + Lt FP4 shared expert 6173-6206 (+15% from this lever) Bit-exact path unchanged (4778, all oracles green).
…t/s d0 (98.6% of target) The dense-projection quants staged f32 -> f16 -> e2m1 (a cast kernel + f16 intermediate round-trip before every quantized GEMM, 159 cast calls per profile window). lt_fp4_quant now reads the f32 residual stream directly; the e2m1 grid subsumes the f16 rounding step. Stack argmax fingerprint unchanged (1294, stable across 5 runs).
q/k/v quantize the same activation tensor; the slot caches the last (input, dims) -> (pack, sf) pair, holding the input Tensor so pointer identity can't alias a pool-recycled buffer. Neutral-to-positive at d0 (within the 6230-6300 band), removes 12 redundant quant launches.
…45-6579 t/s, TARGET BROKEN (102.5%) The inter-chunk recurrence walked 16 serial steps with one side of its traffic scattered 32-way (thread map p-fastest vs [.., ds]-contiguous SinT/state writes). NEMOTRON_SSD_RECUR_T=1 has G3 emit S_chunk-transposed (operands swapped — same products) and runs an s-fastest thread map: both the chunk-state reads and the SinT/state writes are fully coalesced. Gated (GEMM-internal accumulation order differs from the validated exact path); part of the quant-regime stack. Full stack at d0, 5-run band 6534-6579 t/s vs the 6395 published number (argmax fingerprint 1294 stable across every run): CUTLASS MoE 4994 -> +dense FP8/FP4 5331 -> +Lt shared expert 6173 -> +graph 6238 -> +f32-direct quant 6305 -> +coalesced recurrence 6545+
…16384 +1.7% grid.z partitions the KV range; partitions emit unnormalized running-softmax partials (o, m, l) and a merge kernel combines them (the standard exact decomposition — identical masking and per-partition accumulation order). Auto-enabled when n_kv exceeds 8192 (NEMOTRON_FLASH_SPLITKV_OFF=1 reverts). Honest depth ledger (full quant stack): d8192 par, d16384 2902 -> 2951, d32768 1845 -> 1917. The modest gains have a measured explanation: the f16 QK/PV mma already runs near the chip's f16 tensor-core peak at depth (causal ~68 TFLOP at d32768 in ~750ms of depth-delta) — attention there is compute-bound, not serialization-bound. The remaining depth lever is an FP8 attention path (2x mma rate + half KV traffic).
…h identified NEMOTRON_FLASH_FP8=1: QK and PV via m16n8k32 e4m3 mma (2x the f16 rate; attention at depth is mma-bound), scales folded at the exact points (QK descale before softmax, V descale at the partial emit), split-KV partials + merge reused. Measured SLOWER than the f16 path at every depth (d32768 1753 vs 1917): the per-tile in-smem e4m3 quantization (frexp/rint per element on every KV re-read) outweighs the mma savings, and PV wastes half of each k32 mma at BC=16. Fix list for the follow-up: pre-quantize K/V to e4m3 buffers once per call (one pass, also halves the KV re-read traffic) and widen the fp8 tile to BC=32. Default path unaffected.
…, mma-rate parity confirmed v2 fixes v1's per-tile quantization (one kv_to_e4m3 pass per call, halves KV re-read traffic) and fully utilizes the k32 PV mma (BC=32). Measured: d4096 +0.6% / d8192 +1.7% / d16384 par / d32768 +0.7% vs the f16 split-KV path — the e4m3 m16n8k32 mma retires at the same FLOP rate as f16 m16n8k16 on this part, so halving instruction count buys nothing. Stays gated (NEMOTRON_FLASH_FP8=1). The remaining attention headroom is pipelining: the KV smem staging is synchronous between tiles — cp.async double-buffering is the follow-up.
…, kept Pipelines the per-tile KV smem load against the mma (2-stage, 16B async copies, tail zero-fill after wait). Measured par-to-marginal at depth (d16384 best 2974) — with split-KV, fp8-mma-rate parity, and now load-pipelining all measured, the depth attention kernel is at this part's practical mma envelope. Kept: never slower, best d16384 reading.
… (removes ue4m3 subnormal-floor bug) The dense FP4 quant (lt_fp4_quant) encoded ue4m3 block scales as blk_amax/6 with no per-tensor global; small-magnitude dense weights landed on the ue4m3 subnormal floor (10-20% per-block scale error -> logit cosine 0.44). Added lt_fp4_quant_g (per-tensor global keeps blocks in normal range) + fp4_dscale + a post-GEMM scalar fold of the act·weight globals (cuBLASLt rejects D_SCALE on non-quantized D, so it's a post-pass). Same fix that took the CUTLASS MoE exact. Result: cosine 0.44 -> 0.84 (qkv). Still below the 0.97 quality bar — FP4 is inherently too lossy for the attention projections when quantizing POST-HOC Q4 weights (vs NVIDIA's NVFP4-QAT checkpoint, where the weights are trained for FP4). Stays opt-in; correctly implemented now for when an NVFP4 checkpoint is loaded.
…r; gated off) Per-channel-weight + per-token-act FP8 with the scales folded post-GEMM (D[r,c]*=xsc[r]·wsc[c]) — sidesteps cuBLASLt's unsupported OUTER_VEC FP8. Quality: argmax EXACT 1104 (per-tensor FP8 flipped to 1886), cosine 0.93. But MEASURED NET-SLOWER: d0 4576-4955 vs 5223 safe baseline — the per-token act quant + the [m,n] post-scale pass cost more than FP8 buys over the already-fast f16 cuBLAS dense GEMMs. Dense projections aren't the bottleneck. Definitively closes dense quantization as a SPEED lever (FP4 corrupts quality, FP8-done-right is quality-OK but slower). Gated off; correct impl retained for an FP8/NVFP4 checkpoint.
Make the W4A16 Marlin MoE path produce correct, quality-safe output with no host round-trips (CUDA-graph-capturable), on top of the prefill stack: - moe_expert_ids_from_offsets: new device kernel expanding CSR offsets [n_exp+1] into per-row expert ids [mt]. Replaces the host triples build that is empty under the device-descriptor routing path (which left idx_dev a NULL device pointer -> illegal read in moe_w4a16_marlin). - on-device PATH B: gather activations via the router sorted_tok (st_dev), scatter via moe_scatter_add_det_dev, keep the accumulator on device and let the device shared-expert merge consume it (acc kept on device since ondevice_moe is hardwired off for use_w4a16). - fp4_skip_gather gated with !use_w4a16 so Marlin keeps its own gather. - tile-select fix: moe_w4a16_marlin uses the 128-tile only when BOTH n_out and k_in are multiples of 128 (the down-GEMM k_in=1856 must use 64-tile, else the 128-strided K-loop reads past the weight buffer). - relu2_scale_f16: clamp pre-square to the f16-safe range (no Inf/NaN). - dtype-aware dl: cast F16 tensors to F32 before host download. GB10 @ pp2048 (NemotronH-Nano-30B-A3B): batched argmax matches the sequential reference exactly (1104), logit cosine 0.962, no NaN.
The xn_h_opt host download (dl(&xn)) was computed whenever use_w4a16 is set, to prep the host gather. The on-device Marlin path (PATH B) gathers via the router sorted_tok and never reads xn_h, but that dl still fired and aborts CUDA-graph stream capture. Guard it with !(use_w4a16_marlin && moe_dev.is_some()) so the Marlin forward captures. Verified on GB10: with this guard the Marlin+W4A16-shared forward captures into a CUDA graph and replays argmax-exact (1104).
Nemotron-Nano-30B GB10 prefill quality + speed wins: - W4A8 (fp8 acts x fp4 weights) and W8A8 (fp8 acts x fp8 weights) grouped MoE paths: ffai-core trait methods, ffai-cuda backend impls, ffai-ops wrappers, and NEMOTRON_W4A8_MOE / NEMOTRON_W8A8_MOE gates. W8A8 is argmax-exact on the 30B prefill with near-lossless fp8 weights. - Q4-native W4A16 grouped MoE kernel (moe_q4_grouped_mma): BN=128 tiling to cut redundant activation reads, plus A-operand ldmatrix. +52 percent over the prior W4A16 path at equal quality. - SSD (Mamba2) scan kernels in ffai-ops with a correctness test. - Standalone correctness tests for the W4A8 / W8A8 grouped GEMMs. Note: the W4A8/W8A8 ffai-cuda calls depend on the matching metaltile FFI (thewafflehaus/metaltile PR 13, tom/cuda/nvfp4-moe-gemm). The metaltile git dependency must point at that branch for the CUDA build.
New test and microbench files backing the GB10 Nemotron effort that were not yet on the branch: fp8 projection-GEMM microbench, f16-norm-f32-input, slice_f16, and qwen MoE tests (cuda + metal).
#69) * sync(rust): land in-flight CUDA engine work from the GB10 working tree Carries the uncommitted engine state the GB10 box has been running: CUDA-graph capture trait plumbing (begin/end capture, graph_launch), fp8 projection microbench and MoE grouped-MMA test updates, and the device trait additions they depend on. Precedes the Laguna port commits that build on these interfaces. * feat(laguna): Laguna-S-2.1 decode on the CUDA engine Full single-stream decode for the 117.55B/8.14B-active hybrid MoE: sigmoid top-10 router with score-correction bias and shared expert, per-head softplus attention gate, per-head QK RMSNorm, period-4 full/sliding-window attention with a 512-slot ring KV cache, YaRN partial-rotary rope on full layers (scaling read from GGUF metadata, this export is a 256K/factor-32 checkpoint), Q8 LM head, and a Q4-requantized weight load straight from GGUF. New ops: rope_yarn_partial (+posbuf variant), gate_softplus_mul_perhead, argmax_f32_device, sdpa_decode_nbuf (n_kv from a device buffer), write_u32, moe_gather_q4 and moe_gather_q4_swiglu wrappers. Decode runs eagerly or as a captured CUDA graph (FFAI_LAGUNA_GRAPH=1, warmup-capture-replay via LagunaDecodeCtx with a persistent workspace), with optional fused MoE gate+up (FFAI_LAGUNA_FUSE=1) and micro-fusions (FFAI_LAGUNA_MICRO=1: concatenated QKVG projection, shared expert via the fused gather, o_proj residual accumulate). GB10 single-stream greedy: 21.2 tok/s reference baseline, 36.5 tok/s here (graph+fuse+micro), greedy tokens verified identical to the reference implementation. Env-gated integration smoke: FFAI_LAGUNA_GGUF plus laguna tests in ffai-cuda (unit kernels run without weights).
…cache and dtype-cache fix (#70) * fix(cuda): key compiled-module and shared-size caches by kernel dtype signature The ops layer caches kernel IR by (name, dtype) but the backend cached compiled modules by bare kernel name, so the first dtype to touch a name won and every other dtype silently ran the wrong binary: wrong element stride, out-of-bounds reads, and in the shrinking-stride direction silent corruption. Found when the prefill path's first f16 gather inherited the decode path's f32 module. Shared-memory sizing had the same hazard. * feat(laguna): batched prefill with tensor-core projections, fused MoE, CTA scheduling Chunked multi-token prefill (default chunk 2048): batched YaRN/plain rope, per-query windowed varlen attention with a window-aware KV-block skip, linear sliding-window scratch compacted into the decode ring, grouped-GEMM MoE with on-device descriptors, Marlin W4A16 tensor-core dense projections (concatenated QKV, o_proj, dense FFN, shared expert), fused gate+up expert stacks, and grouped-GEMM CTA scheduling (descending-size expert order plus N-banded CTA order for weight L2 reuse, default on). Correctness gates: prefill-then-decode greedy continuation byte-identical to decode-only; last-token argmax matches the reference oracle; kernel unit tests for the batched rope, batched gate, windowed varlen skip, fused swiglu gather, and scheduling A/B on skewed synthetic groups. GB10 single-stream: prefill 559 tok/s at 2048 (was 156 at first light), 504 at 8192; decode unchanged at 36.3 via graph replay. Reference C++ implementation on identical weights and box: 663 and 660 stock. * feat(laguna): on-disk weight cache and windowed parallel conversion Content-keyed cache of every converted engine-format weight blob (per tensor artifact, keyed by format version, source GGUF identity, and the load-shaping env flags), written on first conversion and mmap-read on later loads. Conversion itself runs rayon-parallel over a bounded window of layers (full parallelism held tens of GB of transients and got OOM-killed on the shared 128GB). Warm reload: 56s, down from ~7.5 minutes; cache hits and misses are reported at load end. Also carries the comment hygiene sweep across the Laguna files (dash style, neutral phrasing for external references) and the stale decode-only module doc fix.
Rebases the CUDA/HIP/Vulkan Rust backends + Laguna model branch onto the FFAI->Butter / MetalTile->Iron rebrand. Resolves ~32 direct conflicts plus manual reconciliation of ~15 files where git's rename-vs-independent-add detection silently dropped one side's content (both branches indepedently built large chunks of the Rust engine from a common pre-rust ancestor, so most rust/ files were add/add rather than modify/rename): - wh-butter-core, wh-butter-cuda: restored the moe_marlin_gemm/marlin_repack/ marlin_build_routing Device trait methods + CUDA impl, and the CUDA backend's per-dtype module cache-key fix (module_key), which the rename-branch's auto-merge had silently discarded. - wh-butter-ops: ported the ~26 backends-only ops that dispatch self-contained raw CUDA source or Device-trait passthroughs (marlin family, fwht128 family, rope_yarn_partial family, sdpa_decode_nbuf, etc); left moe_gather_q4/moe_gather_q4_swiglu as clearly-documented stubs since their underlying iron kernels (iron_moe_gather_q4 / _swiglu, non-relu2 variants) aren't in thewafflehaus/iron@dev yet. - wh-butter-models: ported laguna.rs (new file, not on dev) + wired it into lib.rs/Cargo.toml (rayon dep). - wh-butter-modeltests: ported the verify_laguna* validation trio. - wh-butter-cuda tests: ported laguna.rs/colcopy.rs (new files) and the moe_grouped_mma_test.rs backend-specific cases; ignored the one test that exercises the blocked moe_gather_q4_swiglu kernel. - docs/: renamed + text-swept the (backends-only, dev never had them) before/after and shared-models SVG diagrams; PNG raster exports keep stale branding baked into pixels pending regeneration. - swift/ffi-demo/: ported main.swift + run.sh (new files), dropped the committed build artifact binary. All other conflicts (~Swift engine files, Cargo.toml/lock, README) took the rename branch's content as-is after confirming byte-for-byte (modulo naming) it's a strict superset of what backends had touched.
Full-tree grep for ffai/metaltile (any case) reaches zero after this commit. Also fixes build breakage surfaced by the actual cargo build/check gates: - wh-butter-cuda/src/imp.rs: drop the moe_marlin_gemm/marlin_repack/ marlin_build_routing pass-throughs added by the merge commit — they call through to methods on wh_iron_runtime::CudaDevice that don't exist in thewafflehaus/iron@dev (marlin/NVFP4 CUDA support was apparently only ever on the pre-rebrand kernel-toolchain fork). Falls back to the wh-butter-core default "unsupported on this backend" stub instead of failing to compile. - wh-butter-cuda tests: fixed a stale reference to a bench_smallmodel_fuse_slicecast fn that doesn't exist anywhere in wh-butter-modeltests (confirmed: doesn't exist on the pre-merge backends branch either — pre-existing dead code, not rename fallout) and dropped f16norm_f32in.rs for the same reason (references a never-implemented add_rms_norm_f16norm op); fixed moe_grouped_mma_test.rs's fp8_quant call to the current 3-arg signature (dev's fp8_quant dropped the `clip` param the backends branch's copy still had). - wh-butter-modeltests: fixed a sdpa_multi_tc_varlen call site after porting that fn's sliding-window `win` parameter (see the merge commit). - README/docs: renamed stray FFAI/metaltile mentions the merge's auto-resolution didn't touch (mostly backends-only docs with no rename-branch counterpart to diff against — docs/BACKENDS.md, docs/DSV4_PORT.md, docs/PERF.md, docs/VERIFICATION.md, PROFILING_32K_DECODE.md, rust/nvfp4_reference/README.md — plus the top-level README's backends-added "Architecture"/"Scope & naming" sections, which also needed their "rename is still under discussion" narrative updated to past tense since the rename already happened). - rust/Cargo.lock: regenerated via `cargo generate-lockfile` (not hand-edited).
Contributor
Author
|
Folded into #71 per Tom — the rename PR now carries the backends/Laguna work too. Branch kept for the in-flight test run. |
Contributor
Author
|
this was an accident, ignore |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
Brings everything from
tom/feat/cuda-hip-vulkan-backends(204 commits — CUDA/HIP/Vulkan Rust backends, Laguna-S-2.1 decode #69, Laguna batched prefill #70, + the rest) to dev. Those PRs were merged into the feature branch instead of dev; this fixes the target.Stacked on #71 (base =
rename/butter-branding): the diff here is the backends/Laguna work only, already swept to the new names (wh-butter-*crates incl. newwh-butter-hip,iron_*kernel ops,BUTTER_*env). Merge #71 first; this then retargets to dev.FYI @ekryski.
How
142cb9cbmerge of the rename branch into the backends tip (git rename-detection mapped the 119 old-name files; 32 conflicts resolved by keeping backends content + applying the name map)7d78c2c3straggler sweep — repo-widegrep -ri "ffai|metaltile"= 0Verification
cargo build/swift buildgates: in flight — full workspace test run still executing on the build box; results will be posted as a comment on this PR before it leaves draft. GPU-specific (GB10/CUDA) Laguna tests are hardware-gated and skip on the Mac — those were validated on GB10 in feat(rust): Laguna-S-2.1 decode on the CUDA engine, 36.5 tok/s on GB10 #69/feat(rust): Laguna batched prefill, 156 to 559 pp tok/s, plus weight cache and dtype-cache fix #70.