feat: execute FlashKDA on RTX 5090

PROGRESS.md

@@ -8,7 +8,7 @@
 |---|---:|---:|---|
 | 研究框架与规范 | 进行中 | 83% | Scaling Laws 二轮拟合复现与逐图精读 |
 | 网站设计系统 | 进行中 | 89% | 打印样式与更多通用可视化组件 |
-| Kimi K3 深读 | 三轮实证进行中 | 92% | 匹配 CUDA 12.9+ 执行 FlashKDA，并接入真实 hidden-state / expert-load traces |
+| Kimi K3 深读 | 三轮实证进行中 | 94% | 接入真实 hidden-state / expert-load / cache traces，并重绘报告数值图 |
 | 语言模型前史 | 完成首版 | 78% | Kneser–Ney、LSTM、Bahdanau 逐图精读与真实小语料复现 |
 | Transformer 基础 | 完成首版 | 79% | 多头电路、归一化 traces 与真实 kernel / KV 配置 |
 | 表示、位置与残差高速公路 | 完成首版 | 81% | 真实 hidden-state / norm traces、长上下文位置外推与深层稳定性消融 |
@@ -164,10 +164,14 @@
 - [x] 第三轮证据快照、可复现探针脚本与正式审计账本已进入开源树；原始权重字节不提交，真实观测 O、推导 D、执行 X、合成 S 与未决 U 分开标记。
 - [x] K3 新四视图本地真实 Chrome 回归通过：93 层条带、tensor group、参数分布、benchmark/router 切换、键盘 tabs、桌面与 390px 移动端均无异常。
 - [x] K3 三轮开放工件里程碑以源提交 `be2b291`、不可变镜像 `20260729T044605Z-be2b291` 发布；OCI digest `sha256:99aa953e…d00cdf`，NAS、VPS/Tailscale、NPM、DNS、HTTPS、证书、门户与十六套生产 Chrome 回归全链路通过；保留 `20260729T040336Z-b669615` 回滚。
+- [x] FlashKDA 隔离构建闭环：主机 CUDA 13.1 / glibc 2.43 仍在 `rsqrt` declarations 冲突；改用 CUDA 13.0.2 / Ubuntu 24.04 / glibc 2.39 后成功产出 CPython 3.12、`sm_120a` wheel，SHA-256 `14687b6d…2158d`。
+- [x] RTX 5090 正确性闸门通过：one chunk、partial tail、multi-chunk、96 heads 与 varlen 共 6 组，FlashKDA output / final state 和官方 `torch_ref.py` 逐元素完全相等，max absolute error 0。
+- [x] RTX 5090 K3 形状受控计时完成：fixed / varlen、BF16 / no-state / FP32 state 共 1,800 个 CUDA Event samples；fixed BF16 mean 2.6210 ms、P95 2.6437 ms，varlen BF16 mean 2.3335 ms、P95 2.3574 ms。
+- [x] 可复现 Dockerfile、运行脚本与机器可读 JSON 接入开源树；网站第四视图把作者 H20/GB200 表、本机 RTX 5090 值、exact suite、构建链、synthetic router 和 `A_log` 未决冲突分开显示。
 
 ## 正在进行
 
-- [ ] K3 三轮下一闸门：在匹配 CUDA 12.9+ 环境执行 FlashKDA correctness / benchmark，获得真实 token hidden states、expert load 与 cache traces，再做逐图数值重绘和独立小模型复现。
+- [ ] K3 三轮下一闸门：获得真实 token hidden states、expert load 与 cache traces，解释或修订 `A_log [128]` 工件冲突，再做 Figure 3/4/5 数值重绘和独立小模型复现。
 - [ ] DeepSeek 三轮：真实专家负载、MLA kernel、FP8 / pipeline 与 R1-like RL traces，外加独立小模型复现。
 - [ ] 表示、位置与残差二轮：真实 hidden-state / norm traces、长上下文位置外推复现与 mHC / AttnRes 深层稳定性消融。
 - [ ] 评测安全二轮：真实 cross-harness / pass@k 复跑、Judge 元评测、动态污染与过拒案例。
@@ -188,6 +192,8 @@
 
 | 日期 | 决策/发现 | 影响 |
 |---|---|---|
+| 2026-07-29 | K3 FlashKDA 在隔离 CUDA 13.0 / glibc 2.39 构建并回到 RTX 5090 执行 | 主机头文件 ABI 问题与 GPU 架构支持分离；wheel checksum、Dockerfile 与执行 JSON 可审计 |
+| 2026-07-29 | FlashKDA 本机实测永久与作者 H20 / GB200 表分账 | 6/6 exact 与 1,800 个 latency samples 可称 X；未跑本机 FLA 就不计算本机 speedup |
 | 2026-07-28 | 网站命名为 **LLM Atlas / 大模型技术全景** | 既能容纳 K3 深读，也能承载完整 LLM 课程 |
 | 2026-07-28 | K3 作为“汇流点”，不是课程起点 | 初学者可以先学基础，高阶读者可以从 K3 反向跳转 |
 | 2026-07-28 | 优先重绘论文图并标明“简化/改绘” | 图可缩放、可交互，也减少脱离上下文复制论文图片 |

README.md

@@ -21,11 +21,12 @@
 语言模型前史、Transformer 基础、表示/位置/残差、DeepSeek 技术谱系、Scaling Laws、数据工程、长上下文、MoE、指令微调与人类偏好、推理、工具使用与长程 Agent、原生多模态、训练系统、推理服务、数值优化，以及评测与安全深度专题，
 以及 71 个覆盖核心机制的原创交互视图。K3 二轮导读以 32 张问题账、16 图 / 5 表审计、
 8 个交互实验和 100 个一手/官方节点，完整覆盖架构、预训练、后训练、系统、评测、案例与附录。
-第三轮已完成首个开放工件里程碑：固定官方模型与 FlashKDA revisions，审计 96 个 checkpoint shards、
+第三轮已完成开放工件与首个真实 kernel 里程碑：固定官方模型与 FlashKDA revisions，审计 96 个 checkpoint shards、
 497,220 个 tensor entries、真实 KDA / MLA / MoE / MoonViT shapes 与小范围参数统计，并用 4 个新视图
-明确区分官方观测、确定性推导、本机执行、合成探针和未决矛盾。详见
+明确区分官方观测、确定性推导、本机执行、合成探针和未决矛盾；同时用隔离 CUDA 13.0 环境编译
+`sm_120a` wheel，在 RTX 5090 上完成 6/6 官方参考 exact-match 和 K3 fixed / varlen 形状计时。详见
 [K3_ARTIFACT_AUDIT.md](./research/K3_ARTIFACT_AUDIT.md) 与
-[checkpoint_probe.py](./experiments/k3/checkpoint_probe.py)。
+[checkpoint_probe.py](./experiments/k3/checkpoint_probe.py)、[FlashKDA probe](./experiments/k3/flashkda/)。
 DeepSeek 二轮专题以 24 张问题账、10 次技术转向、
 4 个交互实验和 60 个一手/官方节点，串起 Dense、MoE、MLA、V3 协同、R1 与 V4。
 其余专题按进度账本持续扩建。

experiments/k3/checkpoint_probe.py

@@ -370,11 +370,11 @@ def main() -> None:
                 else None,
                 "libc": list(platform.libc_ver()),
             },
-            "local_build": {
+            "baseline_host_build": {
-                "status": "blocked_before_kernel execution",
+                "status": "blocked in the default CUDA 12.8 environment; superseded by the separate CUDA 13 runtime probe",
                 "attempt_1": "system g++ 15 exceeds CUDA 12.8 host compiler range",
                 "attempt_2": "temporary g++ 13 reaches nvcc, then CUDA 12.8 headers conflict with current glibc math declarations",
-                "interpretation": "GPU architecture is listed by the repository, but the local CUDA 12.8 stack is below the official CUDA 12.9 requirement",
+                "interpretation": "this snapshot records the first host path only; see src/data/k3-flashkda-runtime.json for the successful isolated build and RTX 5090 execution",
             },
         },
     }

experiments/k3/flashkda/Dockerfile

@@ -0,0 +1,30 @@
+FROM nvidia/cuda:13.0.2-devel-ubuntu24.04
+
+ENV DEBIAN_FRONTEND=noninteractive
+
+RUN apt-get update \
+    && apt-get install -y --no-install-recommends \
+      build-essential \
+      git \
+      ninja-build \
+      python3 \
+      python3-dev \
+      python3-pip \
+      python3-venv \
+    && rm -rf /var/lib/apt/lists/*
+
+RUN python3 -m venv /opt/flashkda \
+    && /opt/flashkda/bin/python -m pip install --upgrade \
+      pip \
+      setuptools \
+      wheel \
+      ninja \
+      numpy \
+      pytest \
+      --index-url https://pypi.tuna.tsinghua.edu.cn/simple \
+    && /opt/flashkda/bin/python -m pip install \
+      torch==2.11.0 \
+      --index-url https://download.pytorch.org/whl/cu130
+
+ENV PATH=/opt/flashkda/bin:${PATH}
+ENV CUDA_HOME=/usr/local/cuda

experiments/k3/flashkda/README.md

@@ -0,0 +1,85 @@
+# FlashKDA RTX 5090 execution probe
+
+This directory reproduces the local execution evidence shown in the K3
+artifact lab. It does **not** download or run the 1.56 TB Kimi K3 checkpoint.
+All inputs are deterministic synthetic tensors with shapes allowed by the
+official FlashKDA API.
+
+Pinned upstream revision:
+
+```text
+MoonshotAI/FlashKDA@1ce47ea3bb22c84eb9cc665028399cf35e8ffb0b
+```
+
+## Why the build is isolated
+
+The workstation uses glibc 2.43. CUDA 13.1's published Linux support matrix
+currently lists distributions up to glibc 2.41, and compiling this extension
+directly on the host reaches an `rsqrt` / `rsqrtf` exception-specification
+conflict in the CUDA and system math headers. The container fixes the build
+ABI at Ubuntu 24.04 / glibc 2.39 while still targeting `sm_120a`.
+
+## Build the wheel
+
+Clone FlashKDA with submodules and verify the revision before building:
+
+```bash
+git clone --recursive https://github.com/MoonshotAI/FlashKDA.git /tmp/FlashKDA
+git -C /tmp/FlashKDA checkout 1ce47ea3bb22c84eb9cc665028399cf35e8ffb0b
+git -C /tmp/FlashKDA submodule update --init --recursive
+
+docker build -t llm-atlas-flashkda-cu130 experiments/k3/flashkda
+mkdir -p /tmp/flashkda-wheelhouse
+docker run --rm \
+  -e FLASH_KDA_ARCHS=120a \
+  -e MAX_JOBS=12 \
+  -v /tmp/FlashKDA:/src:ro \
+  -v /tmp/flashkda-wheelhouse:/wheelhouse \
+  llm-atlas-flashkda-cu130 \
+  python -m pip wheel /src --no-build-isolation --no-deps -w /wheelhouse
+```
+
+The audited wheel was built for CPython 3.12 and has SHA-256:
+
+```text
+14687b6d84a256d4552f0c73ccf93a601be582aeabcdf49ae3a409266872158d
+```
+
+## Run the probe
+
+Use CPython 3.12 with PyTorch 2.11.0+cu130 and install the wheel. Prebuild the
+small CUDA helper loaded by upstream `tests/torch_ref.py` without needing a
+Docker GPU runtime:
+
+```bash
+mkdir -p /tmp/k3-torch-extensions
+docker run --rm \
+  --user "$(id -u):$(id -g)" \
+  -e HOME=/tmp \
+  -e TORCH_EXTENSIONS_DIR=/cache \
+  -e TORCH_CUDA_ARCH_LIST=12.0a \
+  -e MAX_JOBS=12 \
+  -v "$PWD:/atlas:ro" \
+  -v /tmp/k3-torch-extensions:/cache \
+  llm-atlas-flashkda-cu130 \
+  python /atlas/experiments/k3/flashkda/build_reference_helper.py
+```
+
+Then run the GPU probe on the host:
+
+```bash
+export TORCH_EXTENSIONS_DIR=/tmp/k3-torch-extensions
+export TORCH_CUDA_ARCH_LIST=12.0a
+export CUDA_HOME=/usr/local/cuda
+
+python experiments/k3/flashkda/run_probe.py \
+  --flashkda-source /tmp/FlashKDA \
+  --wheel /tmp/flashkda-wheelhouse/flash_kda-0.0.1+1ce47ea-cp312-cp312-linux_x86_64.whl \
+  --output src/data/k3-flashkda-runtime.json
+```
+
+The correctness suite imports the upstream reference implementation instead
+of copying it into this repository. It checks exact BF16 equality at one
+chunk, a partial tail chunk, multiple chunks, 96 heads, and a variable-length
+batch. The performance cases are local kernel timings, not a comparison with
+the authors' H20 or GB200 tables.

experiments/k3/flashkda/build_reference_helper.py

@@ -0,0 +1,50 @@
+#!/usr/bin/env python3
+"""Prebuild the CUDA helper used verbatim by upstream tests/torch_ref.py.
+
+This runs in the CUDA build container without a GPU. Keeping the extension
+name, sources, generated function, and compiler flags identical lets PyTorch
+reuse the cache when the official reference module is imported on the host.
+"""
+
+from torch.utils.cpp_extension import load_inline
+
+
+CUDA_SOURCE = r"""
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+__global__ void sigmoid_tanh_fp32_kernel(const float* __restrict__ input,
+                                         float* __restrict__ output, int n) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx < n) {
+        float xh = input[idx] * 0.5f;
+        float th;
+        asm("tanh.approx.f32 %0, %1;" : "=f"(th) : "f"(xh));
+        output[idx] = th * 0.5f + 0.5f;
+    }
+}
+
+torch::Tensor sigmoid_tanh_fp32(torch::Tensor input) {
+    auto output = torch::empty_like(input);
+    int n = input.numel();
+    sigmoid_tanh_fp32_kernel<<<(n + 255) / 256, 256>>>(
+        input.data_ptr<float>(), output.data_ptr<float>(), n);
+    return output;
+}
+"""
+
+
+def main() -> None:
+    module = load_inline(
+        name="sigmoid_ext",
+        cpp_sources="torch::Tensor sigmoid_tanh_fp32(torch::Tensor input);",
+        cuda_sources=CUDA_SOURCE,
+        functions=["sigmoid_tanh_fp32"],
+        extra_cuda_cflags=["-O2"],
+        verbose=True,
+    )
+    print(module.__file__)
+
+
+if __name__ == "__main__":
+    main()

experiments/k3/flashkda/run_probe.py

@@ -0,0 +1,454 @@
+#!/usr/bin/env python3
+"""Run auditable FlashKDA correctness and latency probes on one CUDA GPU.
+
+The upstream torch reference is imported from a pinned FlashKDA checkout. The
+script deliberately uses valid synthetic A_log[H] tensors: it does not resolve
+the public Kimi K3 checkpoint's A_log[128] versus H=96 inconsistency.
+"""
+
+from __future__ import annotations
+
+import argparse
+import hashlib
+import importlib.metadata
+import json
+import math
+import platform
+import statistics
+import subprocess
+import sys
+import time
+from datetime import datetime, timezone
+from pathlib import Path
+from typing import Any, Callable
+
+import torch
+import torch.nn.functional as F
+
+import flash_kda
+import flash_kda_C
+
+
+def parse_args() -> argparse.Namespace:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--flashkda-source", type=Path, required=True)
+    parser.add_argument("--wheel", type=Path, required=True)
+    parser.add_argument("--output", type=Path, required=True)
+    parser.add_argument("--warmup", type=int, default=20)
+    parser.add_argument("--iters", type=int, default=100)
+    parser.add_argument("--repeats", type=int, default=3)
+    parser.add_argument("--seed", type=int, default=20260729)
+    parser.add_argument("--captured-at", default=None)
+    return parser.parse_args()
+
+
+def sha256(path: Path) -> str:
+    digest = hashlib.sha256()
+    with path.open("rb") as handle:
+        for block in iter(lambda: handle.read(1024 * 1024), b""):
+            digest.update(block)
+    return digest.hexdigest()
+
+
+def percentile(sorted_values: list[float], fraction: float) -> float:
+    if not sorted_values:
+        return float("nan")
+    index = fraction * (len(sorted_values) - 1)
+    lower = math.floor(index)
+    upper = math.ceil(index)
+    if lower == upper:
+        return sorted_values[lower]
+    weight = index - lower
+    return sorted_values[lower] * (1 - weight) + sorted_values[upper] * weight
+
+
+def nvidia_smi() -> dict[str, Any]:
+    fields = [
+        "name",
+        "driver_version",
+        "memory.total",
+        "power.limit",
+        "clocks.max.sm",
+        "clocks.max.memory",
+    ]
+    output = subprocess.check_output(
+        [
+            "nvidia-smi",
+            f"--query-gpu={','.join(fields)}",
+            "--format=csv,noheader,nounits",
+        ],
+        text=True,
+    ).strip()
+    values = [value.strip() for value in output.split(",")]
+    return dict(zip(fields, values, strict=True))
+
+
+def make_inputs(
+    sequence_lengths: list[int],
+    heads: int,
+    seed: int,
+) -> dict[str, torch.Tensor | float | None]:
+    batch = 1
+    dimension = 128
+    total_tokens = sum(sequence_lengths)
+    sequences = len(sequence_lengths)
+    generator = torch.Generator(device="cuda").manual_seed(seed)
+
+    q = F.normalize(
+        torch.randn(
+            (batch, total_tokens, heads, dimension),
+            dtype=torch.float32,
+            device="cuda",
+            generator=generator,
+        ),
+        p=2,
+        dim=-1,
+    ).to(torch.bfloat16)
+    k = F.normalize(
+        torch.randn(
+            (batch, total_tokens, heads, dimension),
+            dtype=torch.float32,
+            device="cuda",
+            generator=generator,
+        ),
+        p=2,
+        dim=-1,
+    ).to(torch.bfloat16)
+    v = torch.randn(
+        (batch, total_tokens, heads, dimension),
+        dtype=torch.bfloat16,
+        device="cuda",
+        generator=generator,
+    )
+    g = torch.randn(
+        (batch, total_tokens, heads, dimension),
+        dtype=torch.bfloat16,
+        device="cuda",
+        generator=generator,
+    )
+    beta = torch.randn(
+        (batch, total_tokens, heads),
+        dtype=torch.bfloat16,
+        device="cuda",
+        generator=generator,
+    )
+    a_log = torch.rand(
+        heads, dtype=torch.float32, device="cuda", generator=generator
+    )
+    dt_bias = torch.rand(
+        (heads, dimension),
+        dtype=torch.float32,
+        device="cuda",
+        generator=generator,
+    )
+    initial_state = torch.randn(
+        (sequences, heads, dimension, dimension),
+        dtype=torch.bfloat16,
+        device="cuda",
+        generator=generator,
+    )
+    cu_seqlens = None
+    if len(sequence_lengths) > 1:
+        offsets = [0]
+        for length in sequence_lengths:
+            offsets.append(offsets[-1] + length)
+        cu_seqlens = torch.tensor(offsets, dtype=torch.long, device="cuda")
+
+    return {
+        "q": q,
+        "k": k,
+        "v": v,
+        "g": g,
+        "beta": beta,
+        "A_log": a_log,
+        "dt_bias": dt_bias,
+        "initial_state": initial_state,
+        "cu_seqlens": cu_seqlens,
+        "scale": 1 / math.sqrt(dimension),
+    }
+
+
+def run_correctness_case(
+    torch_ref: Callable[..., None],
+    name: str,
+    sequence_lengths: list[int],
+    heads: int,
+    seed: int,
+) -> dict[str, Any]:
+    inputs = make_inputs(sequence_lengths, heads, seed)
+    q = inputs["q"]
+    assert isinstance(q, torch.Tensor)
+    initial_state = inputs["initial_state"]
+    assert isinstance(initial_state, torch.Tensor)
+
+    out_kernel = torch.zeros_like(q)
+    out_reference = torch.zeros_like(q)
+    state_kernel = torch.zeros_like(initial_state)
+    state_reference = torch.zeros_like(initial_state)
+    common = {
+        "A_log": inputs["A_log"],
+        "dt_bias": inputs["dt_bias"],
+        "lower_bound": -5.0,
+        "initial_state": initial_state.clone(),
+        "cu_seqlens": inputs["cu_seqlens"],
+    }
+
+    start = time.perf_counter()
+    flash_kda.fwd(
+        inputs["q"],
+        inputs["k"],
+        inputs["v"],
+        inputs["g"],
+        inputs["beta"],
+        inputs["scale"],
+        out_kernel,
+        final_state=state_kernel,
+        **common,
+    )
+    torch.cuda.synchronize()
+    kernel_ms = (time.perf_counter() - start) * 1000
+
+    start = time.perf_counter()
+    torch_ref(
+        inputs["q"],
+        inputs["k"],
+        inputs["v"],
+        inputs["g"],
+        inputs["beta"],
+        inputs["scale"],
+        out_reference,
+        final_state=state_reference,
+        **common,
+    )
+    torch.cuda.synchronize()
+    reference_ms = (time.perf_counter() - start) * 1000
+
+    output_diff = (out_kernel.float() - out_reference.float()).abs()
+    state_diff = (state_kernel.float() - state_reference.float()).abs()
+    result = {
+        "name": name,
+        "sequence_lengths": sequence_lengths,
+        "heads": heads,
+        "dimension": 128,
+        "kernel_ms_first_measured": kernel_ms,
+        "reference_ms": reference_ms,
+        "output_exact": bool(torch.equal(out_kernel, out_reference)),
+        "state_exact": bool(torch.equal(state_kernel, state_reference)),
+        "output_max_abs_diff": output_diff.max().item(),
+        "output_mean_abs_diff": output_diff.mean().item(),
+        "state_max_abs_diff": state_diff.max().item(),
+        "state_mean_abs_diff": state_diff.mean().item(),
+    }
+    if not result["output_exact"] or not result["state_exact"]:
+        raise AssertionError(f"exact correctness failed: {result}")
+    return result
+
+
+def bench_events(
+    function: Callable[[], None],
+    warmup: int,
+    iters: int,
+    repeats: int,
+) -> list[float]:
+    for _ in range(max(warmup, 1)):
+        function()
+    torch.cuda.synchronize()
+
+    elapsed: list[float] = []
+    for _ in range(repeats):
+        starts = [torch.cuda.Event(enable_timing=True) for _ in range(iters)]
+        ends = [torch.cuda.Event(enable_timing=True) for _ in range(iters)]
+        torch.cuda.synchronize()
+        for index in range(iters):
+            starts[index].record()
+            function()
+            ends[index].record()
+        torch.cuda.synchronize()
+        elapsed.extend(
+            start.elapsed_time(end) for start, end in zip(starts, ends, strict=True)
+        )
+    return elapsed
+
+
+def summarize_latency(
+    values: list[float],
+    total_tokens: int,
+) -> dict[str, Any]:
+    ordered = sorted(float(value) for value in values)
+    mean = statistics.fmean(ordered)
+    return {
+        "samples": len(ordered),
+        "mean_ms": mean,
+        "min_ms": ordered[0],
+        "p50_ms": percentile(ordered, 0.50),
+        "p95_ms": percentile(ordered, 0.95),
+        "max_ms": ordered[-1],
+        "sequence_tokens_per_second": total_tokens / (mean / 1000),
+    }
+
+
+def run_benchmark_case(
+    name: str,
+    sequence_lengths: list[int],
+    heads: int,
+    seed: int,
+    warmup: int,
+    iters: int,
+    repeats: int,
+) -> dict[str, Any]:
+    torch.cuda.empty_cache()
+    torch.cuda.reset_peak_memory_stats()
+    inputs = make_inputs(sequence_lengths, heads, seed)
+    q = inputs["q"]
+    initial_state = inputs["initial_state"]
+    assert isinstance(q, torch.Tensor)
+    assert isinstance(initial_state, torch.Tensor)
+    out = torch.zeros_like(q)
+    state_bf16 = torch.zeros_like(initial_state)
+    initial_fp32 = initial_state.float()
+    state_fp32 = torch.zeros_like(initial_fp32)
+    common = {
+        "A_log": inputs["A_log"],
+        "dt_bias": inputs["dt_bias"],
+        "lower_bound": -5.0,
+        "cu_seqlens": inputs["cu_seqlens"],
+    }
+
+    def invoke(initial: torch.Tensor | None, final: torch.Tensor | None) -> None:
+        flash_kda.fwd(
+            inputs["q"],
+            inputs["k"],
+            inputs["v"],
+            inputs["g"],
+            inputs["beta"],
+            inputs["scale"],
+            out,
+            initial_state=initial,
+            final_state=final,
+            **common,
+        )
+
+    variants = {
+        "bf16_state": lambda: invoke(initial_state, state_bf16),
+        "no_state": lambda: invoke(None, None),
+        "fp32_state": lambda: invoke(initial_fp32, state_fp32),
+    }
+    timings = {
+        variant: summarize_latency(
+            bench_events(function, warmup, iters, repeats),
+            sum(sequence_lengths),
+        )
+        for variant, function in variants.items()
+    }
+    torch.cuda.synchronize()
+    return {
+        "name": name,
+        "sequence_lengths": sequence_lengths,
+        "total_tokens": sum(sequence_lengths),
+        "heads": heads,
+        "dimension": 128,
+        "warmup": warmup,
+        "iters": iters,
+        "repeats": repeats,
+        "timings": timings,
+        "output_abs_mean_after_last_run": out.float().abs().mean().item(),
+        "peak_allocated_mib": torch.cuda.max_memory_allocated() / 2**20,
+        "peak_reserved_mib": torch.cuda.max_memory_reserved() / 2**20,
+    }
+
+
+def main() -> None:
+    args = parse_args()
+    if not torch.cuda.is_available():
+        raise RuntimeError("CUDA GPU is required")
+    if torch.cuda.get_device_capability(0) < (9, 0):
+        raise RuntimeError("FlashKDA requires SM90 or newer")
+    tests_dir = args.flashkda_source / "tests"
+    sys.path.insert(0, str(tests_dir))
+    from torch_ref import torch_ref
+
+    revision = subprocess.check_output(
+        ["git", "-C", str(args.flashkda_source), "rev-parse", "HEAD"],
+        text=True,
+    ).strip()
+    torch.manual_seed(args.seed)
+    torch.cuda.manual_seed_all(args.seed)
+
+    correctness_specs = [
+        ("one_chunk", [16], 1),
+        ("partial_tail", [17], 1),
+        ("two_chunks", [32], 1),
+        ("multi_chunk_multi_head", [65], 2),
+        ("k3_head_count", [17], 96),
+        ("varlen_partial_chunks", [17, 31], 2),
+    ]
+    correctness = [
+        run_correctness_case(torch_ref, name, lengths, heads, args.seed + index)
+        for index, (name, lengths, heads) in enumerate(correctness_specs)
+    ]
+
+    benchmark_specs = [
+        ("teaching_scale", [512], 8),
+        ("intermediate", [2048], 32),
+        ("k3_fixed_shape", [8192], 96),
+        ("k3_varlen_shape", [1300, 547, 2048, 963, 271, 3063], 96),
+    ]
+    benchmarks = [
+        run_benchmark_case(
+            name,
+            lengths,
+            heads,
+            args.seed + 100 + index,
+            args.warmup,
+            args.iters,
+            args.repeats,
+        )
+        for index, (name, lengths, heads) in enumerate(benchmark_specs)
+    ]
+
+    captured_at = args.captured_at or datetime.now(timezone.utc).isoformat()
+    result = {
+        "schema_version": 1,
+        "captured_at": captured_at,
+        "evidence_identity": "X / local execution on deterministic synthetic tensors",
+        "boundary": {
+            "k3_checkpoint_loaded": False,
+            "real_token_hidden_states": False,
+            "a_log_shape_conflict_resolved": False,
+            "benchmark_comparison": "local FlashKDA timings only; author H20/GB200 tables remain separate",
+        },
+        "provenance": {
+            "flashkda_revision": revision,
+            "flashkda_package": importlib.metadata.version("flash-kda"),
+            "wheel_filename": args.wheel.name,
+            "wheel_sha256": sha256(args.wheel),
+            "runner": str(Path(__file__).relative_to(Path.cwd())),
+        },
+        "environment": {
+            "python": platform.python_version(),
+            "platform": platform.platform(),
+            "libc": list(platform.libc_ver()),
+            "torch": torch.__version__,
+            "torch_cuda": torch.version.cuda,
+            "gpu": torch.cuda.get_device_name(0),
+            "capability": list(torch.cuda.get_device_capability(0)),
+            "nvidia_smi": nvidia_smi(),
+            "flash_kda_module": flash_kda.__file__,
+            "flash_kda_extension": flash_kda_C.__file__,
+        },
+        "correctness": {
+            "all_exact": all(
+                case["output_exact"] and case["state_exact"]
+                for case in correctness
+            ),
+            "cases": correctness,
+        },
+        "benchmarks": benchmarks,
+    }
+    args.output.parent.mkdir(parents=True, exist_ok=True)
+    args.output.write_text(json.dumps(result, indent=2) + "\n")
+    print(json.dumps(result, indent=2))
+
+
+if __name__ == "__main__":
+    main()

research/K3_ARTIFACT_AUDIT.md

@@ -22,7 +22,7 @@
 - 没有在单张 RTX 5090 上加载 1.56 TB checkpoint；
 - 没有获得真实 token hidden states、线上 expert load 或生产 cache trace；
 - 没有把随机向量上的 router 行为写成真实数据分布；
-- 没有把未成功执行的 FlashKDA kernel 写成“本机 benchmark”；
+- 没有把合法 synthetic shape 上的 FlashKDA 实测写成“K3 checkpoint 已跑通”；
 - 没有替官方解释下面发现的 `A_log` 形状不一致。
 
 ## 2. 一手工件与校验
@@ -268,7 +268,7 @@ RMS=1 的各向同性随机 hidden vectors，再比较 top-16：
 
 ### 7.1 官方仓库事实
 
-FlashKDA `1ce47ea3`：
+[FlashKDA 官方仓库](https://github.com/MoonshotAI/FlashKDA) `1ce47ea3`：
 
 - CUTLASS kernels；
 - 支持 `90a / 100a / 103a / 120a`；
@@ -288,51 +288,127 @@ FlashKDA `1ce47ea3`：
 
 这些是作者仓库 benchmark，不是本站复跑值。
 
-### 7.2 本机真实构建边界
+### 7.2 从主机构建失败到隔离 wheel
 
-本机：
+GPU 主机：
 
 - RTX 5090，compute capability `12.0`；
-- PyTorch `2.11.0+cu128`；
+- Ubuntu 26.04，glibc `2.43`；
-- PyTorch CUDA `12.8`；
+- 常用环境为 PyTorch `2.11.0+cu128` / CUDA `12.8`；
 - 官方源码明确包含 `sm_120a`，所以不是 GPU architecture 缺失。
 
-两次可复现构建：
+主机路径依次暴露了两个不同问题：
 
 1. 系统 `g++ 15.2`：PyTorch extension 在编译前拒绝，CUDA 12.8 要求 host compiler `<14`；
 2. 临时解包 `g++ 13.4`：成功进入 nvcc，但 CUDA 12.8 headers 与当前 glibc math declarations
-   在 `cospi / sinpi / rsqrt` exception specification 处冲突。
+   冲突；
+3. 新建 PyTorch `2.11.0+cu130` 环境并改用 CUDA 13.1 后，版本主线已经匹配，但仍在
+   `rsqrt / rsqrtf` exception specification 处失败。
 
-结论：
+第三步说明“升级到 CUDA 13”还不够。[CUDA 13.1 官方 Linux support matrix](https://docs.nvidia.com/cuda/archive/13.1.0/cuda-installation-guide-linux/index.html) 当前列出的
+glibc 上限是 2.41，而主机是 2.43；NVIDIA 官方论坛也记录了同类 CUDA 13.1 / 新 glibc
+数学声明[冲突案例](https://forums.developer.nvidia.com/t/fedora-43-and-nvcc-cuda13-1-error-exception-specification-is-incompatible-rsqrt-rsqrtf/354510)。这里能写的是“构建失败与超出验证矩阵一致”，不能写成 CUDA 对未来 glibc
+永久不兼容。
 
-- kernel 尚未在本站机器执行；
+最终构建固定为：
-- 失败与 README 的 CUDA 12.9+ 要求一致；
+
-- 不能把 `sm_120a` 支持写成本机已经跑通；
+| layer | pinned value |
-- 下一次应使用匹配 PyTorch 的 CUDA 12.9+ toolchain 或官方容器后再复跑 correctness + benchmark。
+|---|---|
+| base image | `nvidia/cuda:13.0.2-devel-ubuntu24.04` |
+| container libc | glibc `2.39` |
+| Python | `3.12` |
+| PyTorch | `2.11.0+cu130` |
+| target | `sm_120a` |
+| FlashKDA | `1ce47ea3bb22c84eb9cc665028399cf35e8ffb0b` |
+| wheel | `flash_kda-0.0.1+1ce47ea-cp312-cp312-linux_x86_64.whl` |
+| wheel SHA-256 | `14687b6d84a256d4552f0c73ccf93a601be582aeabcdf49ae3a409266872158d` |
+
+容器没有 NVIDIA runtime，也不需要 GPU 才能为固定 architecture 编译。CUDA 编译器实际产出
+`sm_120a` 对象，最终 wheel 为 3,806,174 bytes。wheel 随后装入主机独立 CPython 3.12.11 /
+PyTorch 2.11.0+cu130 环境；该运行环境能看到 RTX 5090、capability 12.0，并成功导入
+`flash_kda` 与 `flash_kda_C`。
+
+### 7.3 exact correctness：输出与 recurrent state 都必须相等
+
+正确性对照直接调用官方仓库的 `tests/torch_ref.py`，而不是本站重写一份近似参考。
+测试输入满足公开 API 的 `A_log[H]`；它们故意不使用 checkpoint 的 `A_log[128]`，因此不会
+偷偷替未决形状冲突下结论。
+
+| case | sequence lengths | H | output exact | final state exact | max abs error |
+|---|---|---:|---:|---:|---:|
+| one chunk | `[16]` | 1 | yes | yes | 0 |
+| partial tail | `[17]` | 1 | yes | yes | 0 |
+| two chunks | `[32]` | 1 | yes | yes | 0 |
+| multi-chunk / multi-head | `[65]` | 2 | yes | yes | 0 |
+| K3 head count | `[17]` | 96 | yes | yes | 0 |
+| varlen partial chunks | `[17,31]` | 2 | yes | yes | 0 |
+
+六组用 BF16 q/k/v/g、BF16 beta logits、F32 `A_log/dt_bias` 和 BF16 recurrent state；
+output 与 final state 全部逐元素相等，所有 max / mean absolute difference 都是 0。
+这把证据从“源码声称支持 120a”推进成了“本站 RTX 5090 实际执行并通过官方参考”。
+
+### 7.4 本机受控计时：只报本站 FlashKDA，不伪造本机 FLA 对照
+
+协议：
+
+- deterministic synthetic tensors；
+- `D=128`，K3 shape 使用 `H=96, T_total=8192`；
+- warmup 20；
+- 每轮 100 次、3 repeats，即每个 state mode 300 个 CUDA Event samples；
+- output buffer 重用；
+- 分开测 BF16 state、无 state I/O、FP32 state；
+- throughput 是 sequence tokens/s，不是 head-tokens/s；
+- 没有安装并复跑 FLA，所以不计算本站 speedup。
+
+| RTX 5090 case | state mode | mean | P95 | sequence tokens/s |
+|---|---|---:|---:|---:|
+| fixed `[8192]` | BF16 | 2.6210 ms | 2.6437 ms | 3.126 M/s |
+| fixed `[8192]` | none | 2.6154 ms | 2.6362 ms | 3.132 M/s |
+| fixed `[8192]` | FP32 | 2.6267 ms | 2.6500 ms | 3.119 M/s |
+| varlen `[1300,547,2048,963,271,3063]` | BF16 | 2.3335 ms | 2.3574 ms | 3.511 M/s |
+| same varlen | none | 2.3104 ms | 2.3254 ms | 3.546 M/s |
+| same varlen | FP32 | 2.3596 ms | 2.3845 ms | 3.472 M/s |
+
+fixed case 三种 mode 合并的 peak allocated 为 1,755.7 MiB；varlen 为 1,845.7 MiB。
+作者 H20 fixed 的 `2.6220 ms` 与本站 BF16 fixed 的 `2.6210 ms` 数值接近只是一次观测，
+不能据此宣布 RTX 5090 与 H20 等价：软件栈、频率、功耗、输入、测量时段和对照 kernel
+都不足以支持这种外推。
 
 ## 8. 可复现实验入口
 
-脚本：
+checkpoint 工件脚本：
 
 ```text
 experiments/k3/checkpoint_probe.py
 ```
 
+FlashKDA 构建与执行：
+
+```text
+experiments/k3/flashkda/Dockerfile
+experiments/k3/flashkda/build_reference_helper.py
+experiments/k3/flashkda/run_probe.py
+experiments/k3/flashkda/README.md
+```
+
 提交的数据快照：
 
 ```text
 src/data/k3-artifact-snapshot.json
+src/data/k3-flashkda-runtime.json
 ```
 
-脚本会：
+两条实验路径共同完成：
 
 1. 解析 config、index 与 selected headers；
 2. 校验小范围字节长度和 SHA-256；
 3. 统计真实 KDA / router 参数；
 4. 运行明确标注的 synthetic router counterexample；
 5. 解析 FlashKDA 官方 H20 / GB200 benchmark；
-6. 输出本机环境与构建边界；
+6. 在隔离 CUDA 13.0 / glibc 2.39 环境编译 `sm_120a` wheel；
-7. 不提交原始权重。
+7. 用官方 torch reference 跑六组 exact-match；
+8. 在 RTX 5090 跑 fixed / varlen、三种 state mode 的受控计时；
+9. 不提交原始权重或二进制 wheel，只提交 checksum、环境、代码和 JSON 结果。
 
 ## 9. 网站实现合同
 
@@ -341,7 +417,7 @@ src/data/k3-artifact-snapshot.json
 1. **Layer map**：93 层真实 config 条带；显示 KDA/MLA、dense/MoE、AttnRes block。
 2. **Tensor anatomy**：checkpoint / expert / vision / MLA tensor shape 与数量。
 3. **Parameter audit**：真实 Range statistics，并把 `A_log` mismatch 放在主视区。
-4. **Reproduction boundary**：官方 benchmark、本站构建失败点、synthetic router counterexample。
+4. **Reproduction boundary**：作者 benchmark、本站 exact / latency、构建链、未决形状冲突与 synthetic router counterexample。
 
 每个视图必须显示证据类型：
 
@@ -359,7 +435,9 @@ src/data/k3-artifact-snapshot.json
 - [x] selected open-weight ranges 做真实参数统计；
 - [x] 发现并限定 `A_log` shape inconsistency；
 - [x] FlashKDA RTX 5090 构建尝试留下可复现边界；
-- [ ] 使用 CUDA 12.9+ 匹配环境跑 FlashKDA exact correctness；
+- [x] 使用 CUDA 13.0 / glibc 2.39 隔离环境产出可校验 `sm_120a` wheel；
+- [x] RTX 5090 上六组 FlashKDA / official torch reference exact correctness；
+- [x] K3 fixed / varlen shape、三种 state mode 的 1,800 个 CUDA Event samples；
 - [ ] 取得真实 hidden-state / router load trace；
 - [ ] 取得可加载的 reduced checkpoint、官方 trace 或多机资源；
 - [ ] 对 Figure 3 / 4 / 5 做真实数值重绘；

scripts/check-k3-browser.mjs

@@ -221,15 +221,24 @@ const artifacts = await evaluate(`(() => {
     flash: text("[data-benchmark-flash]"),
     fla: text("[data-benchmark-fla]"),
     speedup: text("[data-benchmark-speedup]"),
+    localMean: text("[data-local-mean]"),
+    localP95: text("[data-local-p95]"),
+    localThroughput: text("[data-local-throughput]"),
+    exactSuite: root.textContent.includes("6 / 6 PASS") && root.textContent.includes("MAX ABS ERROR"),
     cv: text("[data-router-cv]"),
     zero: text("[data-router-zero]"),
   };
   input("[data-benchmark-device]", "gb200");
   input("[data-benchmark-case]", "Varlen, \\\`seq_lens\\\`=\\\`1024 x 8\\\`");
+  input("[data-local-case]", "k3_varlen_shape");
+  input("[data-local-state]", "fp32_state");
   input("[data-router-mode]", "bias");
   const reproductionChanged = {
     flash: text("[data-benchmark-flash]"),
     speedup: text("[data-benchmark-speedup]"),
+    localMean: text("[data-local-mean]"),
+    localP95: text("[data-local-p95]"),
+    localMode: text("[data-local-mode-copy]"),
     cv: text("[data-router-cv]"),
     zero: text("[data-router-zero]"),
   };
@@ -253,6 +262,7 @@ await pause(180);
 await screenshot("/tmp/llm-atlas-k3-lab-desktop.png");
 await evaluate(`(() => {
   document.querySelector("[data-k3-artifact-lab]").scrollIntoView({ block: "start", behavior: "instant" });
+  document.querySelector('[data-artifact-tab="reproduction"]')?.click();
   window.scrollBy(0, -82);
 })()`);
 await pause(180);
@@ -297,6 +307,7 @@ await pause(180);
 await screenshot("/tmp/llm-atlas-k3-mobile.png");
 await evaluate(`(() => {
   document.querySelector("[data-k3-artifact-lab]").scrollIntoView({ block: "start", behavior: "instant" });
+  document.querySelector('[data-artifact-tab="reproduction"]')?.click();
   window.scrollBy(0, -64);
 })()`);
 await pause(180);
@@ -334,8 +345,8 @@ if (artifacts.tensors.panel !== "tensors" || artifacts.tensors.groups !== 3 || a
 if (!artifacts.mla.visible || artifacts.mla.rows !== 5 || !artifacts.mla.has576) failures.push("MLA header shape 视图异常");
 if (artifacts.parameterInitial.panel !== "parameters" || artifacts.parameterInitial.shape !== "[128] F32" || !artifacts.parameterInitial.conflict) failures.push("A_log 工件冲突审计异常");
 if (artifacts.parameterChanged.shape !== "[96,128] F32" || !artifacts.parameterChanged.count.includes("12,288")) failures.push("真实 dt_bias 参数切换异常");
-if (artifacts.reproductionInitial.panel !== "reproduction" || numeric(artifacts.reproductionInitial.speedup) !== 1.85 || numeric(artifacts.reproductionInitial.cv) < 2) failures.push("FlashKDA H20 或 router 初始探针异常");
+if (artifacts.reproductionInitial.panel !== "reproduction" || numeric(artifacts.reproductionInitial.speedup) !== 1.85 || numeric(artifacts.reproductionInitial.localMean) < 2.6 || !artifacts.reproductionInitial.exactSuite || numeric(artifacts.reproductionInitial.cv) < 2) failures.push("FlashKDA H20、本机 exact suite 或 router 初始探针异常");
-if (numeric(artifacts.reproductionChanged.speedup) !== 3.27 || numeric(artifacts.reproductionChanged.flash) !== 0.7064 || numeric(artifacts.reproductionChanged.cv) <= numeric(artifacts.reproductionInitial.cv) || numeric(artifacts.reproductionChanged.zero) <= numeric(artifacts.reproductionInitial.zero)) failures.push("GB200 benchmark 或 synthetic router counterexample 未更新");
+if (numeric(artifacts.reproductionChanged.speedup) !== 3.27 || numeric(artifacts.reproductionChanged.flash) !== 0.7064 || numeric(artifacts.reproductionChanged.localMean) >= numeric(artifacts.reproductionInitial.localMean) || !artifacts.reproductionChanged.localMode.includes("FP32 state") || numeric(artifacts.reproductionChanged.cv) <= numeric(artifacts.reproductionInitial.cv) || numeric(artifacts.reproductionChanged.zero) <= numeric(artifacts.reproductionInitial.zero)) failures.push("GB200 benchmark、本机 varlen/state 或 synthetic router counterexample 未更新");
 if (artifacts.keyboardSelected !== "tensors" || artifacts.keyboardVisible !== "tensors") failures.push("开放工件键盘 tab 导航异常");
 if (!mobile.menuVisible || mobile.menuOpen !== "true" || mobile.tabs !== 8 || mobile.artifactTabs !== 4 || mobile.artifactLayers !== 93) failures.push("移动端导航或实验异常");
 if (mobile.offenders.length) failures.push(`移动端越界元素：${JSON.stringify(mobile.offenders)}`);

src/components/K3ArtifactLab.astro

@@ -3,6 +3,7 @@ import {
   k3ArtifactEvidence,
   k3ArtifactLayers,
   k3ArtifactSnapshot as snapshot,
+  k3FlashKdaRuntime as runtime,
   k3ArtifactViews,
 } from "@/data/k3Artifacts";
 
@@ -10,6 +11,9 @@ const checkpoint = snapshot.checkpoint;
 const audit = snapshot.parameter_audit;
 const probe = snapshot.router_stress_probe;
 const flash = snapshot.flashkda;
+const localBenchmarks = runtime.benchmarks.filter((row) => row.name.startsWith("k3_"));
+const localInitial = localBenchmarks[0];
+const localInitialTiming = localInitial.timings.bf16_state;
 
 const bytes = (value: number) => {
   if (value >= 2 ** 40) return `${(value / 2 ** 40).toFixed(3)} TiB`;
@@ -263,8 +267,8 @@ const benchmarkDevices = [
 
   <section class="artifact-panel" data-artifact-panel="reproduction" hidden>
     <div class="panel-intro">
-      <div><span>X + S / WHAT ACTUALLY RAN</span><h4>作者 benchmark、本站编译尝试与合成反例，三者不能写成同一种实测</h4></div>
+      <div><span>X + O + S / WHAT ACTUALLY RAN</span><h4>作者表、本机 RTX 5090 实测与合成 router 反例，三条证据各自归位</h4></div>
-      <p>RTX 5090 是 sm_120，但本机 PyTorch CUDA 12.8 低于 FlashKDA README 的 12.9+；kernel 尚未执行。</p>
+      <p>FlashKDA 已用 CUDA 13.0 容器编译为 sm_120a wheel，并在本站 RTX 5090 上通过 exact-match 与 K3 形状计时；输入仍是合成 tensor，不是 checkpoint hidden state。</p>
     </div>
     <div class="repro-controls">
       <label><span>OFFICIAL DEVICE</span><select data-benchmark-device>
@@ -274,29 +278,54 @@ const benchmarkDevices = [
         <option value="Fixed">Fixed T=8192</option>
         <option value="Varlen, `seq_lens`=`1024 x 8`">8 × 1024 varlen</option>
       </select></label>
-      <label><span>ROUTER STRESS</span><select data-router-mode>
+      <label><span>LOCAL RTX 5090 CASE</span><select data-local-case>
-        <option value="raw">without correction bias</option>
+        <option value="k3_fixed_shape">Fixed · 8192 × 96 × 128</option>
-        <option value="bias">with checkpoint bias</option>
+        <option value="k3_varlen_shape">Varlen · 6 sequences / 8192 total</option>
+      </select></label>
+      <label><span>LOCAL STATE MODE</span><select data-local-state>
+        <option value="bf16_state">BF16 initial + final state</option>
+        <option value="no_state">No recurrent state I/O</option>
+        <option value="fp32_state">FP32 initial + final state</option>
       </select></label>
     </div>
     <div class="benchmark-readout">
       <article><span>FLASHKDA</span><b data-benchmark-flash>2.6220 ms</b><p>O / author repository</p></article>
       <article><span>FLA CHUNK KDA</span><b data-benchmark-fla>4.8388 ms</b><p>O / same author table</p></article>
       <article class="accent"><span>AUTHOR SPEEDUP</span><b data-benchmark-speedup>1.85×</b><p>不能外推到 RTX 5090</p></article>
-      <article class="dark"><span>LOCAL KERNEL</span><b>NOT RUN</b><p>CUDA 12.8 &lt; official 12.9+</p></article>
+      <article class="dark"><span>AUTHOR DEVICE</span><b data-benchmark-device-label>H20</b><p>与本站环境分开比较</p></article>
+    </div>
+    <div
+      data-local-benchmark-json={JSON.stringify(localBenchmarks)}
+      hidden
+    ></div>
+    <div class="benchmark-readout local-readout">
+      <article class="dark"><span>LOCAL MEAN</span><b data-local-mean>{localInitialTiming.mean_ms.toFixed(4)} ms</b><p>X / 300 CUDA-event samples</p></article>
+      <article><span>LOCAL P95</span><b data-local-p95>{localInitialTiming.p95_ms.toFixed(4)} ms</b><p data-local-mode-copy>BF16 state · fixed</p></article>
+      <article><span>SEQUENCE RATE</span><b data-local-throughput>{(localInitialTiming.sequence_tokens_per_second / 1e6).toFixed(3)}M/s</b><p>8192 sequence tokens / latency</p></article>
+      <article class="accent"><span>PEAK ALLOCATED</span><b data-local-memory>{localInitial.peak_allocated_mib.toFixed(1)} MiB</b><p>该 case 三种 state mode 合并峰值</p></article>
+    </div>
+    <div class="execution-gate">
+      <article><span>EXACT SUITE</span><b>{runtime.correctness.cases.length} / {runtime.correctness.cases.length} PASS</b><p>one chunk、tail、multi-head、96 heads、varlen。</p></article>
+      <article><span>MAX ABS ERROR</span><b>0.0</b><p>BF16 output 与 final state 均逐元素相等。</p></article>
+      <article><span>RUNTIME</span><b>Py {runtime.environment.python} · Torch {runtime.environment.torch}</b><p>CUDA {runtime.environment.torch_cuda} · sm_{runtime.environment.capability.join("")}</p></article>
+      <article><span>WHEEL SHA-256</span><b>{runtime.provenance.wheel_sha256.slice(0, 12)}…</b><p>{runtime.provenance.flashkda_revision.slice(0, 9)} · CPython 3.12</p></article>
     </div>
     <div class="local-build">
-      <article><span>ATTEMPT 01</span><b>g++ 15 rejected</b><p>CUDA 12.8 host compiler range要求 &lt;14。</p></article>
+      <article><span>HOST BUILD</span><b>glibc 2.43 blocked</b><p>CUDA math headers 与系统 `rsqrt/rsqrtf` exception declarations 冲突。</p></article>
       <i>→</i>
-      <article><span>ATTEMPT 02</span><b>g++ 13 reached nvcc</b><p>随后在 glibc math declarations 处与 CUDA 12.8 headers 冲突。</p></article>
+      <article><span>ISOLATED BUILD</span><b>sm_120a wheel passed</b><p>CUDA 13.0.2 · Ubuntu 24.04 · glibc 2.39；wheel 3.81 MB。</p></article>
       <i>→</i>
-      <article class="warn"><span>NEXT GATE</span><b>CUDA 12.9+ matched env</b><p>再跑 exact correctness 与本机 benchmark。</p></article>
+      <article class="warn"><span>GPU EXECUTION</span><b>RTX 5090 exact 6/6</b><p>fixed / varlen K3 形状各 900 次计时，三种 state mode。</p></article>
     </div>
     <div class="router-counterexample">
       <div>
         <span>S / REAL WEIGHTS, SYNTHETIC HIDDEN</span>
         <h5>随机 RMS=1 输入为什么不能评价 Quantile Balancing</h5>
         <p>2,048 个固定 seed 向量通过真实 `896×7168` router；它们不是模型 token hidden states。</p>
+        <label><span>ROUTER STRESS</span><select data-router-mode>
+          <option value="raw">without correction bias</option>
+          <option value="bias">with checkpoint bias</option>
+        </select></label>
       </div>
       <div class="router-stats">
         <p><span>LOAD CV</span><b data-router-cv>{probe.without_correction_bias.cv.toFixed(3)}</b></p>
@@ -305,7 +334,7 @@ const benchmarkDevices = [
         <p><span>TOP-16 OVERLAP</span><b>{probe.membership_overlap_mean.toFixed(2)} / 16</b></p>
       </div>
     </div>
-    <div class="boundary"><b>X/S boundary</b><p>编译失败是本站真实执行结果；router counterexample 只证明 hidden distribution 不可省略，不证明真实 QB 变好或变坏。</p></div>
+    <div class="boundary"><b>X/S/U boundary</b><p>本机 kernel 实测只验证公开 FlashKDA API 与合成合法 shape；没有加载 K3 checkpoint，也不解决 checkpoint `A_log [128]` 与 API `[96]` 的冲突。Router counterexample 仍只证明 hidden distribution 不可省略。</p></div>
   </section>
 
   <footer class="evidence-strip">
@@ -425,10 +454,44 @@ const benchmarkDevices = [
       put("[data-benchmark-flash]", `${row.flash.toFixed(4)} ms`);
       put("[data-benchmark-fla]", `${row.fla.toFixed(4)} ms`);
       put("[data-benchmark-speedup]", `${row.speedup.toFixed(2)}×`);
+      put("[data-benchmark-device-label]", device.value.toUpperCase());
     };
     device.addEventListener("input", renderBenchmark);
     benchmarkCase.addEventListener("input", renderBenchmark);
 
+    type LocalTiming = {
+      mean_ms: number;
+      p95_ms: number;
+      sequence_tokens_per_second: number;
+    };
+    type LocalBenchmark = {
+      name: string;
+      peak_allocated_mib: number;
+      timings: Record<"bf16_state" | "no_state" | "fp32_state", LocalTiming>;
+    };
+    const localRows = JSON.parse(
+      $<HTMLElement>("[data-local-benchmark-json]").dataset.localBenchmarkJson ?? "[]",
+    ) as LocalBenchmark[];
+    const localCase = $<HTMLSelectElement>("[data-local-case]");
+    const localState = $<HTMLSelectElement>("[data-local-state]");
+    const localModeLabels = {
+      bf16_state: "BF16 state",
+      no_state: "no state I/O",
+      fp32_state: "FP32 state",
+    };
+    const renderLocalBenchmark = () => {
+      const row = localRows.find((item) => item.name === localCase.value)!;
+      const state = localState.value as keyof typeof localModeLabels;
+      const timing = row.timings[state];
+      put("[data-local-mean]", `${timing.mean_ms.toFixed(4)} ms`);
+      put("[data-local-p95]", `${timing.p95_ms.toFixed(4)} ms`);
+      put("[data-local-throughput]", `${(timing.sequence_tokens_per_second / 1e6).toFixed(3)}M/s`);
+      put("[data-local-memory]", `${row.peak_allocated_mib.toFixed(1)} MiB`);
+      put("[data-local-mode-copy]", `${localModeLabels[state]} · ${row.name.includes("varlen") ? "varlen" : "fixed"}`);
+    };
+    localCase.addEventListener("input", renderLocalBenchmark);
+    localState.addEventListener("input", renderLocalBenchmark);
+
     const routerMode = $<HTMLSelectElement>("[data-router-mode]");
     const routerRows = {
       raw: { cv: 2.0845208168, gini: .8310886025, zero: 558 },
@@ -441,6 +504,9 @@ const benchmarkDevices = [
       put("[data-router-zero]", row.zero);
     };
     routerMode.addEventListener("input", renderRouter);
+    renderBenchmark();
+    renderLocalBenchmark();
+    renderRouter();
   });
 </script>
 
@@ -656,7 +722,8 @@ const benchmarkDevices = [
   .local-build .warn { border-color: var(--signal); background: color-mix(in srgb, var(--signal) 8%, var(--paper)); }
   .parameter-control label { display: grid; grid-template-columns: 1fr; gap: 9px; }
   .parameter-control select,
-  .repro-controls select { width: 100%; padding: 9px; border: 1px solid var(--line); color: var(--ink); background: var(--paper-raised); font: .62rem var(--mono); }
+  .repro-controls select,
+  .router-counterexample select { width: 100%; padding: 9px; border: 1px solid var(--line); color: var(--ink); background: var(--paper-raised); font: .62rem var(--mono); }
   .parameter-control > p { display: flex; justify-content: space-between; gap: 16px; margin: 0; font: .62rem var(--mono); }
   .parameter-control > p b { color: var(--copper); }
   .distribution { margin: 0 28px 24px; padding: 20px; border: 1px solid var(--line); background: var(--paper); }
@@ -684,11 +751,27 @@ const benchmarkDevices = [
   .hypothesis-card p span { margin-top: 8px; color: var(--muted); font-size: .6rem; }
   .repro-controls {
     display: grid;
-    grid-template-columns: repeat(3, 1fr);
+    grid-template-columns: repeat(4, 1fr);
     gap: 12px;
     margin: 24px 28px;
   }
   .repro-controls label { display: grid; gap: 9px; padding: 14px; border: 1px solid var(--line); background: var(--paper); }
+  .local-readout { margin-bottom: 12px; }
+  .execution-gate {
+    display: grid;
+    grid-template-columns: repeat(4, 1fr);
+    gap: 12px;
+    margin: 0 28px 24px;
+    padding: 12px;
+    border: 1px solid var(--ink);
+    background: var(--ink);
+  }
+  .execution-gate article { min-width: 0; padding: 14px; border: 1px solid rgba(255,255,255,.18); }
+  .execution-gate span,
+  .execution-gate b { display: block; font-family: var(--mono); }
+  .execution-gate span { color: var(--copper); font-size: .55rem; letter-spacing: .08em; }
+  .execution-gate b { margin-top: 9px; overflow-wrap: anywhere; color: var(--paper); font-size: .72rem; line-height: 1.35; }
+  .execution-gate p { margin: 9px 0 0; color: rgba(255,255,255,.6); font-size: .6rem; line-height: 1.5; }
   .local-build {
     display: grid;
     grid-template-columns: 1fr 24px 1fr 24px 1fr;
@@ -708,6 +791,7 @@ const benchmarkDevices = [
     background: var(--paper);
   }
   .router-counterexample > div > p { color: var(--muted); font-size: .67rem; line-height: 1.6; }
+  .router-counterexample label { display: grid; gap: 8px; margin-top: 14px; color: var(--muted); font: .58rem/1.4 var(--mono); }
   .router-stats { display: grid; grid-template-columns: repeat(2, 1fr); gap: 8px; }
   .router-stats p { margin: 0; padding: 12px; border: 1px solid var(--line); }
   .router-stats span,
@@ -747,7 +831,8 @@ const benchmarkDevices = [
     .layer-readout,
     .artifact-metrics,
     .benchmark-readout,
-    .tensor-ledger { grid-template-columns: repeat(2, 1fr); }
+    .tensor-ledger,
+    .execution-gate { grid-template-columns: repeat(2, 1fr); }
     .tensor-selector { align-items: stretch; flex-wrap: wrap; }
     .tensor-selector > span { width: 100%; }
     .tensor-table { overflow-x: auto; }
@@ -781,11 +866,13 @@ const benchmarkDevices = [
     .repro-controls,
     .local-build,
     .router-counterexample,
+    .execution-gate,
     .boundary { margin-left: 16px; margin-right: 16px; }
     .layer-readout,
     .artifact-metrics,
     .benchmark-readout,
     .tensor-ledger,
+    .execution-gate,
     .router-stats { grid-template-columns: 1fr; }
   }
 </style>

src/data/k3-artifact-snapshot.json

@@ -590,11 +590,11 @@
         "2.43"
       ]
     },
-    "local_build": {
+    "baseline_host_build": {
-      "status": "blocked_before_kernel execution",
+      "status": "blocked in the default CUDA 12.8 environment; superseded by the separate CUDA 13 runtime probe",
       "attempt_1": "system g++ 15 exceeds CUDA 12.8 host compiler range",
       "attempt_2": "temporary g++ 13 reaches nvcc, then CUDA 12.8 headers conflict with current glibc math declarations",
-      "interpretation": "GPU architecture is listed by the repository, but the local CUDA 12.8 stack is below the official CUDA 12.9 requirement"
+      "interpretation": "this snapshot records the first host path only; see src/data/k3-flashkda-runtime.json for the successful isolated build and RTX 5090 execution"
     }
   }
 }

src/data/k3-flashkda-runtime.json

@@ -0,0 +1,328 @@
+{
+  "schema_version": 1,
+  "captured_at": "2026-07-29T05:22:14.267504+00:00",
+  "evidence_identity": "X / local execution on deterministic synthetic tensors",
+  "boundary": {
+    "k3_checkpoint_loaded": false,
+    "real_token_hidden_states": false,
+    "a_log_shape_conflict_resolved": false,
+    "benchmark_comparison": "local FlashKDA timings only; author H20/GB200 tables remain separate"
+  },
+  "provenance": {
+    "flashkda_revision": "1ce47ea3bb22c84eb9cc665028399cf35e8ffb0b",
+    "flashkda_package": "0.0.1+1ce47ea",
+    "wheel_filename": "flash_kda-0.0.1+1ce47ea-cp312-cp312-linux_x86_64.whl",
+    "wheel_sha256": "14687b6d84a256d4552f0c73ccf93a601be582aeabcdf49ae3a409266872158d",
+    "runner": "experiments/k3/flashkda/run_probe.py"
+  },
+  "environment": {
+    "python": "3.12.11",
+    "platform": "Linux-7.0.0-28-generic-x86_64-with-glibc2.43",
+    "libc": [
+      "glibc",
+      "2.43"
+    ],
+    "torch": "2.11.0+cu130",
+    "torch_cuda": "13.0",
+    "gpu": "NVIDIA GeForce RTX 5090",
+    "capability": [
+      12,
+      0
+    ],
+    "nvidia_smi": {
+      "name": "NVIDIA GeForce RTX 5090",
+      "driver_version": "595.84",
+      "memory.total": "32607",
+      "power.limit": "600.00",
+      "clocks.max.sm": "3105",
+      "clocks.max.memory": "14001"
+    },
+    "flash_kda_module": "/tmp/k3-flashkda-run-cu130/lib/python3.12/site-packages/flash_kda/__init__.py",
+    "flash_kda_extension": "/tmp/k3-flashkda-run-cu130/lib/python3.12/site-packages/flash_kda_C.cpython-312-x86_64-linux-gnu.so"
+  },
+  "correctness": {
+    "all_exact": true,
+    "cases": [
+      {
+        "name": "one_chunk",
+        "sequence_lengths": [
+          16
+        ],
+        "heads": 1,
+        "dimension": 128,
+        "kernel_ms_first_measured": 1.9076780008617789,
+        "reference_ms": 173.11289900681004,
+        "output_exact": true,
+        "state_exact": true,
+        "output_max_abs_diff": 0.0,
+        "output_mean_abs_diff": 0.0,
+        "state_max_abs_diff": 0.0,
+        "state_mean_abs_diff": 0.0
+      },
+      {
+        "name": "partial_tail",
+        "sequence_lengths": [
+          17
+        ],
+        "heads": 1,
+        "dimension": 128,
+        "kernel_ms_first_measured": 0.09024899918586016,
+        "reference_ms": 2.4010630149859935,
+        "output_exact": true,
+        "state_exact": true,
+        "output_max_abs_diff": 0.0,
+        "output_mean_abs_diff": 0.0,
+        "state_max_abs_diff": 0.0,
+        "state_mean_abs_diff": 0.0
+      },
+      {
+        "name": "two_chunks",
+        "sequence_lengths": [
+          32
+        ],
+        "heads": 1,
+        "dimension": 128,
+        "kernel_ms_first_measured": 0.051337992772459984,
+        "reference_ms": 1.9857370061799884,
+        "output_exact": true,
+        "state_exact": true,
+        "output_max_abs_diff": 0.0,
+        "output_mean_abs_diff": 0.0,
+        "state_max_abs_diff": 0.0,
+        "state_mean_abs_diff": 0.0
+      },
+      {
+        "name": "multi_chunk_multi_head",
+        "sequence_lengths": [
+          65
+        ],
+        "heads": 2,
+        "dimension": 128,
+        "kernel_ms_first_measured": 0.0607699912507087,
+        "reference_ms": 5.533189018024132,
+        "output_exact": true,
+        "state_exact": true,
+        "output_max_abs_diff": 0.0,
+        "output_mean_abs_diff": 0.0,
+        "state_max_abs_diff": 0.0,
+        "state_mean_abs_diff": 0.0
+      },
+      {
+        "name": "k3_head_count",
+        "sequence_lengths": [
+          17
+        ],
+        "heads": 96,
+        "dimension": 128,
+        "kernel_ms_first_measured": 0.05647100624628365,
+        "reference_ms": 87.05275200190954,
+        "output_exact": true,
+        "state_exact": true,
+        "output_max_abs_diff": 0.0,
+        "output_mean_abs_diff": 0.0,
+        "state_max_abs_diff": 0.0,
+        "state_mean_abs_diff": 0.0
+      },
+      {
+        "name": "varlen_partial_chunks",
+        "sequence_lengths": [
+          17,
+          31
+        ],
+        "heads": 2,
+        "dimension": 128,
+        "kernel_ms_first_measured": 0.1232630165759474,
+        "reference_ms": 4.736603004857898,
+        "output_exact": true,
+        "state_exact": true,
+        "output_max_abs_diff": 0.0,
+        "output_mean_abs_diff": 0.0,
+        "state_max_abs_diff": 0.0,
+        "state_mean_abs_diff": 0.0
+      }
+    ]
+  },
+  "benchmarks": [
+    {
+      "name": "teaching_scale",
+      "sequence_lengths": [
+        512
+      ],
+      "total_tokens": 512,
+      "heads": 8,
+      "dimension": 128,
+      "warmup": 20,
+      "iters": 100,
+      "repeats": 3,
+      "timings": {
+        "bf16_state": {
+          "samples": 300,
+          "mean_ms": 0.06259957360724608,
+          "min_ms": 0.06054399907588959,
+          "p50_ms": 0.06233600154519081,
+          "p95_ms": 0.06326559744775295,
+          "max_ms": 0.21401600539684296,
+          "sequence_tokens_per_second": 8178969.4481358975
+        },
+        "no_state": {
+          "samples": 300,
+          "mean_ms": 0.0588647465283672,
+          "min_ms": 0.056992001831531525,
+          "p50_ms": 0.05873600021004677,
+          "p95_ms": 0.05920000001788139,
+          "max_ms": 0.0729919970035553,
+          "sequence_tokens_per_second": 8697905.456082527
+        },
+        "fp32_state": {
+          "samples": 300,
+          "mean_ms": 0.0638275195658207,
+          "min_ms": 0.062463998794555664,
+          "p50_ms": 0.0631679967045784,
+          "p95_ms": 0.06521599739789963,
+          "max_ms": 0.08054400235414505,
+          "sequence_tokens_per_second": 8021618.315780101
+        }
+      },
+      "output_abs_mean_after_last_run": 0.0026643681339919567,
+      "peak_allocated_mib": 18.63720703125,
+      "peak_reserved_mib": 28.0
+    },
+    {
+      "name": "intermediate",
+      "sequence_lengths": [
+        2048
+      ],
+      "total_tokens": 2048,
+      "heads": 32,
+      "dimension": 128,
+      "warmup": 20,
+      "iters": 100,
+      "repeats": 3,
+      "timings": {
+        "bf16_state": {
+          "samples": 300,
+          "mean_ms": 0.4393576521674792,
+          "min_ms": 0.4354879856109619,
+          "p50_ms": 0.43958398699760437,
+          "p95_ms": 0.441633602976799,
+          "max_ms": 0.4665600061416626,
+          "sequence_tokens_per_second": 4661350.47357573
+        },
+        "no_state": {
+          "samples": 300,
+          "mean_ms": 0.43100970675547917,
+          "min_ms": 0.42633599042892456,
+          "p50_ms": 0.4307839870452881,
+          "p95_ms": 0.43350398540496826,
+          "max_ms": 0.4801599979400635,
+          "sequence_tokens_per_second": 4751633.125427203
+        },
+        "fp32_state": {
+          "samples": 300,
+          "mean_ms": 0.44364384045203525,
+          "min_ms": 0.44067201018333435,
+          "p50_ms": 0.4436799883842468,
+          "p95_ms": 0.4461440145969391,
+          "max_ms": 0.45952001214027405,
+          "sequence_tokens_per_second": 4616315.641649082
+        }
+      },
+      "output_abs_mean_after_last_run": 0.0026791957207024097,
+      "peak_allocated_mib": 158.26611328125,
+      "peak_reserved_mib": 208.0
+    },
+    {
+      "name": "k3_fixed_shape",
+      "sequence_lengths": [
+        8192
+      ],
+      "total_tokens": 8192,
+      "heads": 96,
+      "dimension": 128,
+      "warmup": 20,
+      "iters": 100,
+      "repeats": 3,
+      "timings": {
+        "bf16_state": {
+          "samples": 300,
+          "mean_ms": 2.620986862977346,
+          "min_ms": 2.6054399013519287,
+          "p50_ms": 2.6176319122314453,
+          "p95_ms": 2.6436815142631533,
+          "max_ms": 2.6599678993225098,
+          "sequence_tokens_per_second": 3125540.274816252
+        },
+        "no_state": {
+          "samples": 300,
+          "mean_ms": 2.615440630118052,
+          "min_ms": 2.5992319583892822,
+          "p50_ms": 2.6129279136657715,
+          "p95_ms": 2.6361759066581727,
+          "max_ms": 2.6476480960845947,
+          "sequence_tokens_per_second": 3132168.211224218
+        },
+        "fp32_state": {
+          "samples": 300,
+          "mean_ms": 2.626678284804026,
+          "min_ms": 2.6074559688568115,
+          "p50_ms": 2.6238080263137817,
+          "p95_ms": 2.650022292137146,
+          "max_ms": 2.667167901992798,
+          "sequence_tokens_per_second": 3118767.931113877
+        }
+      },
+      "output_abs_mean_after_last_run": 0.0026801086496561766,
+      "peak_allocated_mib": 1755.67236328125,
+      "peak_reserved_mib": 2208.0
+    },
+    {
+      "name": "k3_varlen_shape",
+      "sequence_lengths": [
+        1300,
+        547,
+        2048,
+        963,
+        271,
+        3063
+      ],
+      "total_tokens": 8192,
+      "heads": 96,
+      "dimension": 128,
+      "warmup": 20,
+      "iters": 100,
+      "repeats": 3,
+      "timings": {
+        "bf16_state": {
+          "samples": 300,
+          "mean_ms": 2.3335195755958558,
+          "min_ms": 2.312864065170288,
+          "p50_ms": 2.330944061279297,
+          "p95_ms": 2.357441592216492,
+          "max_ms": 2.3840320110321045,
+          "sequence_tokens_per_second": 3510576.935232353
+        },
+        "no_state": {
+          "samples": 300,
+          "mean_ms": 2.3104187711079915,
+          "min_ms": 2.2939839363098145,
+          "p50_ms": 2.3101279735565186,
+          "p95_ms": 2.3254063010215758,
+          "max_ms": 2.332927942276001,
+          "sequence_tokens_per_second": 3545677.563929858
+        },
+        "fp32_state": {
+          "samples": 300,
+          "mean_ms": 2.3596167453130086,
+          "min_ms": 2.3351359367370605,
+          "p50_ms": 2.3561919927597046,
+          "p95_ms": 2.3845438957214355,
+          "max_ms": 2.3907198905944824,
+          "sequence_tokens_per_second": 3471750.2392166285
+        }
+      },
+      "output_abs_mean_after_last_run": 0.0026924293488264084,
+      "peak_allocated_mib": 1845.6728515625,
+      "peak_reserved_mib": 2214.0
+    }
+  ]
+}

src/data/k3Artifacts.ts

@@ -1,6 +1,8 @@
 import snapshot from "./k3-artifact-snapshot.json";
+import runtime from "./k3-flashkda-runtime.json";
 
 export { snapshot as k3ArtifactSnapshot };
+export { runtime as k3FlashKdaRuntime };
 
 const mlaLayerSet = new Set(snapshot.configuration.mla_layers);
 
@@ -24,7 +26,7 @@ export const k3ArtifactViews = [
   ["layers", "01", "LAYER MAP", "93 层配置"],
   ["tensors", "02", "TENSOR ANATOMY", "497,220 entries"],
   ["parameters", "03", "PARAMETER AUDIT", "真实权重小切片"],
-  ["reproduction", "04", "REPRODUCTION", "作者值与本机边界"],
+  ["reproduction", "04", "REPRODUCTION", "RTX 5090 执行证据"],
 ] as const;
 
 export const k3ArtifactEvidence = [

src/pages/index.astro

@@ -132,13 +132,14 @@ const paths = [
           <h2>47 页不再压成摘要：再把 1.56 TB 开放工件接回报告</h2>
           <p>
             在三十二张报告问题账之外，继续审计 96 个 safetensors 分片、497,220 个 tensor entries、
-            真实 KDA / MLA / MoE / MoonViT shape、小范围权重统计、FlashKDA 编译边界与未决形状矛盾。
+            真实 KDA / MLA / MoE / MoonViT shape、小范围权重统计，并把 FlashKDA 推进到 RTX 5090
+            6/6 exact-match、K3 fixed / varlen 计时与未决 checkpoint 形状矛盾。
           </p>
         </div>
         <dl>
           <div><dt>REPORT</dt><dd>16 Figures · 5 Tables</dd></div>
           <div><dt>ARTIFACTS</dt><dd>96 shards · 497,220 entries</dd></div>
-          <div><dt>LAB</dt><dd>8 报告实验 + 4 工件视图</dd></div>
+          <div><dt>KERNEL</dt><dd>sm_120a · exact 6/6</dd></div>
         </dl>
         <span class="release-arrow" aria-hidden="true">从报告目录进入开放工件证据链 →</span>
       </a>

src/pages/k3/index.astro

@@ -854,11 +854,12 @@ const paperGroups = [
         <p class="lede">
           第三轮固定到官方 Hugging Face revision，读取 config、remote code、60 MB tensor index、
           四个 safetensors headers 和两个小范围参数切片。原始权重不进入本站仓库；
-          结构、shape、计数、参数统计与本机编译边界都可以从公开脚本重复生成。
+          结构、shape、计数、参数统计、隔离 wheel 与 RTX 5090 执行结果都可以从公开脚本重复生成。
         </p>
         <div class="artifact-callout">
           <article><span>O / OBSERVED</span><b>1.4196 TiB tensor data</b><p>96 shards、497,220 entries；不是运行显存，也不是参数量口径。</p></article>
           <article><span>D / CLOSED LOOP</span><b>69 KDA · 24 MLA · 92 MoE</b><p>配置、tensor names 与 header shape 三方闭合。</p></article>
+          <article><span>X / RTX 5090</span><b>exact 6/6 · max error 0</b><p>官方 torch reference；fixed BF16 mean 2.6210 ms。</p></article>
           <article class="warning"><span>U / UNRESOLVED</span><b>A_log [128] ≠ expected [96]</b><p>checkpoint 与公开代码 / kernel API 的形状冲突保留在主视区，不擅自解释。</p></article>
         </div>
         <K3ArtifactLab />
@@ -866,6 +867,7 @@ const paperGroups = [
           <a class="button primary" href="https://huggingface.co/moonshotai/Kimi-K3">打开官方开放权重</a>
           <a class="button" href="https://github.com/MoonshotAI/FlashKDA">打开 FlashKDA 官方实现</a>
           <a class="button" href="https://github.com/MoonshotAI/FlashKDA/blob/master/BENCHMARK_GB200.md">核对作者 GB200 benchmark</a>
+          <a class="button" href="https://git.k1412.top/wuyang/llm-atlas/src/branch/main/experiments/k3/flashkda">复跑本站 RTX 5090 探针</a>
         </div>
       </section>
 
@@ -927,7 +929,7 @@ const paperGroups = [
     .anchor-alias { position: relative; top: -88px; display: block; visibility: hidden; }
     .artifact-callout {
       display: grid;
-      grid-template-columns: repeat(3, minmax(0, 1fr));
+      grid-template-columns: repeat(4, minmax(0, 1fr));
       max-width: 1080px;
       margin: 32px 0;
       border-top: 1px solid var(--line);

src/pages/progress/index.astro

@@ -9,7 +9,7 @@ const researching = chapters.filter((chapter) => ["researching", "drafting"].inc
 const workstreams = [
   { label: "研究框架与规范", value: 83, next: "给 Scaling 与推理专题补逐篇图表/实验精读层级" },
   { label: "网站设计系统", value: 89, next: "打印样式与更多通用可视化组件" },
-  { label: "Kimi K3 深读", value: 92, next: "在匹配 CUDA 12.9+ 环境执行 FlashKDA，并接入真实 hidden-state / expert-load traces" },
+  { label: "Kimi K3 深读", value: 94, next: "接入真实 hidden-state / expert-load / cache traces，并重绘报告数值图" },
   { label: "语言模型前史", value: 78, next: "逐图精读 Kneser–Ney、LSTM 与 Bahdanau，并加入真实小语料复现" },
   { label: "Transformer 基础", value: 79, next: "逐图精读多头电路、Pre/Post-LN 与真实 kernel / KV 配置" },
   { label: "表示、位置与残差高速公路", value: 81, next: "加入真实 hidden-state / norm traces、长上下文位置外推复现与更多深层稳定性消融" },
@@ -105,6 +105,7 @@ const workstreams = [
       <article><span>✓</span><h3>DeepSeek 技术谱系二轮深读</h3><p>二十四张问题账、十次技术转向、60 个一手/官方节点，以及稀疏容量—MLA 缓存—V3 协同—RL 偏差四联实验。</p></article>
       <article><span>✓</span><h3>Kimi K3 技术报告二轮深读</h3><p>三十二张问题账、Figure 1–16 / Table 1–5 审计、100 节点阅读链，以及 Delta—Decay—AttnRes—LatentMoE—SiTU—QB—MOPD—Cache 八联实验。</p></article>
       <article><span>✓</span><h3>Kimi K3 三轮开放工件里程碑</h3><p>固定官方 revisions，审计 96 个 shards、497,220 个 tensor entries 与真实 KDA / MLA / MoE / MoonViT shapes；四联实验分开显示层型、tensor anatomy、参数范围和复现边界。</p></article>
+      <article><span>✓</span><h3>FlashKDA RTX 5090 执行闸门</h3><p>隔离 CUDA 13.0 / glibc 2.39 编译 sm_120a wheel；6/6 官方参考逐元素相等，并完成 fixed / varlen、三种 state mode 的 1,800 个 CUDA Event samples。</p></article>
       <article><span>✓</span><h3>Scaling Laws 深度专题</h3><p>九张账、29 个一手节点、DeepSeek/Kimi 双谱系与曲面—部署—复用—涌现四联实验。</p></article>
       <article><span>✓</span><h3>数据工程深度专题</h3><p>十二张账、31 个一手节点、DeepSeek/Kimi 双谱系与流水线—去重—混合—改写四联实验。</p></article>
       <article><span>✓</span><h3>长上下文深度专题</h3><p>五张成本账、26 篇一手论文、10+ 机制图与 8 策略交互实验室。</p></article>
@@ -212,6 +213,8 @@ const workstreams = [
       <div><time>2026-07-29</time><b>K3 图表与实验永久分级</b><p>Figure 1–16 / Table 1–5 建立视觉契约；报告事实、原论文、确定性推导与教学模型使用 R/P/D/T 四种身份。</p></div>
       <div><time>2026-07-29</time><b>K3 开放工件按五种证据身份审计</b><p>真实观测 O、确定性推导 D、本机执行 X、合成探针 S 与未决矛盾 U 分开；作者 benchmark 不冒充本站实测。</p></div>
       <div><time>2026-07-29</time><b>A_log 形状冲突保持未决</b><p>checkpoint 的 [128] 与 config / remote code / FlashKDA API 期待的 [96] 并列展示；不宣布权重损坏，也不把 channel-wise 假设写成真实 forward。</p></div>
+      <div><time>2026-07-29</time><b>FlashKDA 编译与执行永久分两道闸门</b><p>容器产出 sm_120a wheel 只证明可编译；RTX 5090 的 6/6 official-reference exact suite 通过后，才把证据升级为本机执行 X。</p></div>
+      <div><time>2026-07-29</time><b>作者表与 RTX 5090 表永久分账</b><p>H20 / GB200 保持 O；本站只报告独立环境、协议、300 samples/mode 和延迟分布，未跑本机 FLA 就不写本机 speedup。</p></div>
     </div>
   </section>