Start model-agnostic EXL3 SM121 kernel lab

README.md

@@ -935,6 +935,8 @@ After CUDA compile, Python overlay edits (`overlay/exl3.py`, tests) are a cheap
 | `overlay/tp3/` | TP=3-only shape overlays (FlyCockpit MIT): head/vocab/shared-expert/A_log pads, EP loader, SM120 decode pad. Not on the TP=2 path |
 | `files/nfs-share.sh` / `files/nfs-server/` | `NFS_SHARE=1`: workers mount the head's HF cache over NFSv4 instead of holding a copy. Used by `start.sh` and `start-tp3.sh`. This kit has it on |
 | `start-tp4.sh` / `.env.tp4.example` | experimental 4-node TP=4 launch; knobs stay out of `.env` |
+| `kernel_lab/exl3/` | model-agnostic EXL3 K1-K8 oracle, SM121 tactic sweep, and Atlas receipt producer (development-only; not in the production image) |
+| `docs/exl3-sm121-kernel-lab.md` | kernel ABI, measurements, upstream sources, decisions, and next gate |
 | `files/chat_template.jinja` | GLM-5.3 MM template (`<|image|>` / `<|video|>`); checkpoint jinja is language-only |
 | `overlay/qwen3_dflash2.py` | DFlash2 draft (grouped conv + candidate selector) |
 | `overlay/dflash2_speculator.py` | DFlash2 selector walk (V2 speculator) |
@@ -964,6 +966,13 @@ After CUDA compile, Python overlay edits (`overlay/exl3.py`, tests) are a cheap
 Image-build runs `EXL3_SELFCHECK_GPU=0`. `./start.sh` runs the GPU self-check
 (`docker run --gpus all`) before shipping unless `SKIP_OVERLAY_VERIFY=1`.
 
+The kernel lab is development-only source: the production Dockerfile does not
+copy or install `kernel_lab/`, so the served image does not contain it. The lab
+accepts arbitrary `OPERATOR:K:N` shapes or a JSON workload file; GLM/Hy4 names
+are optional historical fixtures, not a supported-model list. See
+[`kernel_lab/README.md`](kernel_lab/README.md) for the portable correctness
+gate and isolated SM121 tuning command.
+
 ## Do not
 
 - Destroy HF weights, requantize, or `docker rm` HF caches. `REFRESH_WEIGHTS=1 ./download.sh` only if you intend to re-fetch

docs/exl3-sm121-kernel-lab.md

@@ -0,0 +1,182 @@
+# Atlas EXL3 SM121 Kernel Lab — engineering log
+
+Last updated: 2026-08-31
+
+## Status
+
+The first optimization stage is implemented but not yet timed on an isolated
+SM121 device:
+
+- reuse ExLlamaV3's packed ABI and direct CUDA kernels;
+- accept arbitrary model/operator shapes rather than a model registry;
+- validate MCG K1-K8 with independent materialized and streaming CPU oracles;
+- enumerate ExLlamaV3's built-in tile shapes and SM quotas;
+- emit native Atlas receipts for every exact measured tactic;
+- keep the production ExLlamaV3 pin unchanged while enabling an isolated
+  current-upstream comparison from this development-only source checkout.
+
+No speedup is claimed until the checked-in sweep runs on SM121.
+
+## Repository and architecture boundary
+
+CUDA/runtime work belongs in this serving repository. Atlas owns profiling,
+candidate decisions, receipt import, and promotion gates. This lab is the
+bridge: it produces `atlas.kernel-benchmark/v1` measurements that Atlas can
+rank without importing CUDA infrastructure.
+
+The lab is development-only source: the production Dockerfile neither copies
+nor installs `kernel_lab/`, and no production launch path imports it. Receipts
+come from a source checkout, never from a served image.
+
+The surrounding deployment remains GLM-specific, but the kernel-lab API is
+not. `model_id`, `operator`, and `phase` are receipt provenance strings.
+Compatibility is based on hardware/software identity, representation ABI,
+M/N/K, tensor-parallel world size, grouped state, and backend commit.
+
+The five unrelated untracked download/cutover scripts discovered in this
+checkout were preserved and are not part of the kernel-lab branch.
+
+## ABI decision
+
+The lab reuses ExLlamaV3's serialized MCG ABI without modification:
+
+| Field | Dtype | Shape |
+|---|---|---|
+| `trellis` | int16 | `[K/16, N/16, 16*bits]` |
+| `suh` | fp16 | `[K]` |
+| `svh` | fp16 | `[N]` |
+| `mcg` | int32 | scalar marker `0xCBAC1FED` / `-877912083` |
+
+The matrix convention is `A[M,K] @ W[K,N]`. Trellis tiles are 16×16 in
+ExLlamaV3 tensor-core lane order. The direct MCG adapter now accepts every
+integer upstream bitrate K1-K8. Fractional model-wide rates remain allocations
+across integer-rate tensors, not an invented fractional tensor format.
+
+Production remains pinned to ExLlamaV3
+`c5d9c657966ffeeaa9353f0cc899f18629da4a13`. An alternate immutable upstream
+commit can still be built and self-checked as an isolated source-checkout
+experiment, but this lab is development-only: the production Dockerfile carries
+no build stamp or copy for it. Measured receipts take the runtime source commit
+from `--runtime-commit` or from the Git identity of the clean checkout that
+produced them, and fail instead of emitting a receipt when no full SHA
+resolves. The `RUNTIME_SOURCE_COMMIT` and `EXLLAMAV3_COMMIT` environment
+values are declarations, not attestation: the first is honored only when it is
+exactly the loaded source, and the second must be supplied before the metadata
+module is imported because the pinned fallback describes only the production
+pin and cannot prove an arbitrary installed extension's identity. Unknown or
+mismatched backend/source identity is not measurement evidence.
+
+## Reference and direct paths
+
+`kernel_lab/exl3/reference.py` implements the authoritative bit order,
+tensor-core permutation, MCG procedural codebook, and 128-wide
+Hadamard/scaling transforms in NumPy.
+
+Two correctness paths are intentionally independent:
+
+1. decode and materialize the complete weight, then multiply;
+2. transform activations, decode one packed 16×16 tile, consume it immediately,
+   and discard it.
+
+The SM121 direct path calls ExLlamaV3 `exl3_gemm` with packed trellis tensors.
+Full reconstruction is used only for separately timed correctness and
+diagnostic baselines. It is outside the direct CUDA event range and receipts
+set `full_precision_materialized=false` for the measured execution path.
+
+## First SM121 tactic surface
+
+The pinned extension already exposes controls that production normally hides:
+
+- automatic dispatch/autotuning;
+- four compiled GEMM tile shapes, filtered by exact K/N compatibility;
+- an explicit number of participating SMs.
+
+The lab now measures automatic dispatch plus each compatible tile at quarter,
+half, and full-device SM quotas. This yields a hardware-derived dispatch table
+instead of rules tied to a model name.
+
+The current upstream ExLlamaV3 candidate at
+`0c49587a7c235e6303a6bbedc8b665272ad3a2ea` additionally contains a QTIP-style
+small-M FP16 GEMV and a fused mul1 int8-activation GEMV. Those are candidate
+experiments, not silently adopted production code.
+
+SM121 is consumer Blackwell. NVIDIA's current CuTe implementation directs
+SM120/121 kernels to warp-level `mma.sync`, not datacenter-Blackwell
+`tcgen05`. Therefore the first experiment stays on ExLlamaV3's CUDA
+`mma.sync` path. CuTe/CUTLASS remains useful for layout and medium-M scheduling
+once measurements demonstrate a gap.
+
+## Workload input
+
+The benchmark accepts any exact shape as `OPERATOR:K:N` or a JSON workload
+list. No code change is needed for a future model. Defaults are generic
+square/expand/contract shape classes. Early GLM-5.2 and Hy4 shapes remain only
+under explicit `--fixture` names for reproducibility.
+
+The default M sweep is `1,2,4,8,16,32,64,128`. The first performance gate
+should focus on M=1/2/4/8; wider M values locate the direct-to-reconstruction
+crossover.
+
+## Evidence so far
+
+The entries below are the author's recorded results for the dates given. They
+were not re-run at the development-only cutover head and are not exact-head
+qualification; the counts belong to their own runs and are not reconciled into
+a single figure.
+
+### Measured — portable host
+
+- Focused kernel-lab suite: 19 passed.
+- K1-K8 deterministic pack/unpack: passed.
+- K1-K8 streaming matmul versus materialized oracle: passed.
+- Worst maximum absolute error across that sweep: `1.08e-6`.
+- No dense tensor exists in the streaming oracle.
+
+### Previously measured — GB10 compatibility
+
+- NumPy K3 packing matched the official ExLlamaV3 CUDA packer byte-for-byte on
+  deterministic and randomized fixtures.
+- NumPy MCG decoding matched the CUDA decoder for all 65,536 states.
+
+### Not measured
+
+- Direct tactic latency on an isolated SM121.
+- Effective bandwidth/TFLOP/s and reconstruction overhead on target shapes.
+- Pinned 0.0.43 versus current-upstream small-M GEMV.
+- mul1 int8-activation crossover.
+
+The current development host exposes no NVIDIA CUDA device. Earlier available
+Sparks were occupied by an unrelated two-node service, so no disruptive timing
+run was attempted.
+
+## Decision gate
+
+Run the automatic path and tactic sweep on an isolated Spark, import the
+generated catalog into Atlas, then choose the next patch from measured data:
+
+- `reconstruction`: optimize bit extraction/register prefetch and K splitting;
+- `memory_bandwidth_or_pipeline`: prioritize packed layout, coalescing, and
+  multi-matrix fusion;
+- `compute_or_launch`: reduce cooperative synchronization/launch work or move
+  next to grouped MoE;
+- direct loses to reconstruct-then-GEMM: investigate a medium-M tile-local
+  decode/MMA kernel or move the crossover rather than forcing direct execution.
+
+Only after that gate should a new CUDA schedule be promoted into the runtime.
+
+## Upstream primary references
+
+- ExLlamaV3 production pin:
+  <https://github.com/turboderp-org/exllamav3/tree/c5d9c657966ffeeaa9353f0cc899f18629da4a13>
+- ExLlamaV3 current small-M kernel:
+  <https://github.com/turboderp-org/exllamav3/blob/0c49587a7c235e6303a6bbedc8b665272ad3a2ea/exllamav3/exllamav3_ext/quant/exl3_gemv_kernel.cuh>
+- ExLlamaV3 current int8-activation kernel:
+  <https://github.com/turboderp-org/exllamav3/blob/0c49587a7c235e6303a6bbedc8b665272ad3a2ea/exllamav3/exllamav3_ext/quant/exl3_gemv_int8_kernel.cuh>
+- QTIP: <https://arxiv.org/abs/2406.11235>
+- FLUTE: <https://arxiv.org/abs/2407.10960>
+- CUTLASS SM121 GEMM:
+  <https://github.com/NVIDIA/cutlass/blob/dc45f979ae336a235da1676b311f35efeb30149a/examples/79_blackwell_geforce_gemm/79a_blackwell_geforce_nvfp4_bf16_gemm.cu>
+- CUTLASS consumer-Blackwell MMA guidance:
+  <https://github.com/NVIDIA/cutlass/blob/dc45f979ae336a235da1676b311f35efeb30149a/python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py>
+- FlashInfer SM12x W4A16 MoE:
+  <https://github.com/flashinfer-ai/flashinfer/blob/2cc51dcf67ee71aade7074c64e84f13b7b7b117b/flashinfer/fused_moe/cute_dsl/blackwell_sm12x/moe_w4a16_kernel.py>

kernel_lab/README.md

@@ -0,0 +1,118 @@
+# Atlas EXL3 SM121 Kernel Lab
+
+This lab measures and tunes the existing ExLlamaV3 packed EXL3 CUDA path. It
+does not define a second weight format, materialize a full weight during direct
+timing, or put CUDA implementation code in the Atlas profiler.
+
+The workload interface is model-agnostic. Model IDs and operator names are
+receipt provenance; arbitrary future models provide their actual `M/N/K`
+shapes without being added to a registry.
+
+This lab is development-only source in this repository. The production
+Dockerfile neither copies nor installs `kernel_lab/`, so the lab is absent
+from the served image and no production launch path imports it. Nothing here
+is a speedup claim for the production pin: the SM121 tactic sweep has not yet
+run on target hardware.
+
+## Portable correctness
+
+The NumPy oracle supports the upstream MCG trellis at every integer bitrate
+K1-K8. It compares a materialized reference with a tile-streaming matmul that
+never constructs the complete decoded weight.
+
+```bash
+python3 -m kernel_lab.exl3.benchmark \
+  --backend cpu-oracle --bits 1 2 3 4 5 6 7 8
+python3 -m pytest tests/test_exl3_kernel_lab.py -q
+```
+
+## Arbitrary workload input
+
+Pass shapes directly as `OPERATOR:K:N`:
+
+```bash
+python3 -m kernel_lab.exl3.benchmark \
+  --backend sm121 \
+  --model-id publisher/new-model \
+  --shape experts.gate:7168:2304 \
+  --shape experts.down:2304:7168 \
+  --tp 4 --bits 2 3 4 --m 1 2 4 8
+```
+
+Or pass a JSON file exported from a profiler/model census:
+
+```json
+{
+  "workloads": [
+    {
+      "model_id": "publisher/new-model",
+      "operator": "layers.0.mlp.experts.gate",
+      "phase": "decode",
+      "k": 7168,
+      "n": 2304,
+      "tp_world_size": 4,
+      "grouped_moe": false
+    }
+  ]
+}
+```
+
+`m` belongs on the command line because one operator is normally swept over
+several routed-row counts. Workload-file objects therefore contain `K/N` and
+provenance, while `--m` controls the tuning sweep.
+
+## SM121 tactic sweep
+
+For every exact shape, the runner measures:
+
+- ExLlamaV3's automatic dispatch;
+- every compatible built-in GEMM tile shape;
+- quarter, half, and full-device SM quotas;
+- dense FP16, reconstruction-only, and reconstruct-then-GEMM diagnostics;
+- correctness, direct latency, effective bandwidth/TFLOP/s, reconstruction
+  overhead, and a bottleneck classification.
+
+Direct timings call `exl3_gemm` on packed trellis tensors. Dense and
+reconstruction baselines are outside the CUDA event range. The output is an
+`atlas.kernel-catalog/v1` document containing exact
+`atlas.kernel-benchmark/v1` receipts:
+
+```bash
+python3 -m kernel_lab.exl3.benchmark \
+  --backend sm121 \
+  --shape linear.square:4096:4096 \
+  --bits 3 --m 1 2 4 8 \
+  --warmup 10 --iterations 50 \
+  --output /tmp/atlas-exl3-sm121.json
+```
+
+Use `--no-tactic-sweep` to time only automatic dispatch. The optional legacy
+fixtures under `--fixture` reproduce early GLM/Hy4 measurements; they are not
+the supported-model set.
+
+## Upstream comparison
+
+The production image remains pinned to ExLlamaV3 `c5d9c657` (0.0.43). Building
+and self-checking an alternate immutable revision stays an isolated experiment
+in a source checkout of this repository: the production Dockerfile carries no
+benchmark build argument, commit stamp or copy of `kernel_lab/` for it, and
+this branch does not silently replace the production pin.
+
+Measured receipts take their runtime-source provenance from `--runtime-commit`
+or from the Git identity of the clean checkout that produced them. The
+`RUNTIME_SOURCE_COMMIT` environment override is honored only when the caller
+set it to the exact source commit that is loaded; it is a declaration, not
+attestation. If no full SHA can be resolved, the runner fails instead of
+emitting a receipt with unknown provenance.
+
+The current upstream candidate at
+`0c49587a7c235e6303a6bbedc8b665272ad3a2ea` adds the QTIP-style small-M GEMV
+and fused mul1 int8-activation GEMV. A developer-provided backend must set
+`EXLLAMAV3_COMMIT` to its actual commit before the metadata module is imported;
+the pinned fallback describes only the production pin and cannot prove an
+arbitrary installed extension's identity. Unknown or mismatched backend/source
+identity is not measurement evidence.
+
+Grouped MoE, prefill, mixed EXL3/NVFP4 dispatch, vLLM/SGLang-wide integration,
+attention, KV, scheduling, and networking remain outside this first measured
+decode gate.

kernel_lab/__init__.py

@@ -0,0 +1 @@
+"""Atlas Runtime Compiler kernel-lab prototypes."""

kernel_lab/exl3/__init__.py

@@ -0,0 +1,21 @@
+"""Stable EXL3 ABI, reference oracle, and direct-runtime adapter."""
+
+from .abi import MCG_MARKER_SIGNED_INT32, PackedExl3Metadata, PackedExl3Weight
+from .reference import (
+    matmul_materialized_reference,
+    matmul_streaming_reference,
+    pack_trellis,
+    reconstruct_weight,
+    unpack_trellis,
+)
+
+__all__ = [
+    "MCG_MARKER_SIGNED_INT32",
+    "PackedExl3Metadata",
+    "PackedExl3Weight",
+    "matmul_materialized_reference",
+    "matmul_streaming_reference",
+    "pack_trellis",
+    "reconstruct_weight",
+    "unpack_trellis",
+]