# Changelog

All notable changes to this project will be documented in this file.

The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.1.0/),
and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html).

This changelog should record more than feature additions. It should also record:

- breaking changes
- deprecated surfaces
- support-matrix changes
- migration notes when user action is required

Related governance docs:

- [docs/stability-policy.md](docs/stability-policy.md)
- [docs/support-matrix.md](docs/support-matrix.md)

## [Unreleased]

### OPD train (CUDA) — new beta surface

- **End-to-end OPD CUDA training stack landed on Qwen3-0.6B.** Single-session
  32-commit arc through kill-or-license-gated wins brings the OPD step at the
  moderate Qwen3.5-like shape to **48.5 ms** on RTX 4070 Ti SUPER —
  **1.71× faster than the like-for-like PyTorch CUDA reference (83 ms)** —
  and the real Qwen3-0.6B checkpoint OPD step to **0.164 s/step** (~170×
  over a naive scratch CPU baseline). CPU/CUDA loss bit-equivalent to
  relerr 1.276e-6. Convergence verified at lr=1e-7 with held-out
  exact-overlap **50 → 82.8 %** by step 5000 (KL/NLL still monotonically
  falling). Five parallel axes killed cleanly via SOLID gates with
  recorded errors entries (forward_last_logits, merge_grad sharing, SDPA
  mask-softmax fusion, high-level CUDA Graph rollout capture, SwiGLU
  silu+multiply fusion). New CUDA op surfaces: `matmul_bt` forward +
  backward, in-place AdamW, KV cache for OPD rollout, device-resident
  RoPE / argmax, fused causal-SDPA decode, fused attention-prepare
  layout, fused grad clip. Usage manual:
  [`docs/projects/2026-05-21-arle-opd-cuda-usage-manual.md`](docs/projects/2026-05-21-arle-opd-cuda-usage-manual.md).
  Cycle wrap:
  [`docs/projects/2026-05-21-opd-cuda-cycle-wrap.md`](docs/projects/2026-05-21-opd-cuda-cycle-wrap.md).
  Industry positioning:
  [`docs/projects/2026-05-21-opd-industry-positioning-best-framework.md`](docs/projects/2026-05-21-opd-industry-positioning-best-framework.md).

### Observability

- Added low-overhead HTTP `request_trace` JSON summaries for streaming and
  buffered requests, including TTFT, total latency, token throughput,
  KV/prefix-cache state, scheduler phase EMA, pipeline, and preprocess
  snapshots. Added `scripts/bench_dsv4_trace_http.py` to run DSv4 HTTP smoke
  cases and collect matching `request_trace` entries from server logs without
  enabling CUDA-synchronizing per-layer tracing.
- Fixed DSv4 distributed HTTP submissions so concurrent client requests keep
  the same logical queue order on every rank. `DistributedSchedulerGroup` now
  serializes cross-rank fanout submission, preventing rank 0 and follower ranks
  from entering different per-request token coordinators under concurrent
  traffic.
- Allowed DSv4 decode to run scheduler batches larger than one via the existing
  per-slot decode path. This keeps multi-slot distributed HTTP fanout alive
  while the vectorized DSv4 B>1 decode kernel work remains pending.
- Added DSv4 HTTP TP/EP axis overrides through the existing `INFER_TP_SIZE`
  / `ARLE_TP_SIZE` and `INFER_EP_SIZE` / `ARLE_EP_SIZE` env vars. The default
  remains the legacy overlapping TP=world, EP=world layout. The first 8xH20
  profiling pass confirms the current runnable DSv4 layout is decode
  communication-bound: default TP=8/EP=8 performs 86 all-reduces per generated
  token per rank, and nsys observed 22016 NCCL all-reduce kernels for a
  32-token decode window. Evidence and industry comparison are recorded in
  [`docs/experience/errors/2026-05-14-dsv4-decode-nccl-bottleneck.md`](docs/experience/errors/2026-05-14-dsv4-decode-nccl-bottleneck.md).
- Added committed DSv4 trace artifacts under
  [`docs/trace-artifacts/2026-05-14-dsv4-decode/`](docs/trace-artifacts/2026-05-14-dsv4-decode/),
  including the compressed raw nsys report/database, `nsys stats`, client JSON,
  server log, and SHA256 manifest. The trace record no longer depends on remote
  `/tmp` files.
- Added DSv4 DeepEP MoE trace artifacts under
  [`docs/trace-artifacts/2026-05-14-dsv4-deepep/`](docs/trace-artifacts/2026-05-14-dsv4-deepep/),
  including compressed BF16 and FP8 combine trace logs, parsed summaries, remote
  build evidence, default trace-off post-checks, and the current bottleneck
  callout for return-side combine exchange plus local expert GEMMs.
- Added a current 8xH20 DSv4 single-token Nsight trace under
  [`docs/trace-artifacts/2026-05-14-dsv4-deepep/nsys-one-token-current/`](docs/trace-artifacts/2026-05-14-dsv4-deepep/nsys-one-token-current/).
  The `max_tokens=2` streaming request returned `霓灯` and produced exactly one
  `step_decode_kernel_launch` wave across 8 ranks. The isolated token takes
  266.020 ms wall; decode-only nsys shows `cuStreamSynchronize`,
  async allocation/free, launch/memset churn, and NCCL send/recv ahead of the
  actual attention and GEMV kernels.
- Added a refreshed 2026-05-15 DSv4 single-token Nsight trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-one-token-current/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-one-token-current/).
  With send/recv route and route-logits scratch reuse in place, the same
  one-token decode shape is now 158.439 ms wall. The remaining ranked costs are
  async allocation/free, launch/memset churn, D2H route readbacks, NCCL
  SendRecv/AllReduce, and local expert FP8/FP4 GEMV.
- Added 2026-05-15 DSv4 padded-dispatch Nsight records under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/).
  The negative first trace (`nsys-single-token-padded-dispatch`) shows that
  padding without removing the dead send-count kernel regresses to 136.908 ms;
  the fixed trace (`nsys-single-token-padded-dispatch-skip-count`) validates the
  shipped B=1 decode path at 123.955 ms and records the remaining ranked costs.
- Added the DSv4 padded peer-combine Nsight trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-token-padded-peer-combine/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-token-padded-peer-combine/).
  The real 8xH20 run keeps the `霓彩` output and shows the single-token decode
  wave at 112.133 ms after pre-summing padded return rows per origin peer.
- Added the DSv4 fused dispatch payload Nsight trace and matching HTTP smoke
  under [`docs/trace-artifacts/2026-05-15-dsv4-deepep/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/).
  The real 8xH20 run keeps the `霓彩` output, cuts decode-window SendRecv
  launches from 1,032 to 688 by exchanging hidden rows and route metadata in
  one BF16 payload, and records a fresh isolated single-token decode wave at
  118.985 ms. The trace-off `decode64` smoke returns normal English content at
  12.22 post-first tok/s and the arithmetic case returns `410`; the nsys run
  makes clear that NCCL exchange/reduction, launch overhead, allocator churn,
  D2H, and local expert GEMV still dominate.
- Added the DSv4 route-grouped pair GEMV Nsight trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-route-pair-gemv/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-route-pair-gemv/).
  The opt-in `ARLE_DSV4_ROUTE_GROUPED_EXPERTS=1` run keeps the `霓彩` output
  and measures a 117.894 ms single-token decode wave. The decode-window top
  costs are now explicit: `ncclDevKernel_SendRecv` at 50.338 ms per rank
  range, FP4 route pair GEMV at 19.616 ms, FP4 route `w2` GEMV at 10.487 ms,
  FP8 GEMV at 9.408 ms, plus allocator/free and launch overhead.
- Added a fresh user-requested single-token `nsys` rerun under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-current-user/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-current-user/).
  The real 8xH20 `/root/DeepSeek-V4-Flash` run returns exact arithmetic `406`
  and measures a 94.841 ms decode wave. The slow stack is reduce-scatter
  combine, local FP8/FP4 expert GEMV, residual all-reduce/send-recv,
  attention/MHC/route kernels, and high per-token launch/alloc/free/D2H
  runtime overhead, not sampler time.
- Added the matching DSv4 single-token `NCCL_PROTO=LL128` negative trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-nccl-ll128/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-nccl-ll128/).
  The arithmetic request still returns `406`, but the isolated decode wave is
  94.936 ms versus 94.841 ms on the current default reference, and
  reduce-scatter combine is slightly worse at 21.371 ms per rank-range.
  Protocol selection alone is therefore not the next default decode fix.
- Added an opt-in DSv4 return-combine overlap experiment behind
  `ARLE_DSV4_COMBINE_OVERLAP=1`. The path creates a second EP NCCL
  communicator on a dedicated communication stream and delays routed-output
  consumption with an explicit CUDA fence so shared expert compute can overlap
  the reduce-scatter. The real 8xH20 run returns exact arithmetic `406`, but
  the trace regresses from 94.841 ms to 104.359 ms because all-reduce timing
  and cross-stream event overhead outweigh the reduce-scatter improvement.
  The matching default-off HTTP smoke still reaches 12.05 post-first tok/s,
  so the overlap experiment remains disabled by default.
- Added a fused DSv4 B=1 padded DeepEP local expert prepare kernel and matching
  trace records under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-small-local-pack-prepare/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-small-local-pack-prepare/)
  and
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-small-local-pack-prepare-smoke/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-small-local-pack-prepare-smoke/).
  The real 8xH20 run returns exact arithmetic `406`, cuts H2D runtime calls
  from 1,040 to 696, cuts `cuMemsetD8Async` calls from 1,232 to 544, and keeps
  trace-off `decode64` at 12.05 post-first tok/s. The single captured nsys wave
  is 92.602 ms due to noisier D2H/AllReduce timing, so this is recorded as
  small-call cleanup rather than a wall-time win.
- Reused per-layer DSv4 incremental attention projection buffers for `c_q`,
  `c_q_normed`, `q_raw`, `kv_raw`, and `kv_normed`. The real 8xH20
  single-token `nsys` run returns exact arithmetic `406` and moves the decode
  wave from 94.841 ms to 90.946 ms, while `cuMemAllocAsync` calls drop from
  6,760 to 5,040 and `cuMemFreeAsync` calls drop from 3,048 to 1,328 inside
  the decode range. The matching HTTP smoke keeps normal Chinese/English
  streaming output and exact math, with `decode64` at 11.89 post-first tok/s.
- Added a direct current-path single decode-token Nsight breakdown under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-current-breakdown/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-current-breakdown/).
  The real 8xH20 `/root/DeepSeek-V4-Flash` run returns exact arithmetic `406`
  and measures a 105.205 ms isolated second-token decode wave. The top stack is
  now explicit: 16,177 CUDA launches, reduce-scatter combine, local FP8/FP4
  expert GEMV, all-reduce, attention/MHC/route kernels, and 347 D2H calls for
  per-layer synchronization. The actual D2H activity payload is only 44,044
  bytes, confirming the current issue is MoE communication/compute plus
  launch/runtime synchronization granularity, not sampler time or copy bandwidth.
- Switched additional full-write DSv4 runtime scratch buffers from zeroed
  allocation to uninitialized allocation: expert/shared/grouped hidden scratch,
  route logits, per-layer hidden scratch, and MHC parameter scratch. The real
  8xH20 single-token `nsys` trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-expanded-uninit/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-expanded-uninit/)
  returns exact arithmetic `406`, moves the isolated decode wave from
  105.205 ms to 88.554 ms, and cuts `cuMemsetD8Async` from 3,640 calls /
  6.932 ms per rank range to 1,920 calls / 2.839 ms. The trace still points at
  reduce-scatter combine and local FP8/FP4 expert GEMV as the main bottlenecks.
  The matching HTTP smoke under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-expanded-uninit-smoke/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-expanded-uninit-smoke/)
  keeps normal Chinese/English multi-token output, exact math `410`, and
  `decode64` at 11.94 post-first tok/s.
- Extended the DSv4 uninitialized scratch cleanup to MoE dispatch, payload,
  recv/local-route, active grouped, and combine buffers. The real 8xH20
  single-token `nsys` artifact under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-moe-scratch-uninit-rerun/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-moe-scratch-uninit-rerun/)
  returns exact arithmetic `406`, moves the isolated decode wave from
  88.554 ms to 87.667 ms after a rerun, and cuts `cuMemsetD8Async` from
  1,920 calls / 2.839 ms per rank range to 1,232 calls / 1.558 ms. The
  matching HTTP smoke under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-moe-scratch-uninit-smoke/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-moe-scratch-uninit-smoke/)
  keeps normal Chinese/English multi-token output, exact math `410`, and
  `decode64` at 12.06 post-first tok/s.
- Moved DSv4 grouped expert weight/scale pointer tables into
  `DeepseekV4MoeBlock` load-time caches for the opt-in grouped/route-grouped
  expert paths and future raw-pointer DeepGEMM integration. On the real 8xH20
  `ARLE_DSV4_ROUTE_GROUPED_EXPERTS=1` trace, exact arithmetic remains `406`,
  H2D activity drops from 1,918 calls / 374,752 bytes to 440 calls / 7,808
  bytes, H2D runtime drops from 5.490 ms to 1.380 ms, and the route-grouped
  single-token wave moves from 105.808 ms to 94.828 ms. The path remains
  default-off because reduce-scatter combine and route-wise FP4/FP8 GEMV still
  dominate; the default DeepEP smoke still returns math `410`, normal Chinese
  writing, and normal English decode text.
- Added the DSv4 default-path warm decode Nsight trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-default-warm-decode/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-default-warm-decode/).
  The run warms a real decode first, then profiles a second single decode token
  on 8xH20. The output remains `霓彩`, the decode wave is 128.130 ms, and the
  trace confirms allocator/free overhead is steady-state rather than only
  first-decode initialization: the decode window still records 8,453
  `cuMemAllocAsync` calls and 6,048 `cuMemFreeAsync` calls. The slow stack is
  NCCL SendRecv/AllReduce, local FP8/FP4 expert GEMV, launch/runtime overhead,
  allocator/free churn, and route-count D2H synchronization.
- Added the DSv4 expert-wise grouped GEMV negative Nsight trace under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-expert-grouped/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-expert-grouped/).
  With `ARLE_DSV4_GROUPED_EXPERTS=1`, the real 8xH20 run keeps the `霓彩`
  output but regresses the warmed single-token decode wave to 145.693 ms.
  The trace shows `ncclDevKernel_SendRecv` at 58.049 ms per rank range, FP4
  grouped gate/up GEMV at 23.162 ms, FP4 grouped `w2` GEMV at 11.428 ms, and
  elevated route-count D2H synchronization. This confirms the opt-in grouped
  GEMV path remains default-off and that the target remains true grouped
  GEMM/DeepGEMM with DeepEP overlap.
- Added the DSv4 route-grouped pair trace-off HTTP comparison under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-route-grouped-pair-vs-default/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-route-grouped-pair-vs-default/).
  Default fused-dispatch decode keeps `decode64` at 11.47 completion tok/s and
  arithmetic at `410`; `ARLE_DSV4_ROUTE_GROUPED_EXPERTS=1` returns normal text
  and the same arithmetic answer but regresses `decode64` to 6.54 completion
  tok/s. Route-wise grouped GEMV remains default-off.
- Added DSv4 incremental stream scratch recycling and captured both the HTTP
  smoke and Nsight follow-up under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-stream-recycle/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-stream-recycle/)
  and
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-stream-recycle/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-stream-recycle/).
  The real 8xH20 run keeps normal text and arithmetic `410`; the isolated
  warmed decode wave improves from 128.130 ms to 111.798 ms, with
  `cuMemAllocAsync` dropping from 8,453 calls / 16.802 ms to 7,757 calls /
  12.574 ms and `cuMemFreeAsync` from 6,048 calls / 13.801 ms to 5,352 calls /
  11.096 ms. HTTP `decode64` stays effectively flat at 11.48 tok/s, so the
  main target remains NCCL plus local expert GEMV.
- Added DSv4 GPU compressor projection scratch reuse for `kv_raw` and
  `score_raw`, with trace artifacts under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-compressor-projection-scratch/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-compressor-projection-scratch/)
  and
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-compressor-projection-scratch/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-compressor-projection-scratch/).
  The real-output checks still pass (`decode64` normal text, arithmetic `410`)
  and alloc/free calls fall again (`cuMemAllocAsync` 7,757 -> 6,765,
  `cuMemFreeAsync` 5,352 -> 4,360), but HTTP `decode64` remains flat at
  11.47 tok/s and the single nsys wave is not a wall-time win because D2H/NCCL
  timing dominates this capture.
- Added DSv4 incremental attention scratch Nsight and HTTP artifacts under
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-attention-scratch/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/bench-attention-scratch/)
  and
  [`docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-attention-scratch/`](docs/trace-artifacts/2026-05-15-dsv4-deepep/nsys-single-decode-token-attention-scratch/).
  The real 8xH20 run returns normal multi-token output and arithmetic `410`;
  the isolated single-token decode wave is 97.042 ms after B=1 attention
  scratch cuts decode-window free calls from 4,360 to 3,048 without retaining
  prompt-sized prefill buffers. The trace directly answers the current
  bottleneck question: sampler is not in the top stack; NCCL SendRecv/AllReduce,
  D2H route-count synchronization, launch/runtime overhead, local expert
  FP8/FP4 GEMV, and attention/MHC kernels dominate.

### CUDA

- Added a default DSv4 B=1 padded BF16 combine reduce-scatter path behind
  `ARLE_DSV4_COMBINE_REDUCE_SCATTER` (default `1`). Expert ranks now sum padded
  route outputs into one row per origin peer and call NCCL `ReduceScatter`
  directly into the owner-rank output hidden row, with `0` preserving the prior
  grouped SendRecv combine. Real 8xH20 DSv4 validation against
  `/root/DeepSeek-V4-Flash` keeps normal Chinese/English streaming output and
  exact arithmetic `410`; `decode64` measures 12.05 post-first tok/s. The
  matching single-token nsys wave moves from 97.071 ms to 94.923 ms, replacing
  the old 23.163 ms SendRecv combine bucket with a 20.443 ms ReduceScatter
  bucket plus 3.259 ms residual SendRecv. This is a modest communication-shape
  cleanup; local expert grouped GEMM/DeepGEMM, DeepEP overlap, launch reduction,
  scratch reuse, and D2H readback removal remain the main performance targets.
- Reused per-layer DSv4 DeepEP dispatch scratch for route setup, rank count
  exchange buffers, packed send hidden rows/metadata, and local expert
  count/offset/cursor buffers. On the 8xH20 default path, trace-off math smoke
  reached 7.7-7.8 tok/s for 12 generated tokens, traced
  `ffn_deepep_dispatch_combine` p50 dropped to 1.552 ms, and the profiled
  `cuMemAllocAsync`/`cuMemFreeAsync` call count fell from 136,825 to 111,531 in
  the 8-token Nsight window. Remaining bottlenecks are still stream sync,
  return-side NCCL send/recv, and local expert GEMV/GEMM.
- Reused DSv4 DeepEP send-route token/slot buffers across decode steps and
  removed the unused `expert_token` output from `dsv4_pack_received_experts`.
  The 8xH20 trace-off math/writing smoke remained normal at 7.94-8.09
  completion tok/s, while the single-token nsys window reduced decode-only
  `cuMemAllocAsync` calls from 11,980 to 11,097 and `cuMemFreeAsync` calls from
  11,988 to 11,105. Remaining allocator pressure now sits in recv/local route
  buffers plus combine scratch and still needs a broader lifetime/graph pass.
- Reused DSv4 DeepEP B=1 decode recv/local route scratch for received hidden
  rows and metadata, local expert packed rows/weights/route slots, and
  route-output rows. Prefill preallocates only a small `ep_world * topk` decode
  capacity so long prompts do not retain prompt-sized route buffers. The real
  8xH20 DSv4 smoke stayed correct at 8.24-8.79 completion tok/s, and the
  single-token nsys window improved from 191.152 ms to 148.253 ms while
  reducing decode-only `cuMemAllocAsync`/`cuMemFreeAsync` calls to
  9,480/9,488 and `cuMemsetD8Async` calls to 10,554.
- Reused the DSv4 B=1 decode MoE route-logits buffer and preallocated its
  one-token scratch during prefill. This is an allocator-count cleanup rather
  than a confirmed wall-time win: the single-token nsys window reduced
  decode-only `cuMemAllocAsync`/`cuMemFreeAsync` calls again to 9,136/9,144 and
  `cuMemsetD8Async` calls to 10,210, while the captured wall time was noisy
  at 162.062 ms versus the prior 148.253 ms.
- Reran the post-scratch DSv4 single-token Nsight capture on 2026-05-15. The
  fresh one-token decode wave measured 158.439 ms wall and confirms the
  remaining cost center is not sampler or KV-cache lookup: runtime allocation,
  launch, memset, D2H routing readbacks, NCCL exchange/reduction, and per-expert
  GEMV still dominate before attention.
- Reused per-layer DSv4 shared expert scratch during DeepEP decode and added an
  in-place BF16 add kernel for accumulating shared expert output into the routed
  MoE output. Real 8xH20 smoke stayed correct at 9.07-9.50 completion tok/s,
  while the single-token nsys wave improved from 158.439 ms to 140.111 ms and
  decode-only `cuMemAllocAsync`/`cuMemFreeAsync` calls fell from 9,136/9,144 to
  7,416/7,424. The same step restored the CUDA `argmax_batch_readback_into`
  re-export required by Qwen3.5 CUDA builds. The scratch is gated to B=1 decode
  so long prefill does not retain prompt-sized shared expert buffers.
- Optimized the gated DSv4 grouped expert prototype behind
  `ARLE_DSV4_GROUPED_EXPERTS=1` by caching per-layer local expert weight
  pointer arrays and launching indexed active experts instead of rebuilding
  active pointer tables every step. The route remains opt-in: 8xH20 trace-off
  smoke improved grouped math latency to 2.37-2.40 s and short writing latency
  to 2.69 s, but traced `ffn_deepep_local_experts` p50 is still 1.196 ms versus
  roughly 0.46 ms on the default scratch-reuse path. The harness is ready for
  the next replacement with real grouped GEMM/DeepGEMM.
- Added a gated DSv4 grouped gate/up pair GEMV launch for the same
  `ARLE_DSV4_GROUPED_EXPERTS=1` harness. The FP8/FP4 pair kernels compute
  `w1` and `w3` in one grouped launch when format, shape, and block-scale
  layout match, otherwise the path falls back to separate grouped GEMV
  launches. 8xH20 nsys with `ARLE_DSV4_MOE_BACKEND=deepep` confirms
  `dsv4_fp4_grouped_gemv_pair_batch_kernel` runs in decode, but the grouped
  harness remains default-off: the decode window is still dominated by NCCL
  send/recv plus allocation/free and launch churn, not by the missing gate/up
  fusion alone.
- Added a gated DSv4 MoE combine exchange experiment via
  `ARLE_DSV4_COMBINE_DTYPE=fp8`. The path quantizes return-route BF16 rows to
  FP8 E4M3 with per-row FP32 scales, exchanges the FP8 payload through NCCL
  `Uint8` send/recv plus scale exchange, and dequantizes back to BF16 before
  the existing route-slot combine kernel. It is validated on 8xH20 but remains
  opt-in because the measured 1,039-token prefill trace is not faster than the
  BF16 combine default.
- Reused per-layer DSv4 HyperConnection/MHC temporary buffers in the
  incremental attention and FFN paths. The 8xH20 trace-off smoke set improved
  from roughly 5.5/5.6/6.0 tok/s to 6.3/6.2/7.3 tok/s for two math cases and
  one short writing case, while traced decode `attn_mhc` and `ffn_mhc` p50
  dropped to 0.088 ms and 0.085 ms respectively.
- Reused DSv4 incremental attention scratch for prepared Q/K, local attention
  output, and the `wo_a` latent projection, gated to B=1 decode so prefill does
  not retain prompt-sized buffers. The real 8xH20 HTTP smoke remains correct
  (`decode64` normal text, writing normal Chinese, arithmetic `410`), while the
  paired single-token Nsight capture reduces warmed decode `cuMemFreeAsync`
  calls from 4,360 to 3,048.
- Added a default DSv4 local expert segment-input path for DeepEP decode. When
  `w1` and `w3` are DSv4 block-scaled FP8/FP4 matrices, the per-expert fallback
  now runs their GEMV directly from the packed `expert_hidden` segment and
  skips the old D2D copy into `scratch.input`; unsupported formats still use
  the original copy fallback. Real 8xH20 nsys against `/root/DeepSeek-V4-Flash`
  kept the `霓虹` streaming output, reduced decode-only `cuMemcpyDtoDAsync_v2`
  from 871 calls / 1.795 ms per rank range to 613 calls / 1.240 ms, and moved
  the isolated single-token wave from 146.448 ms to 145.104 ms. The trace
  confirms this is a small cleanup: allocator/runtime churn, D2H route
  readback, NCCL SendRecv/AllReduce, and per-expert FP8/FP4 GEMV remain the
  dominant costs.
- Reused per-layer DSv4 incremental hidden scratch for attention/FFN
  HyperConnection pre-projection and RMSNorm temporaries. Real 8xH20 nsys
  against `/root/DeepSeek-V4-Flash` kept the streaming `霓虹` output and moved
  the isolated decode wave from 145.104 ms to 135.390 ms. Decode-only
  `cuMemAllocAsync`, `cuMemFreeAsync`, and `cuMemsetD8Async` calls each dropped
  by 1,376, matching four one-token temporary buffers across 43 layers and 8
  ranks. The remaining ranked costs are launch/runtime overhead, D2H route
  readback, NCCL SendRecv/AllReduce, and local expert FP8/FP4 GEMV.
- Removed the default DSv4 DeepEP AllGather route's redundant 32-byte
  `send_rank_counts` host readback. The AllGather count matrix is now
  collected before route packing and reused to derive both send and receive
  counts; the `ARLE_DSV4_COUNT_EXCHANGE=sendrecv` fallback keeps the previous
  readback. Real 8xH20 nsys kept the `霓虹` output, moved the single-token
  decode wave from 135.390 ms to 129.768 ms, and reduced decode-only D2H calls
  from 887 to 543. The remaining D2H cost is the 256-byte all-rank count matrix
  readback, ahead of deeper device-side count-prefix or countless dispatch
  work.
- Added the default DSv4 B=1 padded dispatch fast path for DeepEP decode. When
  the count exchange mode is the default AllGather route, decode now uses fixed
  `ep_world * topk` route slots, initializes unused slots as invalid, skips the
  unused send-rank zero/count kernel, and avoids the count AllGather plus its
  256-byte all-rank D2H readback. Set `ARLE_DSV4_PADDED_DISPATCH=0` to force
  exact-count dispatch. Real 8xH20 nsys kept the `霓彩` streaming output, moved
  the single-token decode wave from 129.768 ms to 123.955 ms, removed
  `ncclDevKernel_AllGather` from the decode window, and reduced decode-only D2H
  calls from 543 to 344. The remaining slow stack is NCCL SendRecv/AllReduce,
  launch/runtime and allocator/memset/free churn, local-count D2H, and local
  expert FP8/FP4 GEMV.
- Optimized the B=1 padded return-side combine exchange by summing valid padded
  route outputs into one BF16 row per origin peer on the expert rank before the
  return send/recv. This keeps the same `霓彩` streaming output, reduces
  returned combine rows by 8x, moves the real 8xH20 single-token decode wave
  from 123.955 ms to 112.133 ms, and drops `ncclDevKernel_SendRecv` time from
  25.211 ms to 23.329 ms per rank range. The local expert FP8/FP4 GEMV timings
  are unchanged, so true grouped GEMM/DeepGEMM remains the next compute target.
- Added and gated the default-path DSv4 single-expert `w1`/`w3` pair GEMV
  experiment behind `ARLE_DSV4_PAIR_EXPERT_GEMV=1`. The real 8xH20 trace kept
  the `霓彩` output and proved the new `dsv4_fp4_gemv_pair_batch_kernel` runs,
  but it regressed the local expert work on the current B=1 decode shape
  (`23.207 ms` per rank range for the pair kernel, 127.412 ms decode wave), so
  the shipped default remains the split GEMV path while the next compute target
  stays true grouped GEMM/DeepGEMM.
- Added and gated a DSv4 route-wise grouped expert experiment behind
  `ARLE_DSV4_ROUTE_GROUPED_EXPERTS=1`. It runs local experts directly from
  padded received route slots and removes the local-count D2H readback from the
  top decode runtime list, but the real 8xH20 nsys trace regressed to a
  145.669 ms single-token wave because `dsv4_fp4_route_gemv_batch_kernel`
  costs 35.895 ms per rank range. The path remains default-off and documents
  why the next compute step needs DeepGEMM-style grouped GEMM rather than
  route-wise GEMV.
- Added a clean 8xH20 decode-only HTTP comparison for the gated
  `ARLE_DSV4_PAIR_EXPERT_GEMV=1` path. The default split expert GEMV path
  reaches 11.79 post-first tok/s on `decode64`, while pair GEMV reaches
  7.70 tok/s; both return normal sequence text and the arithmetic check returns
  `410`. This keeps pair GEMV default-off and confirms the next compute target
  is real grouped GEMM/DeepGEMM rather than single-expert gate/up fusion.
- Added `HiddenStates::uninit` for CUDA call sites that immediately overwrite
  every element and switched DSv4 decode temporaries plus generic GEMM/add/SwiGLU
  outputs to use it where safe. Real 8xH20 DSv4 HTTP smoke remains correct
  (`decode64` reaches 11.99 post-first tok/s and the arithmetic check returns
  `410`), and single-token nsys shows `cuMemsetD8Async` dropping from 8,789
  calls / 11.855 ms per rank range to 2,957 calls / 4.180 ms. The isolated
  decode wave moves from 125.497 ms to 112.724 ms; NCCL exchange, launch
  overhead, async allocation/free, and local expert FP8/FP4 GEMV remain the top
  targets.
- Added the default DSv4 B=1 fused dispatch payload path. Padded DeepEP decode
  now appends the route metadata as raw BF16 words behind each hidden row and
  exchanges hidden+metadata through one BF16 grouped send/recv instead of
  separate BF16 hidden and I32 metadata exchanges. Set
  `ARLE_DSV4_FUSED_DISPATCH_PAYLOAD=0` to force the old split exchange. Real
  8xH20 nsys keeps the output correct, reduces SendRecv launches from 1,032 to
  688, and records the isolated decode wave at 118.985 ms in the latest capture;
  NCCL SendRecv/AllReduce, launch/runtime overhead, allocator churn, D2H, and
  local expert FP8/FP4 GEMV remain the next targets.
- Optimized the gated route-wise grouped expert experiment by pairing its
  route-local `w1` and `w3` GEMV launches for matching DSv4 block-scaled FP8 or
  FP4 weights, falling back to split route GEMV when format or shape differs.
  The real 8xH20 nsys run lowers the prior route-grouped regression from
  145.669 ms to 117.894 ms, but it remains default-off because single-token
  decode is still dominated by NCCL SendRecv, route GEMV work, launch overhead,
  and async allocation/free. The main target remains true grouped
  GEMM/DeepGEMM with DeepEP overlap.
- **🎉 W4-hybrid prefill graph capture closes 4k/c=4 gap — Tier 1 STRONG
  PROCEED** (`a56b7a9`/`c44788f` 2026-05-10). Path B.2 bucketed prefill
  graph allocation key reduces capture key churn from 388 unique → **7
  unique** (98% reduction) with **98.5% LRU dominant key reuse rate**.
  Engine-side TTFT p50 **2000ms → 150ms = -92.5%** improvement on
  4k/c=4 prefill-dominant workload (server-side ground truth via
  `/v1/stats engine_ttft_us`; client-side guidellm 0.6.0 TTFT
  measurement separately broken per `e8d82b0` — bench tool bug, not
  substrate). Throughput **+632%** in matched-control 60s window
  (53 → 388 requests). Codex's "second-order bucketing" insight
  (captured scalar launch parameters use bucket capacity, not exact
  dim from first capture) was load-bearing for the win and added to
  skill v1.7.0 anti-pattern catalog. Followup: n=3 σ-tight re-bench +
  guidellm streaming fix. Evidence:
  [`docs/experience/wins/2026-05-10-bench-40-pathB2-tier1-strong-proceed.md`](docs/experience/wins/2026-05-10-bench-40-pathB2-tier1-strong-proceed.md).
- W4-hybrid Qwen3 paged-prefill **CUDA Graph capture** lands as opt-in
  via `INFER_PREFILL_GRAPH=1` + `INFER_HYBRID_W4A8_PREFILL=1` (`35fc3cf`).
  Phase 1 functional gate: prefill-lifetime `MarlinPrefillScratch`
  lifecycle + multi-key 8-d graph cache (token / page layout / start_pos)
  + W4 graphsafe weight gating for dense BF16, W4A16 Marlin, W4A8 Marlin,
  and W4-hybrid. Default behavior unchanged when env vars unset.
  Throughput license deferred: scout bench A vs B (graph OFF baseline
  TTFT p50 1628.9 ms vs graph ON 1627.8 ms = Δ -0.07%) detected
  capture-key churn — Path A multi-key direction KILLED, Path B
  device-memory `start_pos` re-licensed P0 (`e462c53`). Evidence:
  [`docs/experience/wins/2026-05-10-bench-p24-w4a8-prefill-graph-hoist.md`](docs/experience/wins/2026-05-10-bench-p24-w4a8-prefill-graph-hoist.md),
  [`docs/experience/errors/2026-05-10-37-throughput-bench-killed-pathA-multikey-churn.md`](docs/experience/errors/2026-05-10-37-throughput-bench-killed-pathA-multikey-churn.md).

### Long-context (cross-backend)

- **RoPE scaling support** (YARN / Linear / NtkAware) wired through
  `Qwen3Config::rope_scaling` and `Qwen35Config::rope_scaling` (Phase
  1+2 closed via 7 atomic commits + 51 unit tests). Helpers
  `compute_scaled_inv_freq` and `compute_attention_factor` ship in both
  spec crates. CUDA backend integration via
  `weight_loader::precompute_rope_with_scaling` (qwen3 path) +
  `precompute_rope_with_qwen35_scaling` thin shim. Vanilla path
  (`rope_scaling = None`) is bit-equivalent to the legacy
  `precompute_rope` formula (verified by
  `vanilla_inv_freq_matches_legacy_formula` test). Long-ctx bench
  validation (Qwen3-4B 64k YARN×2 / 128k YARN×4 + FP8 KV) deferred to
  Phase 3; CUDA-side viable on RTX 4070 Ti SUPER 16 GB per
  [`docs/plans/2026-05-10-rope-yarn-phase3-cuda-bench-plan.md`](docs/plans/2026-05-10-rope-yarn-phase3-cuda-bench-plan.md).
  Apply to a model dir via [`scripts/setup_qwen3_yarn_config.py`](scripts/setup_qwen3_yarn_config.py).
  Consolidation:
  [`docs/experience/wins/2026-05-10-m-rope-yarn-scaling-phase1-phase2-landed.md`](docs/experience/wins/2026-05-10-m-rope-yarn-scaling-phase1-phase2-landed.md).

### Structured-output (xgrammar)

- `crates/xgrammar-sys` Rust safe wrapper over upstream
  `mlc-ai/xgrammar` v0.1.34 lands as Phase 1 FFI scaffold (codex's #26).
  Default build is a stub that compiles without native sources or
  network; `--features real` builds a C++ shim against a pinned
  upstream checkout via `cc`. Wrapper surface:
  `GrammarCompiler` / `CompiledGrammar` / `GrammarMatcher` /
  `bitmask_size` / per-step bitmask fill APIs. No HTTP, scheduler,
  sampler, or GPU sampling integration yet — that is follow-up
  tranche work. Plan:
  [`docs/plans/M_xgrammar-ffi-scaffold.md`](docs/plans/M_xgrammar-ffi-scaffold.md).

### Metal

- Qwen3.5-0.8B MLX 4bit single-request step-driver reaches 305.5 tok/s mean
  / 304.7 p50 on M4 Pro 20c for `1024/256`. The matched GGUF Q4_K_M
  exact default remains 202.1 tok/s direct for correctness, while the
  opt-in native-q4 load path reaches 236.7 tok/s direct / 239.8 tok/s
  step-driver, so current status surfaces no longer present the historical
  211.7 tok/s GGUF-only profile as the Metal SOTA headline. Evidence:
  [`docs/experience/wins/2026-04-28-bench-metal-qwen35-0p8b-mlx4bit-qknorm-default.md`](docs/experience/wins/2026-04-28-bench-metal-qwen35-0p8b-mlx4bit-qknorm-default.md).
  Native-q4 GGUF evidence:
  [`docs/experience/wins/2026-04-28-bench-metal-qwen35-0p8b-gguf-native-q4.md`](docs/experience/wins/2026-04-28-bench-metal-qwen35-0p8b-gguf-native-q4.md).

## [0.1.4] — 2026-04-28

Correctness fix release for Metal chat generation. Multimodal Qwen3.5 /
Qwen3.6 MoE configs (HF `Qwen3_5MoeForConditionalGeneration` family)
declare `eos_token_id` as an array at the root of `config.json` while
nesting a single base-LM EOS inside `text_config`. Our Metal config
loader took only the first element of the array and let `text_config`'s
scalar shadow the root, then the C++ generate path passed only that one
id into the stop check — so on `mlx-community/Qwen3.6-35B-A3B-4bit` the
model walked past `<|im_end|>` after the first reply and hallucinated
fresh user/assistant turns, with `skip_special_tokens=true` decoding
leaving "user"/"assistant" plain-text role names visible.

### Metal

- `infer/src/backend/metal/config.rs` — replaced ad-hoc `get_eos` +
  `load_stop_token_ids` with `resolve_stop_token_ids`, a generic
  HuggingFace-precedence resolver: `generation_config.json` (HF
  inference-time authority) → `config.json` root → `text_config`,
  scalar-or-array normalized to `Vec<u32>`, dedup preserves first-seen
  order, fall back to 151645. `MetalModelConfig.eos_token_id` is now
  the first id from the resolved array; `stop_token_ids` is the
  authoritative list.
- `infer/src/backend/metal/{qwen35.rs, generate.rs}` — C++ generate
  paths now extend `stop_ids` with the full `config.stop_token_ids`,
  not just `config.eos_token_id`. Sort + dedup after merge.
- `infer/src/backend/metal/request_state.rs` — scheduler/HTTP path
  also folds `config.stop_token_ids` into `ResumableRequestState`,
  guarded by `params.ignore_eos` so benchmarks that explicitly want
  to generate past EOS still can.
- Three new HF-precedence unit tests in `metal::config::tests`.

Verified on Apple M4 Max with `mlx-community/Qwen3.6-35B-A3B-4bit`:
multi-turn REPL produces clean, on-topic, EOS-terminated replies; tool
execution works on the third turn.

### CUDA · breaking — SM env-var policy

- `INFER_CUDA_SM` and `CUDA_SM` are removed. The new build-time env
  var is `TORCH_CUDA_ARCH_LIST` (alias `CMAKE_CUDA_ARCHITECTURES`),
  matching the PyTorch / vLLM convention. Migration: replace
  `INFER_CUDA_SM=89` with `TORCH_CUDA_ARCH_LIST=8.9`.
- T1 (sm_80/86/89/90) is the default fat-binary set when no env is
  set. T2 (sm_100/120) is opt-in. T3 (sm < 80) is rejected at build
  time with a hint. Unknown SMs hard-fail rather than silently
  falling back.
- Plan + tier policy live in `docs/plans/sm-coverage.md`; four
  per-SM bench stubs are tracked under `docs/experience/wins/` as
  `pending-remote` and will be filled in alongside the multi-cubin
  AOT dispatch (Phase B/C).

## [0.1.3] — 2026-04-27

Packaging fix release: macOS bottles and release tarballs were missing
`mlx.metallib`, so every `metal_serve` model load on a fresh install
hit "Failed to load the default metallib". v0.1.3 ships the metallib
alongside the binaries on every distribution channel. Also folds in
the scheduler `cuda/runtime.rs` + `cuda/core.rs` splits and the Phase 2
KV swap revert. Pre-built artifacts on the
[GitHub Release page](https://github.com/cklxx/arle/releases/tag/v0.1.3)
and on GHCR (`ghcr.io/cklxx/arle:0.1.3`, `:0.1`, `:latest`).

### Tooling / packaging

- `scripts/package_macos_metal_artifact.sh` bundles `mlx.metallib`
  into the macOS tarball (auto-discovers from the cargo build dir;
  override via `ARLE_MLX_METALLIB`). MLX's `load_default_library`
  searches binary-colocated paths, so without this the brew bottle
  and curl-installed binaries fell straight through to the
  compile-time `METAL_PATH` (a cmake build dir absent in
  distribution).
- `scripts/install.sh` installs `mlx.metallib` next to the binaries
  on macOS-arm64.
- `crates/mlx-sys/build.rs` stages `mlx.metallib` into
  `target/<profile>/` after the cmake build, so locally-built
  `metal_serve` works without colocating the metallib by hand.
- `cklxx/homebrew-tap` formula picks up `mlx.metallib` via
  `bin.install "mlx.metallib" if File.exist?("mlx.metallib")` — no
  effect on installs of older tarballs that don't ship the file.

### Scheduler

- `infer/src/scheduler/cuda/runtime.rs` split into
  `runtime/{admission,fetch,helpers,scheduler_loop,swap,tests}.rs` to
  contain the per-iteration scheduler loop. Behavior unchanged.
- `infer/src/scheduler/cuda/core.rs` extracted into
  `helpers/state_types/emit_worker/construction/warmup` siblings —
  same flat-module-no-`mod.rs` shape used elsewhere.
- Phase 2 KV swap path deleted entirely (vLLM V1 / SGLang precedent
  — tier demote/promote already covers the workload that motivated
  swap; the dual path was not paying its complexity cost).
- Host pool for the now-removed swap path right-sized while the
  revert was in flight; entry recorded under `errors/`.

### Metal

- Qwen3.5 prefill: Rust fallback path now materializes its hidden
  state correctly when the C++ prefill bails out (previously
  returned a stale handle).
- Qwen3.6 MoE wired through the C++ prefill alongside Qwen3.5 so the
  large model uses the same fused step path.
- `mlx-sys` FFI guard shared across the `infer` and `autograd`
  crates so MLX process-global state has one Rust synchronization
  boundary regardless of which side enters MLX first.

### HTTP

- `infer/src/http_server.rs` split into
  `http_server/{types,handlers,router,tests}.rs`. Behavior unchanged.

### Refactors

- `infer/src/backend/metal/request_state.rs` extracted into helper
  modules + standalone tests.

## [0.1.2] — 2026-04-27

Engine consolidation, Metal GGUF perf, train binary migration, and a
v2 landing site. Pre-built artifacts on the
[GitHub Release page](https://github.com/cklxx/arle/releases/tag/v0.1.2)
and on GHCR (`ghcr.io/cklxx/arle:0.1.2`, `:0.1`, `:latest`).

### Runtime

- `ModelInferenceEngine` and per-model aliases removed; everything now
  routes through `LoadedInferenceEngine`. The scheduler grew a unified
  `Cuda(RequestHandle)` variant so CUDA drives the same path as Metal.
- E2E tests migrated to the `LoadedInferenceEngine` scheduler path,
  matching the runtime that production serves use.
- Unused `PreemptionMode` enum and config field dropped from the
  scheduler — the unification project chose tier demote/promote over
  the separate preemption type.
- `kv-tier` gained a `WholeKv` variant on `ReadmissionPlan` plus a
  `PlanKind` discriminator; the Plan path is documented as
  reserved-for-distributed scaffolding (no runtime change).
- Metal Qwen3.5-0.8B GGUF Q4_K_M decode now crosses 200 tok/s on M4 Pro
  (211.7 tok/s for 512 prompt / 1024 decode) after Q5_K/Q8_0 affine
  repack and Q6/group16 qmv tile tuning.
- Qwen3.6-35B-A3B Metal rechecked locally with a short load/execute
  diagnostic. This is not DFlash acceptance evidence; future DFlash
  optimization claims for Qwen3.6 should use long-context /
  ultra-long-sequence workloads only. Evidence:
  [`docs/experience/wins/2026-04-27-bench-metal-qwen36-a3b-dflash-quick-check.md`](docs/experience/wins/2026-04-27-bench-metal-qwen36-a3b-dflash-quick-check.md).

### HTTP

- `session_id` now threads through `CompletionRequest`; the hardcoded
  `None` in the OpenAI-compatible handlers is gone.

### Metal

- Qwen3.5 GGUF decode path tuned (perf); Q4/Q6 weights are now repacked
  into MLX affine layout at load time so decode hits the fast kernel.
- `qwen35` GGUF projections kept packed end-to-end (avoids a redundant
  unpack/repack round-trip).
- Metal GGUF decode status doc refreshed under
  `infer/src/backend/metal/AGENTS.md`.

### Train

- `crates/train/src/bin/*` moved to `crates/train/src/commands/*` —
  `arle train …` and `arle data …` are the only entry points; the
  `[[bin]]` declarations in `crates/train/Cargo.toml` are gone.

### Docs & web

- Landing site v2: 3-pane docs shell + Quickstart page; homepage
  split-layout hero with dark footer and framed sections; logo
  redrawn as a two-story humanist 'a'; wordmark swapped for topology
  mark with v2 design tokens.
- Install cards on the landing site now lead with `brew install
  cklxx/tap/arle` and the `curl | sh` installer (Docker / source
  demoted to fallbacks).
- README, roadmap, support matrix, and maintainer docs now share the
  same Metal GGUF support and benchmark wording.
- Intra-doc links cleaned up for `rustdoc -D warnings`: links from
  `kv-tier` and elsewhere now resolve to public items only.
- `CONTRIBUTING.md` Getting Started leads with an end-user pointer to
  the README install section so contributors don't assume source
  build is the canonical path.

### Tooling

- Homebrew formula in `cklxx/homebrew-tap` simplified to macOS-arm64
  only. The `mislav/bump-homebrew-formula-action` only bumps one URL,
  so a multi-platform formula went stale on every release; Linux
  brew users for a CUDA-driven tool are vanishingly rare (Docker /
  install.sh / source cover that path).

## [0.1.1] — 2026-04-27

Install ergonomics + a batch of TileLang, KV-tier, and Metal/Qwen3.5
follow-ups. Pre-built artifacts on the
[GitHub Release page](https://github.com/cklxx/arle/releases/tag/v0.1.1)
and on GHCR (`ghcr.io/cklxx/arle:0.1.1`, `:0.1`, `:latest`).

### Install

- **Homebrew tap**: `brew install cklxx/tap/arle`
  ([cklxx/homebrew-tap](https://github.com/cklxx/homebrew-tap)). The
  `bump-homebrew` job in `release.yml` keeps the formula in lockstep
  with each `v*` tag.
- **One-line installer**: `curl -fsSL
  https://github.com/cklxx/arle/releases/latest/download/install.sh
  | sh`. Detects platform, SHA256-verifies the tarball, installs to
  `~/.local/bin` (override via `INSTALL_DIR`).
- New `docs/install.md` documents the full matrix, env-var overrides,
  and uninstall steps; `docs/release-checklist.md` §4a covers the
  `HOMEBREW_TAP_TOKEN` secret needed for tap automation.

### Runtime

- TileLang prefill HD128 path behind the `tilelang-attn` feature
  (Experimental); per-Qwen3 head-config specialization. Prefill HD256
  + decode HD256 + tc-decode AOT kernels added under the same flag.
  L4 floor verified; H100 verification runbook landed.
- macOS Metal link fix: `compiler-rt` now linked so mlx-sys resolves
  `__isPlatformVersionAtLeast`; release tarballs and Metal CI build
  cleanly on hosted runners.
- Metal Qwen3.5 GGUF Q4 path: aligned `gdr` and prefill matmul; gated
  GGUF DFlash; preserved greedy C++ generate. Long bench coverage
  recorded under `docs/experience/wins/`.
- `release.yml` now installs `flashinfer-python` (in lockstep with
  Dockerfile) so Linux CUDA tarballs build without the
  `flashinfer/pos_enc.cuh` not-found regression.
- KV-tier P3 follow-ups: typed `FailureClass` in events with explicit
  emit policy; `OrchestratorEvent` unified into `CoordinatorEvent`;
  `emit_observability` switched to `try_send` while `Store*` paths
  retain required delivery.
- `cargo build` no longer requires an explicit backend feature on
  macOS — `default = ["cuda"]` was dropped from the workspace,
  `agent-infer`, and `crates/cli` so a flag-less build no longer pulls
  cudarc on platforms without nvcc.
- Standalone `pretrain` / `train_sft` / `train_grpo` /
  `train_multi_turn` / `eval_lm` / `download_dataset` /
  `convert_dataset` binaries removed from `crates/train`;
  `arle train …` and `arle data …` are the single front door
  (sources still in `crates/train/src/bin/*.rs` as in-process
  dispatch modules under `autobins = false`).

### Docs & web

- `docs/support-matrix.md` §4b documents the multi-turn KV reuse /
  tiered-KV stability tiers (T0 GPU Supported, T1 host-pinned Beta,
  T2 NVMe Beta, T3 cluster-shared Experimental — NIXL stub-only).
- `docs/experience/wins/README.md` explains the `*pending-remote*.md`
  filename convention so external readers do not mistake stubs for
  published claims.
- `docs/troubleshooting.md` and `docs/comparison.md` added.
- Astro/Vite landing site at `web/` (deploys via `pages.yml` to
  <https://cklxx.github.io/arle/>); homepage IA refactor — runnable
  hero, dated bench, dropped chrome; switched to minimal-white
  aesthetic.

### CI

- `cuda-ci.yml` gated behind `workflow_dispatch` until the
  self-hosted CUDA runner returns.

## [0.1.0] — 2026-04-26

First tagged release. Pre-built artifacts live on the
[GitHub Release page](https://github.com/cklxx/arle/releases/tag/v0.1.0)
and on GHCR (`ghcr.io/cklxx/arle:0.1.0`, `:0.1`, `:latest`).

### 2026-04-26 — Open-source usability and `arle` front door cleanup

#### CLI / DX
- Added `arle serve`, a unified front door that launches the matching serving
  binary (`infer`, `metal_serve`, or `cpu_serve`) from the release artifact or
  PATH.
- Added `--no-tools` for the local agent runtime so one-shot and REPL prompts
  can explicitly disable built-in shell/python tool execution.
- Extended `arle --doctor --json` to schema version 3 with tool/sandbox
  diagnostics, including the detected sandbox backend.

#### Packaging
- Renamed release tarballs to `arle-<version>-<platform>.tar.gz`.
- macOS release artifacts now include both `arle` and `metal_serve`; Linux
  artifacts include `arle`, `infer`, and `bench_serving`.
- The Docker image now uses `ghcr.io/cklxx/arle` and enters through `arle`
  instead of exposing only `infer`.

#### Docs and examples
- Added copyable examples under `examples/` for curl, stdlib Python,
  Docker Compose, Apple Silicon local serving, and the tiny train fixture.
- Updated README, Chinese README, support matrix, release checklist, and
  security guidance for the unified front door and tool safety controls.

### 2026-04-25 — Truth-surface cleanup

Documentation-only refactor that collapses `docs/` to a single source of
truth per [`docs/plans/2026-04-20-project-constitution-and-refactor-plan.md`](docs/plans/2026-04-20-project-constitution-and-refactor-plan.md)
§2. No code or behavior change.

Net effect: the documentation tree shrinks by ~330 markdown files. After
this commit series, `docs/index.md` lists every document that counts as a
source of truth; anything not on that index is not.

Retired surfaces:

- `docs/archives/` and `docs/areas/` removed; the surviving "Workspace
  governance rules" (PR discipline + crate-admission criteria) inlined
  into `docs/architecture.md`.
- `docs/plans/` collapsed from 58 entries to 10 — the 8 active plans
  listed in `docs/index.md` plus the canonical bench-parameter
  (`guidellm-integration.md`) and kernel-crate-extraction blueprints.
  Six tiered-kv `*-remote-acceptance.md` checklists folded into the
  `docs/projects/tiered-kv-cache.md` milestone ledger as one-line
  "completed YYYY-MM-DD; see wins/<file>" pointers.
- `docs/projects/` collapsed from 8 to 5; the dropped three
  (`kv-quantization-long-context`, `qwen35-batched-decode`,
  `xma-future-research`) are either superseded by `docs/resources/
  kv-cache-quantization.md` or off-roadmap.
- `docs/research/` collapsed from 6 to 1 (only the
  `mni-ml-framework-notes.md` reference referenced from the agent-RL
  project survives).
- `docs/reviews/` collapsed from 4 to 2 (cuda-kernel-six-principles +
  metal-ecosystem-route-correction; both still cited from active docs).
- `infer/docs/` parallel tree retired; `profiling-guide.md` consolidated
  into `docs/resources/`.
- 45 `pending-remote` / `pending-local-rerun` stub wins/ entries that
  never converted to real measurements deleted.
- 44 pre-2026-04-15 micro-cleanup wins/ + early errors/ retired (history
  preserved in git log + this CHANGELOG).
- 150 superseded bench wins/ entries retired (intra-step iterations of
  CUDA c1–c16 closure, Qwen3.5 paged-prefill landing, Qwen3.6 MoE
  DFlash bring-up, and per-step scheduler / KV-tier redesigns) — kept
  the milestone summaries and the latest-per-topic entry only.

Survival criteria for future cleanups: `docs/index.md` lists every
source-of-truth file. Adding a second index, a parallel doc tree, or a
sibling status matrix is a regression and must be rejected at PR time.

### 2026-04-15 — Workspace consolidation + CUDA layer hygiene

Coordinated refactor round finishing Route-A and pre-staging the internal seams for a future CUDA kernel crate extraction. No user-facing behavior changes — all work is structural.

#### Workspace consolidation (Route-A)
- Folded four shell crates back into `infer`. The workspace is now a flat `infer` crate again, with submodules as the only seam between backends.
- Collapsed the `agent_engine` duplicate façade into `server_engine`. There is now one engine trait (`InferenceEngine`) and one loaded enum (`LoadedInferenceEngine`) covering CUDA, Metal, and CPU backends through a single dispatch path.
- Renamed types for unambiguous semantics: `ServerEngine` → `InferenceEngine`, `CompleteRequest` → `CompletionRequest`, `Usage` → `TokenUsage`, etc. Imports across the HTTP server, agent CLI, and scheduler updated in lock-step.
- `LoadedInferenceEngine` is now the single entry point for both the HTTP server and the agent CLI — no more backend-specific façades above the engine trait.

#### Chat naming disambiguation
- Renamed the OpenAI wire-format chat types to `OpenAi*` (`OpenAiChatMessage`, `OpenAiToolCall`, `OpenAiFunctionCall`, …) so they no longer collide with the internal protocol names re-exported from `infer_chat::protocol`. The HTTP layer now consistently uses `OpenAi*` on the wire and converts to the internal types before handing work to the engine.

#### CUDA layer hygiene
- Split `backend/cuda/ffi.rs` (1500 lines) into ten domain submodules (attention, gemm, kv, quant, graph, stream, etc.) with a clean re-export surface — no behavioral change, just navigability.
- Introduced `backend::cuda::prelude` as the proto-API contract for downstream modules. Only genuinely universal CUDA handles land in the prelude; per-discipline writeup, `TokenKVPool` explicitly stays out (see `prelude.rs` doc comment for the rule).
- Deleted four dead Triton kernels that had no live callers and removed the vestigial `replaced_cuda_files` bookkeeping directory left over from an earlier migration.
- `build.rs` Triton `cargo:rerun-if-changed` list is now auto-derived from a directory walk instead of a hand-maintained constant, so new Triton kernels don't silently skip rebuilds.

#### Kernel crate extraction (option B landed same day)
- Followed the hygiene round by executing the option B extraction locked in [`docs/plans/cuda-kernel-crate-extraction.md`](docs/plans/cuda-kernel-crate-extraction.md) — the `backend::cuda::prelude` seam was the staging point, and commits `a4e12f5` → `0ab2cd1` → `081cf32` landed the one-day mechanical refactor. `backend/cuda/` now contains only `bootstrap.rs`; all kernel sources, FFI, paged KV, FlashInfer wrappers, graph pool, tensor primitives, KV quant, and TurboQuant live under [`crates/cuda-kernels/`](crates/cuda-kernels/) with a one-way `infer → cuda-kernels` dependency. CUDA kernel C++ sources moved from `infer/csrc/cuda/` to `crates/cuda-kernels/csrc/{attention,gemm,kv,misc,quant}/`.
- The `mlx-sys` bridge was promoted from `infer/mlx-sys/` to [`crates/mlx-sys/`](crates/mlx-sys/) as part of the same Route-A flattening so both native layers (CUDA, Metal) sit peer-level under `crates/`.

### Governance
- Added a formal stability policy
- Added a support matrix document
- Added a compatibility and deprecation policy
- Added performance and correctness gate guidance
- Documented the CPU backend as a development-oriented smoke and validation path

### Added
- Radix-tree prefix cache for cross-request KV reuse
- Paged KV block manager with copy-on-write sharing
- Token-level KV pool (page_size=1, FlashInfer-compatible)
- GPU-CPU KV offload for contexts beyond VRAM capacity
- CUDA Graph batched decode (per batch size 1-32)
- Continuous batching scheduler with decode-priority and chunked prefill
- OpenAI-compatible API (`/v1/completions`, `/v1/chat/completions`, SSE streaming)
- Prometheus metrics (`/metrics`) and stats endpoint (`/v1/stats`)
- Qwen3 (0.5B-72B) and Qwen3.5-4B model support
- Built-in agent runtime with tool calling (shell, python)
- Benchmark suite (throughput, agent, multi-request)
- macOS Metal backend (experimental)
- Development-oriented CPU backend path for non-GPU request validation

### Performance
- TTFT 4.6x faster than SGLang v0.5.9 (8.6ms vs 39.3ms at C=1)
- Throughput parity: 0.99x at C=1, 0.92x at C=8
- 100% KV cache hit rate on multi-turn agent benchmarks
