我开源了 automegakernel —— 将任意 HuggingFace 模型编译成一个持久的单一巨型内核。批量大小为 1 的解码受带宽限制。正常执行会为每个算子启动一个内核，并且每层在 HBM 中来回传输激活值数十次。这种开销就是全部问题所在。automegakernel 将整个前向传播融合到一次启动中。一次启动 = 一次前向 = 一个 token。难点在于一个跨越所有 SM 的单一内核，仅靠计数器同步，这是一个死锁/竞态地雷区。因此核心是一个静态验证器，在启动前证明任何调度是无死锁且无竞态的。Agent 可以自由编辑调度，但无法交付一个挂起的内核。7160 个对抗性调度中，6091 个不安全，零误接受。同一份源码可重新目标到 sm_80 / sm_90 / sm_120。在 HuggingFace greedy decode 中逐 token 复现真正的 SmolLM2-135M 模型。搜索发现的 int8 megakernel 在 batch-1 下击败了 CUDA-graphed cuBLAS bf16：L4 最高 1.33 倍，L40S 1.25-1.27 倍。在 A100/H100 上则落后，我们如实说明。目前仅限 Llama 系列 :p sc: http://github.com/RightNow-AI/AutoMegaKernel… 论文: https://arxiv.org/abs/2606.09682

RightNow-AI/AutoMegaKernel

源码: https://github.com/RightNow-AI/AutoMegaKernel

AutoMegaKernel (AMK)

一个用于 GPU 巨型内核合成的通用 Agent 框架。 一个编码 Agent（Claude Code / Codex） 驱动 AMK 将模型编译成一个可证明正确、可自重新目标、整个前向传播融合为单个持久启动的巨型内核， 然后自动调优，并且在每次运行时变得更好。

仓库: https://github.com/RightNow-AI/AutoMegaKernel · 许可证: MIT · 开放研究框架（Enterprise / Forge）

amk compile --gpu --regime single-stream # 导入 -> 降级 -> 验证（无死锁无竞态）-> 与 eager 对比验证 -> 构建 GPU 巨型内核
# -> 测量其相对于 HBM 屋顶线的性能 -> 输出一个正确的巨型内核 + 报告

AMK 是 AutoKernel (https://github.com/RightNow-AI/autokernel) 的姊妹项目：AutoKernel 自动生成单个最优内核；AMK 自动生成整个模型的最优巨型内核。它继承了 AutoKernel 的自动研究循环（提议 → 固定评估 → 保留/回退，数小时不间断），并增加了一个新的搜索维度：调度。

当前覆盖范围: HuggingFace Llama 系列在 CUDA 上（sm_75 到 sm_120）。 路线图（未来工作）: 将导入器和后端推广到更多架构、语言和目标。框架（验证器、预言机、搜索循环）并非特定于模型；Agent 提供了这种通用性，而拓宽它是工作的核心方向。

结果：int8 在整个推理集群上击败 cuBLAS

AMK 自动调优的 int8（W8A16，近无损）巨型内核在 NVIDIA 数据中心推理类 GPU 上的 batch-1 解码中击败了 CUDA-graphed cuBLAS bf16，这是由 AMK 自己的搜索自主发现并通过正确性门控（argmax 精确）的。比率 = cuBLAS / AMK，所以 > 1 意味着 AMK 更快。

| GPU | 类别 | int8 对比 cuBLAS bf16（最佳） | 结论 | |—|—|—|—:|—| | L4 (sm_89, 300 GB/s) | 推理 | 1.18× → 1.33× (1.3B→4B) | ✅ 胜出，随规模增长 | | L40S (sm_89, 864 GB/s) | 推理旗舰 | 1.25× → 1.27× (4B→6.7B) | ✅ 胜出 | | A10G (sm_86, 600 GB/s) | 推理 | 1.04× → 1.08× (≥3.5B) | ✅ 在规模上胜出 | | RTX 5090 (sm_120) | 消费级（本地） | 1.19× → 1.23× | ✅ 胜出 | | A100 (sm_80, 1.4 TB/s) | 训练 | 0.79× → 0.55× (1.3B→13B) | ✗ 落后（随规模下降） | | H100 (sm_90, 3.1 TB/s) | 训练 | 0.72× → 0.60× | ✗ 落后 |

推理/数据中心 GPU 在 Modal 上测量（可复现；文件存储在 paper/results/int8_scale_*.json 中）。RTX 5090 本地测量（int8_search_multisize.json；Modal 没有 RTX 5090 芯片，因此不在 Modal 扫描中）。

决定胜负的是推理类与训练类架构的差异，而非带宽排序： 864 GB/s 的 L40S 胜出幅度大于 600 GB/s 的 A10G。AMK 读取的重量字节数大约减半（int8），这个优势必须克服一个固定的 tile 间跨 SM 同步开销；更大、以 GEMV 为主的模型可以摊薄这个固定成本。训练类的 A100/H100 从未突破（比率随规模下降，这是同步不足的特征；cp.async int8-GEMV 探针甚至在 A100 上倒退到 0.82×，确认了跨 SM 同步而非加载延迟才是限制因素）。胜出场景是 batch-1、pos-0 / 低上下文。完整数据和分析：DATACENTER_RESULTS.md。

同等精度差距（明确说明）

在同类 bf16 路径上，AMK 落后于 cuBLAS。在一个 622.9 MB 的模型上，bf16 内核运行在 ~1.38 ms/token，比 CUDA-graphed cuBLAS 慢约 1.24 倍；优化的 bf16 GEMV 达到 ~451 GB/s = 理论值的 ~51% / 实测 HBM 峰值的 ~63%（cuBLAS 上限约 90%）。因此 int8 的胜出来自于流式传输更少的字节，而非更快的内核；同一模型上的 int8 内核约为 1.22 ms/token（比 cuBLAS bf16 慢约 1.09 倍）。剩下的杠杆是一个带宽饱和的 cp.async GEMV 配合更粗粒度的跨 SM 同步，这是内核质量的提升而非重新设计，而保障安全与自动化的正确性承载架构已经就位。我们并未在 bf16 batch-1 上击败 cuBLAS/vLLM，我们如实说明。

复现（GPU）：uv run pytest tests/test_cuda_perf.py；10 分钟自我改进运行：uv run python amk_cli.py autoresearch small --gpu rtx5090 --minutes 10。

已构建并验证的内容

• 整个前向传播作为一个协作内核启动运行。 持久巨型内核（每个 SM 一个线程块，计数器同步）执行完整的 Llama 风格解码，并在 fp32 / bf16 容差内匹配 eager PyTorch 以及 CPU 参考虚拟机，精度约为 1e-7。
• 自重新目标，已在三种架构上验证。 同一份源码在 sm_120 (RTX 5090)、sm_80 (A100) 和 sm_90 (H100) 上构建并运行了正确的巨型内核，nvcc 生成代码源自实时设备（H100 正确运行了一个 3 GB / 3202 任务的 Llama-1B 形解码）。参见 DATACENTER_RESULTS.md。
• 多 token 生成匹配 eager。 AMK 贪婪解码一个序列，在步骤间串起一个持久 KV 缓存；生成的 token ID 与 eager 贪婪解码相同。复现：uv run amk generate toy --gpu rtx5090 --prompt-ids "1,2,3" --max-tokens 32 --verify。
• 一个真实的已训练检查点，端到端。 AMK 导入 HuggingFaceTB/SmolLM2-135M（真实权重 + tokenizer）并逐 token 复现 HuggingFace 自己的贪婪 generate。复现：uv run python examples/run_hf_model.py（另见 tests/test_hf_checkpoint.py）。
• 一个静态检查的调度验证器。 在 7160 个对抗性调度上，验证器实现了零误接受；一个不安全的 Agent 提议调度会在验证时被 REJECTED，而不会挂起 GPU。
• 一个原生编码 Agent 框架。 可通过一个结构化的编辑界面由任何编码 Agent（Claude Code / Codex）驱动：MCP 服务器 + Claude Code 技能/命令/子 Agent/工作流 + Codex AGENTS.md。参见 docs/AGENT_HARNESS.md 和 HARNESS.md。
• 一次 10 分钟无人值守的自动研究运行 使巨型内核相对于其自身起始调度性能提升了 1.47 倍。
• 98 个测试通过（uv run pytest）：78 个在 CPU 上通过；20 个 CUDA 测试在无 GPU 时自动跳过。

为什么需要巨型内核

单流解码是带宽受限的：每个 token 必须将整个权重集通过 SM 流式传输一次。理论下限是 weights_bytes / HBM_bandwidth。正常的 PyTorch / cuBLAS 执行会每个算子启动一个内核，并在算子之间通过 HBM 来回传输激活值，每层付出启动延迟和内存气泡数十次。而巨型内核只启动一次，将持久线程块驻留在每个 SM 上，并原地遍历模型的依赖图：激活值存在于片上页面中，下一层的权重在当前层计算时被预取，算子之间没有内核启动气泡。胜出场景是单流/低 batch 解码延迟：语音、实时、Agent 循环。我们并未声称在高 batch 下击败吞吐优化的服务；那属于计算受限，并非本战场。

工作原理

生成过程被限制在一个经过验证的结构内，因此正确性是架构的属性，而非模型输出的属性。

四层

无死锁，由构造保证

前向传播是一个 DAG。生产者只增加计数器；消费者只等待静态已知的阈值；执行是一个具有单调计数器的拓扑遍历：没有锁，没有任意信号。虚拟机拒绝加载任何不是有效 DAG 的调度：无效调度会在验证时被干净地 REJECTED，而不是挂起 GPU。这正是使自动生成的调度可以安全无人值守运行的原因。

两个自动研究循环

• 循环 1，指令优化（这就是 AutoKernel）：编辑一个符合 ABI 的微内核，隔离的正确性然后延迟评估（~秒级），保留/回退。错误的指令会在其自己的单元测试中失败；没有持久内核，不会挂起。
• 循环 2，调度优化（新循环）：Agent 的编辑表面是调度 IR，一个结构化对象 {tiling, fusion_grouping, sm_assignment, pipelining_depth, page_allocation} 加上内核参数，而非巨型内核代码。冻结的虚拟机确定性地下放它；每个提议在启动前都经过静态 DAG 验证。完整约定见 HARNESS.md。

四个特性（产品规格）

1. 通用性：一条命令将模型编译成经过验证的巨型内核，无需为每个模型编写手工 CUDA（目前：HF Llama 系列；扩大覆盖范围是路线图）。
2. 自重新目标：当新芯片发布时，AMK 通过搜索 + 硬件上验证在数天内重新目标。已在 sm_120 / sm_80 / sm_90 上得到验证。这是护城河。
3. 标准 IR：AMK 拥有权威的巨型内核 IR：SM 级任务 DAG、指令 ABI、调度格式。参见 docs/IR_SPEC.md、schedule/ir.py、vm/abi.h。
4. 数据飞轮：每次运行记录 (model, gpu, schedule, instruction, measured result)；这个语料库训练一个学习先验，使未来的每次运行都从更高的起点开始。

原生编码 Agent 集成

AMK 将经过验证的循环基底原生暴露给编码 Agent，具有相同的行为和相同的诚实规则。单一指南是 docs/AGENT_HARNESS.md。

• MCP 服务器 (amk_mcp.py)：amk_doctor / amk_propose / amk_eval / amk_loop / amk_autoresearch / amk_orchestrate_*。启用：uv sync --extra agent；注册：通过 .mcp.json (Claude Code) 或 ~/.codex/config.toml (Codex)。
• Claude Code：megakernel-optimization 技能；/amk-optimize、/amk-autoresearch、/amk-compile 斜杠命令；amk-megakernel-optimizer 子 Agent；.claude/ 下的工作流和目标。
• Codex：AGENTS.md + 相同的 MCP 服务器。

快速开始

AMK 作为一个真正的包（hatchling）安装，并暴露一个真正的 amk 控制台命令。uv sync 配置完整环境（torch cu128、numpy、用于 HF 导入器的 transformers、用于 CUDA JIT 构建的 ninja、pytest）；uv 是推荐路径。 使用 pip：pip install "automegakernel[models,cuda]"，注意 cu128 torch 固定仅适用于 uv，因此 Blackwell/sm_120（或任何特定 CUDA 构建）上的 pip 用户应首先安装匹配的 torch wheel，例如 pip install torch --index-url https://download.pytorch.org/whl/cu128。

uv sync # 配置环境 + 安装 `amk` 命令（可编辑）
# --- 无需 GPU（可在全新仅 CPU 机器上运行）---
uv run pytest # 完整套件（98 个测试；78 个在 CPU 上，20 个 CUDA 无 GPU 时自动跳过）
amk doctor # 环境 + GPU + nvcc + 注册目标
amk eval toy --device cpu # 一个结构化的正确性判断（基于 CPU 参考 VM）
# --- 需要 CUDA GPU + nvcc ---
amk compile toy --gpu rtx5090 --regime single-stream # 模型 -> 验证的巨型内核 + 报告
amk generate toy --gpu rtx5090 --prompt-ids "1,2,3" --max-tokens 32 --verify # 多 token 解码 == eager
amk eval toy --gpu rtx5090 # 正确性 + 测量的 GPU 延迟

每个子命令也可以通过 uv run python amk_cli.py ... 运行（行为相同）。一个真实的已训练检查点通过 uv run python examples/run_hf_model.py 端到端运行。

诚实规则（由框架强制执行，而不仅仅是声明）

• 决不以缺少配对正确性结果的方式报告延迟数字。 eval/bench.py 拒绝在没有从 eval/oracle.py 得到结论的情况下输出延迟。
• 正确性 = 全模型 logit 等价性在容差内 加上 与 eager PyTorch 在一个序列上的生成 token 一致性。
• 始终报告相对于 weights / HBM_bandwidth 屋顶线的距离（eval/roofline.py）。
• 测量的数字是在飞轮语料库 / results.tsv 中命名的硬件上产生的；DATACENTER_RESULTS.md 中的数据中心数字确实是测量的。我们不转录未测量的数字。
• 我们在 bf16 路径上还远未达到带宽受限，我们如实说明。诚实的当前声明是：int8 推理集群 cuBLAS 胜出、通用性、自重新目标、信任（无死锁且无竞态，由构造保证）、以及诚实的到屋顶线距离。

仓库布局

vm/                   第 0 层，可信的巨型内核虚拟机（CUDA）+ CPU 参考模拟器 + 验证
instructions/         第 1 层，符合 ABI 的微内核（Triton + CUDA）+ 生成器 + 验证
schedule/             第 2 层，图导入、降级、标准 IR、成本模型、搜索
dynamism/             第 3 层，连续批处理、动态形状、MoE（路线图占位）
eval/                 预言机（logit 等价性）· 基准（延迟）· 基线 · 屋顶线
amk_cli.py            `amk` 控制台命令（doctor/compile/generate/eval/propose/loop/...）
compile.py            产品：amk compile --gpu
generate.py           自回归多 token 解码（跨步骤串起 KV 缓存）
harness.py            编码 Agent 集成界面（循环 2，调度搜索）
examples/             run_hf_model.py，一个真实的 HF 检查点端到端（SmolLM2-135M）
docs/                 IR_SPEC.md（标准 IR）· AGENT_HARNESS.md（Agent 集成）HARNESS.md  编码 Agent 框架约定
DATACENTER_RESULTS.md 测量的推理集群 + sm_80/sm_90 自重新目标结果
program.md            自主操作大脑（让 AMK 无人值守运行）
models/               自包含测试模型（小稠密 -> MoE）

状态

正确性承载核心、GPU 巨型内核、自重新目标、多 token 生成、真实检查点路径以及 Agent 框架都已构建并验证。当前的重点是内核质量性能向屋顶线推进（参见差距）。参见 program.md 了解路线图和自主循环纪律；IR（[schedule/ir.py](schedule/ir.py