Part 6: GPU 加速

Metal、CUDA、Flash Attention

1. GPU 后端 + 总结

CPU decode 速度约 5-10 t/s，对话体验卡顿。GPU 后端将矩阵运算卸载到 Metal/CUDA/ROCm，decode 速度提升 5-10 倍。今天学习 ObjC 互操作、Metal 计算管线和 Flash Attention，最后用一个完整的请求生命周期串联 15 天的知识。

GPU 后端将 CPU 注意力卸载到 Metal/CUDA/ROCm，大幅提升推理速度。

C 知识点

1. Objective-C 互操作

ds4_metal.m 用 Objective-C 调用 Metal API，ds4_cuda.cu 用 CUDA C++ 调用 NVIDIA API。ds4.c 是纯 C，通过 ds4_gpu.h 统一抽象层与 GPU 后端交互：

// ds4_gpu.h — GPU 后端无关的抽象接口
// 所有函数签名使用 ds4_gpu_tensor * 不透明指针
// 不暴露 Metal (MTLBuffer) 或 CUDA (CUdeviceptr) 类型

typedef struct ds4_gpu_tensor ds4_gpu_tensor;

// 张量生命周期
ds4_gpu_tensor *ds4_gpu_tensor_alloc(ds4_gpu_graph *g, const char *name, uint64_t nbytes);
void ds4_gpu_tensor_free(ds4_gpu_tensor *t);
void ds4_gpu_tensor_read(ds4_gpu_tensor *t, void *dst, uint64_t offset, uint64_t len);
void ds4_gpu_tensor_write(ds4_gpu_tensor *t, const void *src, uint64_t offset, uint64_t len);

// 命令队列
void ds4_gpu_begin(ds4_gpu_graph *g);
void ds4_gpu_flush(ds4_gpu_graph *g);
void ds4_gpu_end(ds4_gpu_graph *g);
void ds4_gpu_synchronize(ds4_gpu_graph *g);

// ~60 个推理内核: embedding, RoPE, attention, MoE, steering 等
void ds4_gpu_embedding(ds4_gpu_graph *g, ...);
void ds4_gpu_flash_attn(ds4_gpu_graph *g, ...);
// ...

后端枚举 (ds4.h):

typedef enum {
    DS4_BACKEND_METAL,
    DS4_BACKEND_CUDA,
    DS4_BACKEND_CPU,
} ds4_backend;

编译选择:

macOS: 默认 Metal 后端 (CORE_OBJS = ds4.o ds4_metal.o)
Linux: 默认 CUDA 后端 (CORE_OBJS = ds4.o ds4_cuda.o，用 nvcc 链接)
make cpu: 两平台通用，-DDS4_NO_GPU 编译纯 CPU 版本
CUDA 构建目标:
- make cuda-spark: DGX Spark / GB10 专用（无需指定 arch）
- make cuda-generic: 通用 CUDA GPU
- make cuda CUDA_ARCH=sm_N: 显式指定架构（如 sm_80）
ROCm 构建目标（AMD Strix Halo / gfx1151）:
- make strix-halo（别名 make rocm）：用 hipcc 链接 ds4_rocm.o + rocm/*.cuh，定义 -DDS4_ROCM_BUILD，默认 ROCM_ARCH=gfx1151
- ROCm 后端现已并入 main 分支（此前长期只在独立的 rocm 分支，由社区因 antirez 无对应硬件而维护）
质量评分工具 gguf-tools quality-score 也支持跨平台：macOS 用 $(CC) 链接 Metal，Linux 用 $(NVCC) 链接 -lcudart -lcublas，通过 #ifdef __APPLE__ 编译时选择后端

关键设计：

ds4_gpu.h 对 C 完全隐藏具体 GPU API 类型（Metal/CUDA/HIP）
新增 GPU 后端只需实现 ds4_gpu.h 中的 ~60 个函数，不改动 ds4.c
CPU 路径通过 DS4_NO_GPU 宏完全绕过 GPU 代码
ROCm 后端靠编译期 -DDS4_ROCM_BUILD 选择（而非运行期 enum 值）：它在 ds4_gpu.h 里用 #ifdef DS4_ROCM_BUILD 暴露若干独有钩子（layer-major 批量加载、q8_f16 缓存释放等），其它后端提供 no-op stub 保持链接。ds4_backend 枚举仍是 METAL/CUDA/CPU——ROCm 复用了非 Apple（Linux）的代码路径，只是把实现文件从 ds4_cuda.cu 换成 ds4_rocm.cu。

2. Metal 计算管线

Metal GPU 执行模型推理的核心流程：

ds4_metal.m 中的流程：

objc

id<MTLDevice> device = MTLCreateSystemDefaultDevice();
// 启动时打印设备名称和内存
// ds4: Metal device Apple M2 Ultra, 192.00 GiB RAM
id<MTLCommandQueue> queue = [device newCommandQueue];

// 编译 .metal shader
id<MTLLibrary> lib = [device newLibraryWithSource:shader_source ...];
id<MTLFunction> fn = [lib newFunctionWithName:@"kernel_name"];
id<MTLComputePipelineState> pipeline = [device newComputePipelineStateWithFunction:fn ...];

// 执行
id<MTLCommandBuffer> cmd = [queue commandBuffer];
id<MTLComputeCommandEncoder> enc = [cmd computeCommandEncoder];
[enc setComputePipelineState:pipeline];
[enc setBuffer:input offset:0 atIndex:0];
[enc setBuffer:output offset:0 atIndex:1];
[enc dispatchThreadgroups:... threadsPerThreadgroup:...];
[enc endEncoding];
[cmd commit];
[cmd waitUntilCompleted];

可选 buffer 绑定占位

encoder 绑定 buffer 时，可选参数即使不使用也必须绑定占位值，否则 Metal debug layer 会报校验错误：

objc

// ds4_metal.m — router finalize
float zero_f32 = 0;
if (has_bias)
    [enc setBuffer:bias_buf offset:0 atIndex:3];
else
    [enc setBytes:&zero_f32 length:4 atIndex:3];  // 零值占位

if (hash_mode)
    [enc setBuffer:hash_buf offset:0 atIndex:5];
else
    [enc setBytes:&zero_f32 length:4 atIndex:5];  // 零值占位

Metal debug layer 要求 shader 声明的每个 [[buffer(N)]] 都有对应绑定，不用的可选参数用 setBytes:&zero 填充即可。

张量句柄的双重释放守卫

C 代码把 ds4_gpu_tensor 句柄当作 retained Objective-C 对象持有（__bridge_retained 转出）。ds4_gpu_tensor_free 用 __bridge_transfer 把所有权还回 ARC 释放。如果同一个不透明句柄被释放两次，第二次 __bridge_transfer 会把一个已经释放的对象再交给 ARC——macOS 会报 malloc corruption 崩溃（issue #404）。

修复是在 ds4_metal.m 里维护一张活跃张量句柄集合（一张真正的开放寻址哈希表），在触碰任何 Objective-C 状态前先校验：

static pthread_mutex_t g_tensor_mu = PTHREAD_MUTEX_INITIALIZER;
static uintptr_t *g_tensor_live_slots;   // 开放寻址槽
static size_t g_tensor_live_cap, g_tensor_live_count, g_tensor_live_tombs;

// 释放前先从集合里删除；不在集合里 → 双重释放，忽略
static int ds4_gpu_tensor_prepare_free(ds4_gpu_tensor *tensor, ...) {
    pthread_mutex_lock(&g_tensor_mu);
    if (!ds4_gpu_tensor_live_remove_locked(tensor)) {   // 不在集合里
        pthread_mutex_unlock(&g_tensor_mu);
        fprintf(stderr, "ds4: Metal tensor free ignored for unknown handle %p\n", tensor);
        return 0;                                       // 直接拒绝，不碰 ObjC
    }
    ...
}

这张表用线性探测开放寻址、删除时写**墓碑（tombstone，UINTPTR_MAX）**而非清零（清零会断开探测链），负载因子到 70% 时翻倍 rehash。alloc/view 时插入，free 时先删再交给 ARC。诊断用的 g_tensor_alloc_live_bytes/peak_bytes 计数器也走同一把 g_tensor_mu，避免内存报告读到半更新的值。

为什么用哈希表而不是引用计数：句柄的所有权模型已经是"一个 C 句柄 = 一次 ObjC 引用"，没有共享持有者；真正的风险是重复释放而非**过早释放"。集合成员资格（"这个句柄还活着吗"）正是 O(1) 哈希查找能廉价回答的问题，墓碑机制又让删除不影响后续探测——这是一个教科书级的开放寻址 + 墓碑实战用例。

3. Metal Shading Language

.metal 文件用 MSL 编写 GPU kernel：

metal

// metal/softmax.metal (简化)
kernel void softmax(
    device float* input  [[buffer(0)]],
    device float* output [[buffer(1)]],
    uint gid [[thread_position_in_grid]])
{
    // 每个 GPU 线程处理一个元素
    output[gid] = exp(input[gid]) / sum;
}

ds4 项目有 19 个 .metal 文件，覆盖：

dense: 矩阵乘法
flash_attn: Flash Attention
moe: MoE 专家计算
norm: RMSNorm
softmax: Softmax
dsv4_kv/dsv4_rope: KV cache 和 RoPE
argsort: 排序（用于 top-k）
等等

Metal Neural Acceleration (NAX)

M5 芯片引入了 MPP/Neural Accelerator 硬件（NAX），ds4 在检测到 M5 时自动利用两个 NAX kernel：

Indexer Scores NAX — 压缩 KV 注意力的索引评分，将评分计算卸载到 Neural Accelerator
Tensor Matmul NAX — MoE 专家的 Q8_0 矩阵乘法，支持 32/64/128 tile sizes

检测方式：g_metal4_m5_neural_accelerators_hint（从系统属性读取的 M5 标志）。当 n_tokens >= 16 时启用 NAX 路径，低于此阈值时退回标准 GPU 计算。

Compressed KV in F16

Metal 后端新增 comp_kv_f16 模式——将压缩 KV cache 以半精度（float16）而非 float32 存储。每行 KV 内存减半（从 DS4_N_HEAD_DIM × 4B 降为 DS4_N_HEAD_DIM × 2B），在大上下文场景下节省大量内存。对应的 Metal kernel 位于 metal/dsv4_kv.metal，在读入时将 F16 解压缩为 F32 用于注意力计算。

GPU 功耗节流

GPU 后端支持 power_percent（1-100）参数控制推理速度。实现方式是在每个 prefill 层和 decode token 后调用 graph_power_sleep()：

sleep_us = (100 - power) / power × elapsed_us

agent 模式下通过 ds4_engine_set_power() 动态调整——等待用户确认时降低功耗，活跃推理时恢复满速。

Typed Pointer vs void Pointer

Metal shader 参数应使用 device const char * 而非 device const void *，让 Metal 反映正确的只读访问语义：

metal

// 改前：void * — 不表达访问语义，debug layer 可能报校验错误
kernel void get_rows(device const void* x [[buffer(0)]], ...)

// 改后：char * — 明确只读，满足 Metal debug layer 校验
kernel void get_rows(device const char* x [[buffer(0)]], ...)

get_rows.metal 和 set_rows.metal 等文件中的参数已从 void * 改为 char *。void * 在 Metal 中不传达内存访问意图，char * 配合 const 明确表达只读语义。

LLM 知识点

1. GPU 加速原理

CPU vs GPU 的根本区别：

CPU: 4-12 个强核心，每个核心处理复杂任务
GPU: 数千个弱核心，每个核心做简单计算但并行执行

矩阵乘法 A[n,m] × B[m,k]:
CPU: 逐元素计算，串行（或少量并行）
GPU: 每个元素一个线程，数千元素同时计算

为什么 LLM 推理需要 GPU？

矩阵乘法是高度并行的（每个输出元素独立）
GPU 的内存带宽远高于 CPU
Metal GPU 可以直接引用 mmap 内存（零拷贝）

2. Flash Attention

标准注意力：先算完整的 Q·K^T 矩阵，再 softmax。内存 O(n²)。

Flash Attention：分块计算，每次只加载一小块 Q 和 K 到高速缓存，在片上完成 softmax。内存 O(n)。

ds4 的 metal/flash_attn.metal 实现了这个优化。

3. 零拷贝 MTLBuffer

// Metal 可以直接引用 mmap 的内存，不需要拷贝到 GPU 专用内存
id<MTLBuffer> buf = [device newBufferWithBytesNoCopy:
    (void*)(mmap_base + tensor_offset)
    length:tensor_bytes
    options:MTLResourceStorageModeShared
    deallocator:nil];

这就是为什么模型加载用 MAP_SHARED——Metal GPU 直接从文件映射读取权重，不需要额外的 GPU 内存拷贝。81GB 的模型只需要 81GB 的磁盘空间和 mmap 映射，不需要额外的 GPU 显存。

模型视图与张量覆盖不变量

当模型文件超过 Metal 设备的 maxBufferLength 时，引擎创建多个重叠的 MTLBuffer 视图来覆盖整个文件。相邻视图必须有足够大的重叠区域，确保每个张量至少完整包含在一个视图中：

// 最大单张量大小限制，决定视图重叠区域
#define DS4_METAL_MODEL_MAX_TENSOR_BYTES (2ull * 1024ull * 1024ull * 1024ull) // 2 GiB

// 重叠 = round_up(MAX_TENSOR_BYTES, page) + page
// 视图步进 = max_buffer - overlap

Q4_K 量化模型的 MoE 专家张量约 1.125 GiB，之前的 672 MiB 限制不足以覆盖。2 GiB 阈值确保所有量化格式（包括 Q4 专家张量）都能完整映射到至少一个视图中。

4. 分布式推理（Distributed Inference）

PRO 模型（~430GB IQ2_XXS 或 ~838GB Q4 分片）太大，单机无法运行。ds4 新增 coordinator/worker 分布式推理架构，将模型层切片分配到多台机器上执行：

核心设计：

组件	职责
Coordinator	接收请求、执行前半层、转发激活、接收 logits
Worker	注册层范围、执行后半层、返回结果
传输协议	自定义 TCP 二进制协议（magic `0x44533444` = "DS4D"）
KV 快照	拓扑无关：保存时聚合所有 worker 的层张量，加载时按当前路由分发

Worker 通过 HELLO 消息注册（包含 layer_start、layer_end、model_id、quant_bits），Coordinator 据此构建路由表。激活传输支持可配置的 activation_bits（默认 32-bit float，可压缩为更低精度减少网络开销）。

bash

# Coordinator（加载前半层）
./ds4 --role coordinator -m pro-q4-layers00-30.gguf -p "Hello"

# Worker（加载后半层）
./ds4 --role worker -m pro-q4-layers31-output.gguf

模型分片加载（Model Span Loading）：分布式引擎使用 ds4_gpu_set_model_map_spans() API，只加载分配给本机的层范围，而非整个模型。这通过 offset/size 对描述需要映射的字节范围，配合 chunked preloading 异步拷贝到设备内存。

// 只加载本 worker 负责的层切片
ds4_gpu_set_model_map_spans(
    model_map, model_size,
    offsets, sizes, span_count,
    max_tensor_bytes
);

详见分布式推理。

4a. SSD Streaming（SSD 流式推理）

除了分布式推理，ds4 还提供另一种突破内存限制的方式：SSD Streaming。核心洞察是 MoE 层有 256 个路由专家，但每个 token 只激活 6 个——路由专家占模型大部分空间，但大部分时间都在"睡觉"。

缓存计划：ds4_ssd_cache_plan 根据可用 RAM 预算自动计算能缓存多少专家。自动预算取推荐工作集的 80%、扣除非路由权重，剩下的用于路由专家；显式 --ssd-streaming-cache-experts 32GB 则直接给字节预算（换算成完整专家数），过大时推理前自动封顶以保持可锁定（lockable）。详见 SSD Streaming 术语表。

热度排名（启动）：ds4_streaming_hotlist.inc（13,334 行）按流行度预排列所有 (layer, expert_id) 对。启动时预加载最热门的专家，自动预热默认封顶 4096 个。

运行期淘汰（路由热度）：决定"淘汰谁"的是一张运行期热度表 route_hotness[layer][expert]，每处理 16 个 decode token 整表右移衰减一位（>>= 1），淘汰时驱逐热度最低者。这张表是提示局部的——新 prompt 开始时清零（ds4_gpu_stream_expert_cache_reset_route_hotness()），但常驻缓存本身跨会话保持热：易变的启发式状态要重置，昂贵的物化缓存要复用。

三路 Prefill：默认 prefill 调度器按 prompt 长度分流——极短走逐 token decode 式、中等走 selected-expert 批量、长 prompt 走全层 layer-major。切分点按量化格式调优，避免短 prompt 落到"加载一堆用不上的专家"的慢路径。

混合精度路由专家：当 GGUF 跨层不统一量化（per-layer boosted quants，例如全局 IQ2_XXS 但有几层升到 Q4_K）时，提升层的 ffn_exps 张量被纳入 decode-span 构建器走 mapped-view 兜底，且 slab 分配器在启动时预置尺寸类、对超大尺寸 freeze+reject 而非 last-writer-wins，避免 slab 被毒化。

多后端实现：SSD streaming 不再是 Metal 独享——CUDA（ds4_cuda.cu，有界缓存 + 进度上报）和 ROCm（ds4_rocm.cu，selected-expert 缓存 + 重叠读取 + 全层 streaming prefill）都已实现。streaming API 已重构为传一张 ds4_gpu_stream_expert_table 表：

typedef struct ds4_gpu_stream_expert_table {
    const void *model_map;   uint64_t model_size;
    uint32_t layer, n_total_expert;
    uint64_t gate_offset, up_offset, down_offset;
    uint64_t gate_expert_bytes, down_expert_bytes;
} ds4_gpu_stream_expert_table;

// 异步加载缺失专家（重构后参数更简洁）
int ds4_gpu_stream_expert_cache_begin_selected_load(
        const ds4_gpu_stream_expert_table *table,
        const int32_t *selected_ids, uint32_t n_selected);
uint32_t ds4_gpu_stream_expert_cache_current_count(void);

mlock pinning（让缓存专家驻留 RAM）

SSD streaming 模式下模型 >> RAM，缓存的专家随时可能被 OS 在内存压力下换出——换出后每个 token 都得从磁盘重读，Flash decode 退化约 2-3×、首 token ~6s。mlock 把已缓存的专家钉在物理 RAM（不可换出）。GLM 5.2 重构时曾把这块 pinning 整段删掉（mlock 调用、slab 锁、预算上限全归零，#532），现已适配重构后的 slab 分配器移植回来：整缓冲 mlock、per-slot fill 锁 / relief 解锁、routed-memory 预算上限优雅降级、失败时 cap-on-failure、压力下解钉最冷的 ~10%。M5 Max 128GB 实测从 0.3-2.5 → ~7.0 gen t/s，反超之前基线。

Metal 4 prefill 加速

Metal 4 后端新增一批 prefill 内核加速：路由 MoE prefill（532ec8b）、indexed prefill（222b2cb）、Q4 attention 输出 prefill（96c3ba4）、partial Q4 prefill tiles（d69a017）；以及让 streamed prefill maps 与专家内核一致（短 SSD prefill 不读活跃 Metal view 外的专家张量，d14ce35）、从映射 prefill 层播种专家缓存（4893e0c）。

GLM 全量检查点 streaming 成熟化（2026-09）

177.77 GiB 的 GLM 5.3 Q4_K（或 196.58 GiB IQ2_XXS）在单台 128 GB Mac 上 streaming 到"可用"花了这一周的打磨，主题是让缓存优先级可被运行期夺回：

有界的内置热度加成（b6af0ad）：metal_graph_streaming_builtin_hotness 给 hotlist 预加载专家的"新鲜度加成"原来无界——decode 的按需加载永远挤不掉种子。加成封顶后，冷种子可以被热需求替换，缓存才真正"活"起来。
从第一个路由层开始老化（715ed7a）+ 超越内置排名的自适应（91dda0b / 85837d0）：老化不再等 dense 前缀结束；运行期观测到的路由频率可以推翻 hotlist 的静态排名，把更多安全工作集让给实际被选中的专家。
前瞻一层读（437dbcb）：续 prefill chunk ≥256 token 时提前读下一层专家——8K/4K append 从 63.23/61.56 提到 90.70/84.31 t/s；小 prefill（<256 token）则复用 IQ2 选中专家缓存读（d7ade9f）。
首段 prefill 与可用缓存重叠（6ccc1e7）、并行 prefill 读不再被串行预读卡住（d6580a5）、路由器回读复用（edcafe4）。
内存视图卫生：流式模型视图有界化但保留视觉/草稿模型的辅助映射（660e1d4）；静态权重在运行时预算内保持常驻（22e4172，DS4_METAL_DISABLE_STREAMING_STATIC_LOCK）；注意力 scratch 对任意上下文长度对齐（fb2abb2）。
正确性：完整 GLM 紧凑 prefill 现在保留 RoPE（e965352，回归 tests/test_glm_rope_prefill.c）；专家缓存回归 make test-metal-ssd-experts / test-ssd-cache。

实测（上游 docs/SSD_STREAMING.md）：GLM Q4_K 初始/续 prefill 121/104 t/s、生成 11.9/14.9 t/s；完整 IQ2_XXS 的 16-token append 从 30.8 s 降到 2.9 s（61.35 GiB 专家缓存），生成 4.09→5.11 t/s。已知未解回归：完整 GLM 走 server（多会话）时仍会失速（31d1992 记录在案）。

命令行选项：

选项	说明
`--ssd-streaming`	启用 SSD streaming 模式
`--ssd-streaming-cache-experts N`	缓存字节预算（过大时推理前封顶）
`--ssd-streaming-preload-experts N`	显式指定启动预加载专家数（默认自动封顶 4096）
`--ssd-streaming-cold`	冷启动（不预加载 hotlist；仅用于测量）

与分布式推理的区别：SSD streaming 是单机方案，适合内存不足但有大容量 SSD 的场景；分布式推理是多机方案，适合 PRO 模型等超出单机存储极限的情况。两者可以结合使用——但组合时 streaming 的流式映射要从"每步设一次"降级为"每层设一次"（分布式把层栈拆成切片，而映射是 per-layer 的），详见 SSD Streaming §分布式层切片下的流式映射。

4b. DGX Spark / GB10 后端（HBM 缓存）

NVIDIA DGX Spark（GB10 芯片，121 GiB UMA）有一个特殊的 CUDA 推理挑战：ATS（Address Translation Service）允许 GPU 直接消费 host mmap 权重，但有效带宽远低于 HBM 驻留数据。

ds4 的解决方案：启动时 HBM 缓存——将"热"张量（注意力投影、MoE 共享专家、输出投影）拷贝到 GPU HBM，冷 MoE 路由专家保持 ATS 映射：

默认缓存预算 96 GiB（cuda_model_cache_limit_bytes()），通过 DS4_CUDA_WEIGHT_CACHE_LIMIT_GB 环境变量可覆盖。IQ2 模型（~81GB）和混合 Q2/Q4 模型（~91GB）可完整缓存；完整 Q4 模型（~153GB）超出预算，需使用分布式层加载。

HBM 缓存查找优先于 UVA 映射指针——直接 HBM 读取比通过 host 页表映射快约 10%。cuda_model_range_ptr() 使用 hash-keyed 查找实现单次读取命中。

效果（DGX Spark / GB10，ds4flash，n=256）：

优化前：~13.9 t/s（纯 ATS 映射）
优化后：~16.13 t/s（HBM 缓存热张量 + 小 batch kernel 融合）

编译目标：make cuda-spark（自动检测 GB10 架构，无需指定 CUDA_ARCH）。

Metal View Cap 修复

分布式推理的 span 映射使用 ds4_gpu_add_model_view_range() 创建 Metal 视图。之前的实现给所有大映射都应用 128 GiB 视图上限，包括普通的单次完整模型映射，导致 Metal VM 验证变慢。

修复：新增 use_default_view_cap 参数——普通完整模型映射（ds4_gpu_map_model_views()）不设上限，分布式 span 映射（ds4_gpu_set_model_map_spans()）保留 128 GiB 分割。这样普通映射保持单次映射的高效性，只有需要切分的分布式场景才做视图分割。

模型缓存 FD 绑定

CUDA 后端使用直接文件 I/O（非 mmap）加载路由专家权重，需要正确的文件描述符和 mmap 基地址来计算文件偏移。当 MTP 模型在主模型 fd 绑定和缓存准备之间加载时，fd 可能被错误覆盖。

修复：新增 ds4_gpu_set_model_fd_for_map(int fd, const void *model_map) 接受显式的 fd 和 model_map 指针，替换隐式使用 g_model_host_base 的旧 ds4_gpu_set_model_fd()。MTP 模型设置完成后，重新绑定主模型的 fd：

// 缓存前绑定正确的 fd
ds4_gpu_set_model_fd_for_map(e->model.fd, e->model.map);
accelerator_cache_model_tensors(...);
// MTP 设置后重新绑定主模型 fd
ds4_gpu_set_model_fd_for_map(e->model.fd, e->model.map);

2026-09：常驻会话 + 小 batch 对齐内核

常驻 Spark 会话复用种子验证器（df80a87），小 batch 换对齐的 Q8 向量内核与专家向量内核（2803820 / 45f54e9）。DSpark 在常驻 Q2 上的三连跳（aa1e71c）：C hash table temp 0 从 19.72 → 23.44 → 31.41 t/s；而散文 prompt 停在 ~18.8 t/s——推测解码只对可预测文本回本，这是同一个引擎上"加速取决于文本分布"的最干净对照。

2026-09 中旬：V4.1 与 Qwen3.8 登陆 Spark

V4.1 Flash 文本（a04f46f，见 §4r）：Q2 SSD streaming 单机即可（341 GiB 文件、189 GiB 盘上 Engram），或双 Spark RDMA TP 常驻（每台 ~81 GiB）。自动缓存调优后 3,241-token /read README.md prefill 93–96 t/s（此前 49–54）；64 GiB 缓存预算 + 64K ctx 下 32K 初始 prefill 384 t/s、3.2K 追加 88 t/s；自动缓存的 256-token 回复 decode 慢约 5%，但 prefill+回复总时长从 ~87s 降到 59s——prefill 吞吐换 decode 微降的净赚交易。双 Spark RoCE：单会话 21.9 t/s，8 个就绪会话聚合 ~28 t/s（5 个起批处理），长跑 32K prefill ~400 t/s、8K 追加 ~290 t/s。Q2 SSD serving 8 行批处理 11.0 vs 8.4 聚合 t/s。

Qwen3.8 Flash Next：Q2（41.73 GiB 常驻）/ Q4（69.74 GiB）单 Spark 都装得下，95.37 GiB n-gram 表留 SSD。Q2：16K prefill 250 t/s、扩到 30K 243 t/s、decode ~17.3 t/s；256-token 散文 17.1（普通）→ 19.9 t/s（MTP，temp 0）。视觉、agent、服务器 API、文本会话检查点在 CUDA 全可用（make cuda-spark 构建，无 CUDA 专属转换或特性开关；Spark 上 --prefill-chunk 不设即用更快的 8192 默认）。

4c. Strix Halo / ROCm 后端

AMD Strix Halo（如 Framework Desktop，Radeon 8060S，gfx1151，128 GB 统一内存）是 ds4 的第三个 GPU 后端。ROCm 后端现已并入 main 分支（此前长期只在独立的 rocm 分支，由社区维护，因为 antirez 没有对应硬件）。

构建：make strix-halo（别名 make rocm），用 hipcc 编译 ds4_rocm.o + rocm/*.cuh（约 18K 行新代码），定义 -DDS4_ROCM_BUILD。后端基于 rocWMMA（AMD 的矩阵乘加速库），需要 libhipblas/libhipblaslt/librocblas/librocwmma 等。

关键工程难点（详见上游 STRIXHALO.md）：

ROCm 安装补全：Ubuntu 26.04 的 librocwmma-dev 缺 rocwmma/internal/ 头文件，需手动 clone rocWMMA 7.1.0 补齐；若 ROCm 装在 /usr 而工具链期望 /opt/rocm，要建符号链接。
显存扩展（GT TTM aperture）：128 GB Strix Halo 系统默认只暴露约 62 GB GPU 可见内存，不够 80.76 GiB 模型 + 运行时缓冲。需设内核参数扩大 GTT：
```
amd_iommu=off amdgpu.gttsize=126976 ttm.pages_limit=32505856 ttm.page_pool_size=32505856
```
重启后 rocminfo 应报 gfx1151 pool: 130023424 KB（约 124 GB）。
GGUF 选择：在 Strix Halo 上用标准 IQ2XXS/Q2K/Q8 imatrix GGUF；避免 mixed IQ2/IQ4 或 IQ2/Q4 GGUF——它们给 ROCm 路径更大的内存压力，可能触发系统 OOM 而非干净的 ds4 失败。

ROCm 特定修复（并入本次同步）：

Fix ROCm distributed slice mapping — 分布式层切片映射
Fix ROCm MTP model residency — MTP 模型驻留
Fix ROCm mixed streaming expert fallback — 混合精度 streaming 专家兜底

2026-08 DeepSeek ROCm 稳定性：IQ2 SSD prefill 死锁（2f49b27）、Q4 SSD 专家 staging（9e1988b）、大模型 arena 收紧（6882a0f，避免 HSA 分配空间耗尽）、prequant decode 内核恢复（1e8f16c）、两 token MTP 验证优化（d250a7c）。

DSpark 登陆 ROCm（84cc882）：DSpark 推测解码此前缺 ROCm 侧的草稿/验证内核。新增三个头文件补齐：rocm/ds4_rocm_attention.cuh 的 attention_noncausal_raw_batch_heads_kernel（草稿块用的非因果批量注意力——每个 query 行看完整 raw-KV 环而非因果前缀）、rocm/ds4_rocm_indexer.cuh 的 dspark_markov_argmax_kernel、以及 rocm/ds4_rocm_attention_launch.cuh 把它们接进 ds4_gpu.h 的注意力 API。至此 DSpark 在 Metal / CUDA / ROCm 三个后端都可用。

更安全的 ROCm prefill 默认（#387）：ROCm 构建曾对 >4096 token 的长 Flash prompt 默认 8192-token prefill 工作区（ds4_prefill_cap_for_prompt 里的 #ifdef DS4_ROCM_BUILD 分支）。在 Strix Halo 上，IQ2 全权重缓存 + 8192 工作区会在生成开始前耗尽 HSA 分配空间。现在移除 ROCm 专用的 8192 默认，让它与普通 Flash 一样用 4096（PRO 仍保留 8192），--prefill-chunk 显式覆盖不受影响——是"留够内存余量"对"更大单块吞吐"的取舍。

MTP 启动驻留（#425）：MTP 载入第二个 GGUF 映射后，cuda_model_range_ptr 之前对非主映射一律走 cuda_model_range_copy_uncached（无界拷贝路径），把主模型也逼进无界拷贝，启动即失败。修复是调整查找顺序：先查已准备好的 range 表、再查 fd 支持的缓存拷贝，只有都落空才回退到无界拷贝——让主模型 GGUF 和 MTP GGUF 共存而不走慢路径。同时，注册第二个模型映射时禁用可选的扩展 Q8→F16 缓存（g_q8_f16_disabled_for_multi_model）：该缓存只是加速路径，但在 UMA 系统上会吃掉 MTP session/context 张量所需的内存余量；普通 Q8 kernel 照常可用。新增 DS4_CUDA_NO_Q8_F16_CACHE 环境变量可手动关闭该缓存。

可选后端钩子（Fused GPU Ops as Backend Hooks）：融合算子（如 fused RMSNorm/RoPE、fused MoE）从 ds4.c 里写死的 Metal/CUDA 分支重构为 ds4_gpu.h 上的可选后端钩子。每个后端按需实现（Metal 131 行、CUDA 94 行新增），ds4.c 侧减为统一的钩子分发（-40 行）。这让 ROCm 这类新后端能选择性提供融合内核而不必复制所有分支逻辑。

ROCm MXFP4 路由专家内核（PR #656，2026-09 合入）：MXFP4 终于有了自己的 ROCm 原生路径（rocm/ds4_rocm_moe.cuh / ds4_rocm_moe_launch.cuh，mxfp4_path = gate/down 都是 GGUF 类型 39）。这是一套值得整段读的内核设计课：

解码布局：17 字节块 = 16 字节 nibble（前 16 个值在低半、后 16 个在高半——不交错）+ 1 字节 UE8M0 指数。dev_e8m0_to_f32() 一行完成幂转换（e==0 映射到位型 0x00400000 的次正规数 2⁻¹²⁷，数值上视作零）；dev_mxfp4_unpack4() 在 AMDGCN 上用 __builtin_amdgcn_perm（V_PERM_B32）把码表两半放 SGPR 立即数、一条指令展开 4 个码字——没有发散的 LUT 读
dot-product 辅助函数族（核心抽象）：从 dev_dot_mxfp4_q8_K_block（结构体布局）到 _half_block（一对 lane 分摊一个 32 值块，wave 就能按 16 个连续 17 字节块合并访存）、_q8a_block8（对齐暂存的激活 + lo/hi 整数累加器分离——重结合但逐位一致）、_q8a_pair_block8（融合 gate+up：一次激活装载喂两张权重矩阵，砍掉最大的 LDS 流量项）。激活在 GPU 上量化成 Q8_K（q8_K_quantize_kernel），一个 Q8_K 块跨 8 个 MXFP4 块
寄存器/LDS 预算（f52dc1c）：把 per-tile pair 设置从运行时计数循环改成固定 8 次全展开 + 统一谓词，pair[]/xqb[] 标量进寄存器；静态 [8][16] Q8_K 暂存改 launch 大小的动态共享内存。gfx1151 代码生成从 99 VGPR / 48 B scratch / 12 waves → 83 VGPR / 无 scratch / 16 waves
向量化 + 融合（bfd0682）：内核原在 ~3% dp4a 峰值、发射受限——激活改 16 字节对齐量化切片（scale 单独成数组）、权重每块一条（错位的）b128 装载。gate/up tile 106.0→50.1 ms（2.11×）、down 69.6→34.0 ms
r9 L2 友好网格序（df1a49a）：profiling 显示 tile 内核被专家权重 DRAM 重读卡住（每个 8-token tile 都要读整段 gate+up 行，计算只占峰值 ~0.6%）。把 gate/up tile8 网格交换为 (tile, row-block)——相邻 block 沿同一专家的 token tile 重走同一段 32 行权重切片，切片留在 L2、专家权重从 DRAM ~只流一次（−10.3%）。同一交换在 down 内核上测得负收益、随即回退——"测量驱动"而非"直觉驱动"
占用率变体与 env 门（affa9c3）：全部逐位精确（同累加序）、env 门控、默认关：DS4_ROCM_ENABLE_MXFP4_TILE4 / _ROW64（512 线程、2 块/WGP）、DS4_ROCM_MXFP4_DOWN_RGROUP=1..8。测试在 n=128/512 快照默认路径输出，逐变体 memcmp 断言
make test-mxfp4-rocm / tests/test_mxfp4_rocm.c 按 DS4 Flash 真实形状（4096/2048、256 专家取 6）测——5d90eba 特意把 down 阶段的判定改成"对 GPU 自己产出的 mid 做 CPU 再量化再比"（真实形状下 1-LSB 舍入翻转会超出固定容差）

gfx1151 prefill 调优战役（2026-09，~226 → 290.79 tok/s）：一批以 bench-commit 节奏推进的优化（每轮数字记入 speed-bench/gfx1151-prefill-results.md）。方法论本身就是看点——kernel lab 工作流（61e9b99）：cuda/mmq/test/proto_*.cu 是独立原型基准（proto_gemm_lib.cu 用 cuBLASLt 16 算法 best-of 做上限参照），赢家经 ds4_rocm_gfx1151_flag("DS4_ROCM_*")（默认开、=0 回滚）接进生产内核。为了让 CUDA MMQ 代码在 HIP 下编译，vendored 了 cuda/mmq/vendors/hip.h（+308 行兼容层）。主要主题：

WMMA 注意力 tile：attention_mixed_heads16_wmma_kernel（rocwmma 16×16×16 片段、KEYS = 80、fp32 累加器）+ 稀疏注意力的 online 变体；注意力输出 GEMM 专门化 attention_output_b_f16_wmma_64x64_kernel；索引注意力 KV 暂存向量化为 b128 装载
IQ2 MoE 经 MMQ：moe_gate_up_mid_iq2_hotlist_wmma_n2_kernel / down 侧 Q2K 版、紧凑 launch（__launch_bounds__(32)）、成对 gate/up tile（一次激活装载喂两矩阵）、MMQ 里缓存 IQ2 符号掩码（sign_masks uint2 复用而非重算）
small-M F16 prefill GEMM 专门化：matmul_f16_smallm_wmma_kernel<M,N> 按输出形状模板化（prefill 的 M 通常极小，通用 GEMM 全在浪费 tile）
Q8 tiling：按输出形状 tile 的 matmul_q8_0_f32_batch_wmma_rowtile_kernel<256,16>；tiny batch 用 matmul_q8_0_f32_batch_sharedx_exact8_kernel<2..5>（24 VGPR 的内核胜过 72 VGPR 的占用率悬崖，2.4×）
MMQ 加固：c012d5c 稳定 gfx1151 MMQ 路由映射；f9c459c 修 MMQ scratch 生命周期与尾部读取；4b321cf release 构建零警告
DSpark 验证加速：DS4_ROCM_DSPARK_Q8_MMVQ / _ATTN_A_MMVQ（验证批的 Q8 vec 路径与分组 Q4 MMVQ）、verify 形状的 MoE 派发、非因果验证路径（DS4_DSPARK_VERIFY_NONCAUSAL 调试门）；rocBLAS 解方案钉扎：verify 形状 pin 手调 algo solution 索引（rocBLAS 5.6 集 DS4_ROCBLAS_F16_SOLUTIONS_5_6_8D1AE90E），0efd822 修了"解方案被 rocBLAS 拒绝后永远重试"——原子置位 g_rocblas_f16_solutions_disabled、永久回退默认 cublasGemmEx
工程化：make test-rocm 目标让无 NVIDIA卡的 ROCm 主机也能跑核心回归（7891f27 + babca3d Darwin 链接守卫 + 7470dd6 MMQ 对象链接）；9254c3d 让无 GPU 的 ROCm 测试构建跳过 GPU 路径；STRIXHALO.md 现推荐维护的 ROCm 10.0 容器（kyuz0/strix-halo-ds4-toolbox）

2026-09 第二轮（PR #992）：hipBLASLt 版本匹配 + DSpark 验证器注意力 + 内存预算

版本匹配的 Lt prefill 方案（1ef9bba）：rocm/ds4_rocm_hipblaslt.cuh 维护一张按（形状、输出类型、solution 索引）键控的 hipBLASLt 方案缓存，方案经 hipblaslt_ext::getAlgosFromIndex 枚举 + matmulIsAlgoSupported 校验；全局 g_hipblaslt_prefill_state（0 未查 / 1 版本匹配 / −1 禁用）保证手调 solution 索引只在匹配的 hipBLASLt 版本上生效——版本漂移时安全回退，杜绝"升级后索引指向别的算法"。同族还有 gfx1151 专属的 64e490b（IQ2 prefill 输入 staging 专门化）。
DSpark 验证器注意力投影（e83b106）：hipblaslt_dspark_attention_a 为验证器宽度 2-6 各建一个定宽方案（DS4_ROCBLAS_F16_SOLUTIONS_5_6_8D1AE90E 门 + ds4_rocm_is_gfx1151()），5/6 行算子 ~2.55× 于通用路径。
分配前的内存预算：新头 ds4_rocm_memory.h 把裸 hipMalloc* 包成带预算的 ds4_rocm_malloc/_malloc_host/_malloc_managed——ds4_rocm_allocation_fits() 排除 CMA 页并保留 2 GiB OS 余量，ds4_rocm_uses_host_ram() 识别 APU 整机内存，ds4_rocm_model_arena_bytes() 按模型大小选 arena（256 MiB 阈值，否则 1792 MiB）；ds4_linux_memory.h 的 ds4_linux_nonmovable_memory() 提供不可移动内存查询。配套 make test-rocm-memory / test-linux-memory——先问"装得下吗"再分配，把 OOM 从运行中段失败提前成启动期拒绝。
死代码清除（96f547a / 4a14eb8）：删掉不再被引用的 ds4_rocm_q8.cuh（−449 行）、ds4_rocm_moe.cuh、ds4_rocm_attention.cuh 旧版本及 CUDA/ROCm Q8 matvec 内核——内核战役的收尾是把没胜出的原型移出去。
跨后端 GLM 注意力掩码修复（0e9cc2d）：GLM 的 dense 边界（2051 token，由模型定义）此前硬编码在各后端、且归约存储未同步。统一为模型定义边界 + 填充选择 + 同步的归约存储后，Strix Halo 续 prefill 从 21 → 68 t/s（此前掩码错误触发多余计算路径）。回归 make test-glm-attention（387 行，CPU + ROCm 双跑）。

4d. 帮助系统（Help System）

新增结构化帮助模块 ds4_help.c/h，替代各工具中散乱的 usage 打印。API：

// ds4_help.h
typedef enum {
    DS4_HELP_DS4,      // ds4 CLI
    DS4_HELP_SERVER,   // ds4-server
    DS4_HELP_AGENT,    // ds4-agent
    DS4_HELP_BENCH,    // ds4-bench
    DS4_HELP_EVAL,     // ds4-eval
} ds4_help_tool;

void ds4_help_print(FILE *fp, ds4_help_tool tool, const char *topic);

关键设计：

主题过滤：topic 参数可查询特定子主题（如 --help distributed）
TTY 感知：isatty() 自动检测，终端中显示 ANSI 256-color 彩色输出，非终端输出纯文本
格式化助手：opt() 打印彩色选项行、para() 打印黄色段落、title() 打印章节标题
分布式模式指导：帮助文本中包含 distributed mode 的用法说明和示例

4e. 模型下载脚本更新

download_model.sh 新增 PRO 模型下载目标：

目标	说明	大小
`pro-q2-imatrix`	PRO IQ2_XXS 单文件	~430 GB
`pro-q4-layers00-30`	PRO Q4 前半层（coordinator）	~426 GB
`pro-q4-layers31-output`	PRO Q4 后半层（worker）	~412 GB
`pro-q4-split`	下载两个 Q4 分片	~838 GB

2026-09 追加 V4.1 Flash 目标（HF 仓库 antirez/deepseek-v4.1-flash-gguf，五个工件全部钉死尺寸 + SHA-256，断点续传路径同样校验）：ds41f-q2（341 GiB）、ds41f-q4（483 GiB，两段式：480 GB + ~36 GB，flock 保护的 .assembling.lock 下拼接——中断会把输出截断到边界只重拼尾部，全程另需 ~37 GiB 空闲盘）、ds41f-vision（0.9 GiB）。新增 make test-download-model 离线测 8 个用例。

PRO 文件过大，脚本自动检测并使用 Hugging Face CLI（hf from huggingface_hub）下载，支持断点续传。未安装时提示 python3 -m pip install -U huggingface_hub hf_xet。

非 PRO 目标下载后自动符号链接到 ./ds4flash.gguf，PRO 目标打印分布式用法指引。

4f. 方向性激活引导（Directional Steering）

基于论文 Refusal in Language Models Is Mediated by a Single Direction 的思想，在推理时对每层激活做低秩编辑：

y = y - scale * direction[layer] * dot(direction[layer], y)

引导文件：43 × 4096 的 f32 矩阵（43 层，每层一个 4096 维归一化方向向量）
可在 attention 输出后（--dir-steering-attn）和/或 FFN 输出后（--dir-steering-ffn）施加
正 scale 移除该方向，负 scale 放大该方向
提取工具：dir-steering/tools/build_direction.py（用好/坏两组 prompt 取激活差异，归一化后生成 .f32 文件）
示例：dir-steering/ 包含预构建的 verbosity 方向向量，负 FFN scale 使回答更简洁

2026-09 扩展（详见方向性引导术语表）：GLM 5.3 也支持（8adb221；GLM 5.2 仍不支持）——方向形状 45×4096（MTP 层除外），build_direction.py 加 --profile glm-5.3-flash；流式 GLM 图按 tier 索引方向（directional_steering_dirs_by_tier）。Live steering（87495f6）：新库级 API ds4_session_set_directional_steering_ffn() 让 scale 成为会话运行属性——不重建 KV、运行中可调；REPL 里 /steer（显示 / 0 关 / F ∈ −100..100），agent 里要求模型空闲。改 scale 会使 DSpark 捕获状态失效（需重建）；分布式/TP 会话拒绝。

4g. CUDA Q8→F16 缓存显存预算守卫

PRO 模型 GPU 注意：DeepSeek V4 PRO（1.6T 总参数，49B 激活参数）的专家张量更大（更多专家 × 更高维度），Q4 量化模型需要 512GB 内存的 Mac Studio。GPU 后端在加载 PRO 模型时会自动调整 KV cache 容量和 buffer 分配。

Q8 权重在推理时被扩展为 F16 存入 GPU 显存，加速后续 cuBLAS 矩阵乘法。但显存有限的 GPU 上，缓存可能占满 VRAM 导致后续分配失败：

// 显存预算检查流程
static int cuda_q8_f16_cache_has_budget(uint64_t request_bytes, const char *label) {
    // 1. 检查用户设定的硬上限
    const uint64_t limit = cuda_parse_mib_env("DS4_CUDA_Q8_F16_CACHE_MB");
    // 2. 查询当前显存空闲量
    cudaMemGetInfo(&free_bytes, &total_bytes);
    // 3. 计算保留量：max(4 GiB, 5% VRAM)
    reserve = max(4ull * 1024 * 1024 * 1024, total_bytes / 20);
    // 4. 分配后剩余空间必须 >= reserve
    return (free_bytes - request_bytes >= reserve);
}

三层回退机制：

关键设计：缓存是纯加速路径，失败不影响正确性。环境变量控制：

DS4_CUDA_Q8_F16_CACHE_MB=0：完全禁用
DS4_CUDA_Q8_F16_CACHE_MB=8192：限制缓存 8 GiB
DS4_CUDA_Q8_F16_CACHE_RESERVE_MB=8192：覆盖默认保留量

4h. CUDA 长上下文修复

长上下文（>100K token）暴露了多个 CUDA kernel 的固定缓冲区假设：

Shared Memory 分数溢出 → 在线注意力 kernel 回退：

// 编译时常量限制了 shared memory 中的分数缓冲区大小
enum { DS4_CUDA_ATTENTION_SCORE_CAP = 8192u };
// 当压缩行数超过缓冲区容量时，切换到在线注意力 kernel
// 在线 kernel 不需要预分配所有分数，逐块处理

多级 Tree-Merge TopK → 替代单层 chunk+merge：

// 旧：单层 chunk(2048) → merge，无法处理 >4096 压缩行
// 新：chunk(4096) → 多级 tree-merge，每 8 组合并一次
// n_sets: n_chunks → ceil(n/8) → ceil(n/64) → ... → ≤8 → 最终合并

回归测试 tests/cuda_long_context_smoke.c：模拟长上下文场景（多个 prefill frontier + decode），验证 CUDA 路径在边界条件下的正确性。

MoE Down 路径清理

移除了 block16 MoE down 诊断路径（-172 行），因存在 top-logit 不稳定性和贪心首 token 行为差异。默认保留更快的 tile16/row2048 路径，block16 改为 DS4_CUDA_MOE_DOWN_BLOCK16 环境变量 opt-in。

CUDA 内存优化

三项改进减少不必要的 GPU 开销：compressor state 清零改用 cudaMemsetAsync（替代 fill kernel 调用）；graph/session tensor 改用 cudaMalloc 设备内存替代 managed memory（避免多余的页面迁移）；新增 backend fill_f32 hook 让 CUDA 设备端直接初始化张量。

CUDA Managed KV Cache

百万 token 级上下文时，KV cache 张量改用 cudaMallocManaged（统一内存按需分页），避免 DGX Spark 统一内存系统中 GPU 显存分配饿死 CPU 侧。普通上下文仍用 cudaMalloc 设备分配保持性能。通过 ds4_gpu_should_use_managed_kv_cache() 自动判断是否启用。

Metal Q4 Model View 回退

Metal Q4 expert tensor model views 的扩展（DS4_METAL_MODEL_MAX_TENSOR_BYTES 从 2 GiB 回到 ~672 MiB）已被回退——该修复由 AI 生成且未经用户确认，上游坚持不接受未经人工验证的修复。

4i. CUDA Prefill 性能优化

大幅优化长上下文 prefill 速度（32K prefill 255→346 tok/s，全曲线均值 316→370 tok/s），核心改动：

更宽的 WMMA 评分 kernel： 原有 indexer_scores_wmma_kernel 每 tile 处理 16 个压缩 token。新增 wmma32/64/128 变体（默认 128 宽，256 线程），将压缩索引预加载到 shared memory 后迭代所有注意力头，减少 kernel 启动开销并提高内存吞吐。

CUB BlockRadixSort 替代手写 bitonic sort： 新 indexer_topk_8192_cub_kernel 用 NVIDIA CUB 库的 BlockRadixSort（512 线程 × 16 items = 8192/block），将 float key 打包为 uint64_t 实现降序排序。通过 cudaFuncSetAttribute 请求额外 shared memory。

uint16_t 索引的 bitonic sort： 对 ≤65536 压缩 token 的场景，用 uint16_t 代替 uint32_t 索引，减半 shared memory 占用。

Top-K 预排序提升注意力访存： 新 indexed_topk_sort_512_asc_kernel 将 top-k 选出的 512 个索引按升序排列后再送入注意力 kernel，把对压缩 KV cache 的随机访问转变为顺序访问，显著改善内存合并。

Templatized 注意力 kernel： attention_indexed_mixed_heads8_online_kernel 重构为模板，参数化 ROWS_PER_STAGE 和 HEADS_PER_GROUP，启动配置从 <8, 8> 改为 <8, 16>（16 head/group, 512 线程），配合动态 shared memory。

Shared memory 复用： 原有 wmma kernel 交换加载顺序，压缩索引 tile 只加载一次（在头循环之前），避免每头重复读取全局内存。

// ds4.h — 引导选项
typedef struct {
    // ...
    const char *directional_steering_file;
    float directional_steering_attn;
    float directional_steering_ffn;
} ds4_engine_options;

4j. 张量并行（Tensor Parallelism）

流水线并行按层切片、把模型拆到多机装配线上，主要为了"塞下更大的模型"。张量并行走另一条路：两台机器在同一个 token 上同时工作，把每层最重的矩阵运算切给两边，在图的同步门处交换部分和——它降低的是单 token 延迟，decode 反而更快。详见张量并行术语表。

实现集中在 ds4_tp.c / ds4_tp.h（~3K 行）。它有两种部署形态，共用同一套 lockstep 协议：

形态一：Mac-to-Mac（RDMA over Thunderbolt）。Leader 是一个普通前端 session，把每一次 ds4_session_sync() / ds4_session_eval() 原样镜像给 Worker，于是两个引擎执行完全相同的图序列。每台机器常驻连续的一半路由专家；dense、注意力、共享专家、embedding、output 在两边复制——路由专家内核永不碰对方那一半。在一个 token 内部，部分和在每层的同步门处交换：

// ds4_tp.h — 每层两个同步门
enum {
    DS4_TP_GATE_ATTN = 0,          // 注意力门
    DS4_TP_GATE_FFN  = 1,          // FFN 门
    DS4_TP_GATES_PER_LAYER = 2,
    DS4_TP_BATCH_MAX_ROWS = 8,     // 推测验证块最多 5 行
};

// 把本机 partial 发给对方，等对方 partial 到位（GPU gate 服务线程调用）
int ds4_tp_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t gate, uint64_t seq);

部分和是 f32、从不在网络上量化（vec_bytes = n_embd * 4）。DeepSeek 的 gate 向量 16 KB 一条 RDMA 消息；GLM 的 6144 宽 24 KB 向量拆两条。传输优先 RDMA two-sided SEND/RECV，回退全双工 TCP（--transport tcp）；一个 GPU 可见的共享 slab 充当收发缓冲，NIC 注册它并交换 remote key。

握手时交换 ds4_tp_identity（GGUF 字节、model_id、层数、embd、词表、量化位、上下文、解码门调度）。门调度决定 RDMA recv 落在 slab 哪个槽——DeepSeek 每层先 ATTN 后 FFN；GLM 只在稀疏层发一个 FFN 门，调度跳过 dense 前缀和 ATTN 槽。两边必须一致，否则推理前 abort。

形态二：CUDA 张量并行（单机多卡）。--cuda-tensor-parallel 把 Flash 的张量和路由专家工作切到偶数块 GPU，独立于 Mac 模式（不用 --role、不用 RDMA）。设备顺序很关键：N 块卡时前 N/2 个是"流水线层之家"、后 N/2 个是张量并行搭档，先列所有 home 再列所有 partner，配对位置是最亲近的 P2P 对（8×L40S 写 0,2,4,6,1,3,5,7）。每对存 50/50 路由专家、词表头按行分片（不复制）；dense/router/共享专家在每对内复制。

// ds4.h — 两种并行的角色与传输
typedef enum { DS4_TP_NONE=0, DS4_TP_LEADER, DS4_TP_WORKER } ds4_tp_role;
typedef enum { DS4_TP_TRANSPORT_AUTO=0, DS4_TP_TRANSPORT_RDMA, DS4_TP_TRANSPORT_TCP } ds4_tp_transport;

typedef struct {
    ds4_tp_role role;
    ds4_tp_transport transport;
    const char *rdma_device; int rdma_gid_index;
    bool glm_token_prefill;
    int  debug_hash;   // 每 N token 跨机核对隐藏状态
    // ...
} ds4_tp_options;

TP 要求 Metal 后端，且不能与 SSD streaming、分布式模式、层切片（--dist-* / load_slice）同时启用（ds4_tp_validate_engine_options()）。推测解码已于 2026-09 解禁：DSpark 的验证块经 DS4_TP_FRAME_VERIFY 镜像给 worker（旧单阶段 MTP 在 TP 下退回逐 token decode）——见下文实测（TP 下 DSpark 反而更慢）。GLM 5.2 在 CUDA 上走普通层放置（非 TP）；DGX Spark 是单卡，不能用 --cuda-tensor-parallel。

双 Mac TP 延迟工程战役（2026-09，`4886c56` 合入 `macm5us/mac-tp-1-sep-2026`）

TP 协议升到 version 10（新增帧 DS4_TP_FRAME_RDMA_WARM=19 / RDMA_POSTED=20 / GLM_MTP=21；两边必须同版本，否则握手 abort）。一周 ~25 个提交几乎全是延迟工程——两个 rank 每 token 要在每层的门上碰面，任何一侧的微秒级抖动都直接加进单 token 延迟：

优化	机制	关键锚点
轮询门（poll gates）	CPU→GPU 共享事件唤醒要 27-43 µs，改成 16 个 flag 区环形轮询（~600 ms 预算）	`DS4_TP_DISABLE_POLL_GATES` 回退
校验和 flag	flag 行可能比 partial 的 store 早 ~35 µs 落地——裸 flag 会读到半成品；整数校验和防伪	`kernel_dsv4_tp_flag_set_checked`
flag 折叠	把"发 flag"折进产出 partial 的载荷内核本身，省一次独立 dispatch	`kernel_dsv4_add2_f32_tp_flag_checked`、`kernel_dsv4_mul_mv_q8_0_f32_tp_flag_checked`、`ds4_gpu_tp_flag_fold_request`
RDMA 预热	Apple UC verbs 上向没有 posted receive 的 QP 发送会永久杀死该方向；先握手预热，失败则重建 QP/CQ（新 PSN，至多 4 次）	`tp_rdma_warm_up` / `tp_rdma_recreate_qp`
分级大门	prefill 每层 32 MB 的 big gate 改成 256×16 KB 窗口 + 链式 posted receive（~26 ms → 4-6 ms）	`tp_rdma_window_barrier`；FFN 子门流水线 opt-in `DS4_TP_FFN_SUBGATE`
队列保活	M5 Max 上空闲 ~3 s 后提交的命令缓冲晚 600-800 ms 才启动——每秒提交一个 1 KB 内核保温	`ds4_gpu_warm_command_queue`、`DS4_METAL_DISABLE_QUEUE_KEEPALIVE`
GPU 决定的负载均衡	由本 rank 选中的专家 id 在 GPU 上算出共享专家车道互补切分（多领一个路由专家就把共享专家让出一半），消除"静等慢 rank"	`ds4_gpu_parallel_ffn_start_sliced/_split`（`shift_q16` = 路由/共享字节比 Q16）、`kernel_dsv4_shared_gate_up_swiglu_q8_0_split` / `kernel_dsv4_shared_down_q8_0_split`
推测验证预投递	验证块的 RDMA 窗口提前投递，验证提交从布尔 `full_accept` 升级为三态	`ds4_tp_batch_block_begin/end` + `ds4_tp_verify_commit_mode {ROLLBACK_REPLAY, COMMIT_FULL, COMMIT_PREFIX}`（partial accept 也能直接提交前缀快照）

注意力 split count 走了"32 → 自适应 8 → 回到 32"的弯路（DS4_METAL_FLASH_NWG 覆盖仍可用）；配套还有 kv-norm 与 q_b 并发的延迟任务族（ds4_gpu_kv_norm_task_*）与两 rank 融合路由选择。bbec200 修复混合 2/4/6-session TP 批的延迟残差后，混合批 TP 的 logits 与串行完全一致。计时驱动的车道偏置 DS4_TP_LANE_BIAS 因让输出依赖时序而保持默认关闭——"更快但不确定"不做默认。

实测（两台 M5 Max 128 GB、Thunderbolt 5、MXFP4 配方，来自 README 7fc640c 时点）：

	双 Mac TP	单 Mac SSD streaming
decode @128 ctx	~47-48 t/s	~11 t/s
decode @2048 ctx	~48 t/s	—
prefill	~250 t/s @128 token、~850 t/s @2048	~10 t/s
启动	~10 s/台（钉住分片）	—
DSpark	30-40 t/s，比纯 decode 慢（每个验证行都重流它的专家）	—

prefill 从 ~93 t/s 起步，经队列保活、分级 RDMA 窗口、TensorOps 路由 GEMM（§4p）逐级抬到上表数字。注意 coordinator 上的 GPU 抢占：远程屏共享或动态壁纸可吃掉 10% decode。部署细节（rdma_ctl status / ibv_devinfo -v、成员接口 IP、--rdma-device / --rdma-gid-index、worker 先启、不要 curl 探测等待中的 coordinator——探测会被当成 worker 握手）全部移到了上游 docs/DISTRIBUTED.md；TP 磁盘缓存恢复现在在两个 rank 上重建保存的 token 前缀（恢复时预期会 prefill）。

2026-09 协议加固 + V4.1 入场（bd66c40，协议版本 10→12）：

同步检查点帧：新 DS4_TP_FRAME_SYNC_CHECKPOINT (22) 与 ds4_tp_sync_checkpoint()——两 rank 在匹配的 prefill 边界交换 {seq, point, current, total, cancelled}，取消由双方协商一致，绝不在一方还在 GPU 门里时就各自采样决定；worker 把 DS4_SESSION_SYNC_INTERRUPTED 视为非致命、只在 sync 完整成功后才发 logits
RDMA 语义修复：所有发送改 IBV_SEND_SIGNALED + 逐发送计数（Apple 提供方会把 unsignaled 发送也完成，之前暂存内存被提前释放）；tp_rdma_window_barrier()（标签 0xD1）修 UC 模式下"发送先于对端 post receive 到达被丢弃"的竞态；大门的逐 chunk CQE 用位图核对、拒绝重复
V4.1 Flash TP：Q2 双 128 GB Mac（每 rank ~81 GiB 主权重，两台都要完整 341 GiB GGUF——Engram 表两 rank 各读各的行，永远在盘上）；输出头改为复制而非切分（tp.vocab_split 仅保留给 V4 Flash）；第 2 张 Engram 表要在 il==13 排空再复用输入缓冲；续 prefill 的批处理下限 TP 是 32 token（resident 8）。原生会话解码 3–8 路、2 路有序兜底（§4q）

2026-09 中旬：双 Spark TP（a04f46f，协议版本 → 14）。V4.1 Q2 文本推理支持两台 DGX Spark 网络 TP：每 rank 一块 GPU、专家 50/50 对半（每台 ~81 GiB），不与 --ssd-streaming / --cuda-tensor-parallel 混用。传输是 Linux 侧 RoCEv2 + RC 队列对（libibverbs 运行时 dlopen，控制 socket 的直连地址自动选 RoCEv2 GID，--rdma-device/--rdma-gid-index 可覆盖；TCP --transport tcp 兜底）。与 Metal 路径的关键差异：CUDA 结果经 pinned host 内存中转——这是 RoCE 不是 GPUDirect；且交换模型不同——Metal 把每个门排在服务线程上异步等待，CUDA 等自己的本地流然后同步调交换（ds4_gpu_tp.h 的 gate/batch/big-exchange 回调族，big exchange 走 CUDA→pinned staging→传输→写回，staging 按需增长）。注意力输出每 rank 发未舍入的部分和、图侧求和后才做 BF16 舍入；共享专家与路由专家并发跑；decode 用 CUDA 图捕获每 rank 每层 3 个区域 × 40 层 = 120 个区域。协议 14 要求所有 peer 一起升级（含 Metal peer）。实测见 §4b（单会话 21.9 t/s、8 会话聚合 ~28、5 个起批）。

4k. 多 GPU 放置与层打包（Multi-GPU Placement）

CUDA 多卡和分布式都需要回答"哪层放哪块卡"。这块能力拆成了三个新模块：

GPU 参数解析（ds4_gpu_args.c/h）—— --gpu-vram 和 --gpu-devices 的统一解析器，ds4 / server / agent / bench 共用：

// ds4_gpu_args.h
// vram_arg: NULL | "0"(跳过 CUDA) | "auto"(查空闲显存) | "N[,N...]"(每卡 GB 预算)
// devices_arg: NULL(0..n-1) | "N[,N...]"(CUDA 设备号)
int parse_gpu_vram_arg(const char *vram_arg, const char *devices_arg,
                       ds4_gpu_config *out, bool *out_skip_cuda,
                       char *errbuf, size_t errbuflen);

auto 仅 CUDA 构建有（查 cudaMemGetInfo），Metal/CPU 构建返回错误。两个参数都给且数量不等、负数、超过 DS4_MAX_GPUS 都报错。--gpu-vram 0 表示完全跳过 CUDA。

层打包器（ds4_layer_pack.c/h）—— 把层序列单调连续地放进各卡预算，算出每个层 → 设备的映射：

// ds4_layer_pack.h
#define DS4_LAYER_PACK_MAX_GPUS 16
#define DS4_LAYER_PACK_CPU      (-1)   // 溢出到 CPU 层级

// entry 顺序: 0=embedding 伪层, 1..n_layers=transformer 层, n_layers+1=output 伪层
// 输出 device_for_entry[]; 超过所有 GPU 预算的层溢出到 CPU，
// 且据单调性其后所有层也落 CPU——没有"层太大"的错误路径，CPU 层级永远兜底
int ds4_compute_layer_placement(const size_t *entry_bytes, int n_entries,
                                const ds4_layer_pack_config *cfg,
                                int *device_for_entry);

启动时打印的布局形如 GPU0: layers 0-21 + embedding (38.4 / 40.0 GB) / CPU: layers 32-42 + output head。

多 GPU 管线类型（ds4_gpu_mgpu.h）—— 这是多 GPU 工作的唯一真相源（因 ds4_gpu.h 历史上不被 ds4_cuda.cu 包含）。它给出了完整的 ds4_gpu_tensor（带 device_id）、ds4_gpu_config、全局 g_gpu[16] / g_n_gpus / g_gpu_peer_ok[][]，以及跨设备拷贝原语：

// ds4_gpu_mgpu.h（节选）
struct ds4_gpu_tensor { void *ptr; uint64_t bytes; int owner; int device_id; };
typedef struct ds4_gpu_config {
    int    device_indices[DS4_MAX_GPUS];
    size_t vram_bytes[DS4_MAX_GPUS];   // 每卡字节预算；0 = 零预算（会逼出 CPU 溢出）
    int    n_gpus;
    size_t safety_margin_bytes;        // 每卡保留量
} ds4_gpu_config;

int ds4_gpu_init_multi(const ds4_gpu_config *cfg);
// 跨设备拷贝: 同卡 cudaMemcpyAsync / 可 P2P 的 cudaMemcpyPeerAsync / 否则 pinned host 中转
int ds4_gpu_tensor_copy_xdev(ds4_gpu_tensor *dst, const ds4_gpu_tensor *src, uint64_t bytes);
// 张量并行的跨设备浮点加法（部分和归约）
int ds4_gpu_add_xdev_tensor(ds4_gpu_tensor *out, const ds4_gpu_tensor *local,
                            const ds4_gpu_tensor *remote, ds4_gpu_tensor *remote_tmp, uint32_t n);

引擎入口 ds4_engine_create_with_gpu_config() 接受可选的 ds4_gpu_config；传 NULL 与 ds4_engine_open 完全等价。层打包的 KV 定价用 placement_ctx_hint 估算每层 KV 存储。--gpu-vram auto 用空闲显存减去保留量（max(4 GiB, 5% VRAM)）做预算。

4l. GLM 5.2 / 5.3 后端支持

DwarfStar 不再是 DeepSeek 专用——现在支持 GLM 5.2 与 GLM 5.3（后者详见 Part 3 §6 与 KDA 术语表）。GLM 在三个图后端（Metal / CUDA / ROCm）上都能跑，关键差异在量化布局和批处理能力：

路由专家支持 IQ2_XXS / Q2_K / Q4_K / Q5_K（gate/up）和 Q2_K/Q4_K/Q5_K/Q6_K（down）；dense / 模型控制张量走现有 Q8/F32 路径
批处理：Metal GLM 只能走有序精确兜底（Flash 才有原生共享专家 + QKV 批处理）；CUDA 上 GLM 走普通层放置，不用 --cuda-tensor-parallel
限制：GLM 5.2 不支持方向性引导、--power < 100、显式 --prefill-chunk、外部 --mtp 文件；MTP 块在主 GGUF 内部，用 --glm-mtp / GLM 5.3 的 --mtp 启用实验性贪心推测（--mtp-timing 看接受计数）
SSD streaming：GLM 也支持（Metal + ROCm），路径保留能装下的最大全层前缀常驻，余量给动态专家缓存；GLM 5.3 的 glm53-full-q2 把它扩到全量检查点（常驻层区间整层映射 + 自适应内存预算守卫 glm_graph_memory_guard_budget）
两机张量并行：要求 ownership-aware 的 IQ2_XXS 或 Q2_K 路由布局；路由 Q4 的 GLM 必须在评估前拒绝

ROCm 侧为 GLM 做了大量调优（rocm/ds4_rocm_glm.cuh，GLM 5.3 后已 ~5.3K 行）：selected/causal 注意力 GEMM、wave value projection、slice decode、routed IQ2 down 等，多数是 opt-in 实验路径，验证有效后才转为默认。GLM 5.3 的 KDA 层在 ROCm 上有整套内核（glm53_rocm_kda_decode_kernel + prefill 三步 prepare/recurrence/output，常量 GLM53_ROCM_KDA_DIM = 128、GLM53_ROCM_KDA_HISTORY = 3）；稀疏 DSA 层走 glm_attention_full_kernel / glm_attention_indexed_lora_kernel 及 decode split 归约族。fea665c 修了一个教科书级竞态：glm_attention_indexed_lora_kernel 的块归约写 max_score = red[0] 后缺一道 __syncthreads()——快的线程会在别人还在读共享 red 缓冲时就覆写它。

Metal 上的 GLM 5.3：metal/glm53_kda.metal（decode + prefill 三内核）与 metal/glm53_bf16.metal。两条性能路径值得记：

M3 Ultra BF16 融合 QKV（6cf658a）：Q4 工件的 kda_q/k/v 是 BF16，ds4_gpu_glm53_matmul_bf16_qkv 一个内核产出三个矩阵（kernel_glm53_mul_mv_bf16_f32_qkv，每输出行一个 simdgroup、8 路独立 load 隐藏延迟）——每 KDA 层砍掉两次 dispatch，M3 Ultra 512GB decode 24.74 t/s；DS4_METAL_DISABLE_M3_ULTRA_GLM53_DECODE / _GLM53_BF16_QKV 可回滚
FP32 紧凑 KV cache（81c8c39）：GLM 5.3 的 DSA cache 行没有 RoPE 尾巴（n_rot = 0），FP32 存储因此无旋转混合的复杂性——DS4_GPU_GLM_COMPACT_CACHE_F16 编译期开关让同一套内核按 cache_f16 选 f16/f32 元素宽度

续 prefill 正确性（639d4eb）：GLM 5.3 常驻图会在 prefill 中途切换注意力切片，短续写后缀必须强制 chunk = 1 并跳过索引/batch prefill 路径（glm53_graph_resume_prefill_min_tokens）——带 199 行回归测试。

填充选择的解码掩码修复（582d809，2026-09）：GLM decode 的 selected_rows_valid 旗标向内核承诺"所有选中行都是合法 cache 行"，从而走免边界检查的快内核。但 GLM 5.3 的池化扩展只保证 [0, indexer_top_k) 槽位有效，因果尾部的填充槽是 0xffffffff 哨兵——承诺落空就是越界读/垃圾注意力。修复引入 last_indexer_guaranteed_prefix 追踪"保证有效的连续前缀"，有填充时退回现成的边界检查内核（干净选择仍走快内核）。这是"快路径的谓词必须由生产者证明，不能由调用者许愿"的典型教训；9775e26 用独立双精度参考把全部池尾长度（0–3）钉进回归。

CUDA：vendored MMQ 的 GLM 移植（b1b4ea0）在 DGX Spark 上加速 GLM prefill（见 §4m）。

4m. CUDA 性能（vendored MMQ + decode-island graphs + Blackwell）

本轮把 batched-serving fork 的工作接入上游，CUDA 性能大幅提升，分三块：

Vendored MMQ prefill tier（cuda/mmq/，~3K 行）从 Entrpi/ds4 fork 引入，pin 自 llama.cpp 的 Multi-Marlin Quantization 内核（MIT，cuda/mmq/VENDOR.md 记录版本）。只接 RAW-layout 层：n_tok≥2 时 dense Q8_0 GEMM + IQ2_XXS gate/up + Q2_K down 的路由 MoE 管线；decode（n_tok==1）不动。代价是 MMQ 的 FP32 归约顺序与 cublas+dequant 不同 → prefill logits 在 ULP 尺度漂移，故仅在 quality 关闭时启用，DS4_CUDA_MMQ=0 回退到旧调度。

构建陷阱（6f087f4）：make cuda-spark 之前没设 CUDA_ARCH，nvcc 退到默认架构，把 MMQ 张量核路径整个编掉——实测 39 vs 423 prefill tok/s。现在 cuda-spark 用 native arch。

Decode-island CUDA graph 捕获（e50104f）：串行 decode 每个 token 都重启一堆 kernel，launch 开销累积。把每个 decode 层的两个位置无关岛屿——hc-pre 前缀（恰好 TO_QKV 阶段前缀）和注意力输出到层尾的尾部（恰好 FROM_ATTN_TO_FFN 阶段后缀）——按激活缓冲变体捕获一次、后续 token 直接回放。捕获/回放复用现有 phase 机制，无重复编码路径：

首次见某 key → eager 跑（预热惰性分配器）
二次见 → 捕获 + 实例化 + 启动
之后 → 回放
任何失败 → 退役该条目并 eager 重编码，不会丢 token 的工作

TP / 多级 / SSD streaming / debug dump / profiling / reference kernel / hash-router 层（token 相关的路由参数）都自动 disqualified 保持 eager；DS4_CUDA_DECODE_GRAPHS=0 全关。DGX Spark GB10 上 decode +1-2%——说明 2k 上下文下 launch 开销不是串行 decode 的主成本，这个捕获层是为后续覆盖位置相关中段（device-scalar substrates）铺的地基。

Blackwell sm_120a（b1f8893）：RTX 50 / PRO Blackwell 编为 compute_120a/sm_120a 才能接受 block-scaled MXFP4 指令（见 MXFP4）；源码同时接受 sm_120 与 sm_120a 两种写法。SM121（GB10）侧也用 MXFP4 做宽 indexer 评分、为 heads8 注意力 pin 占用率。

配套内核：一批 HMMA / indexed-attention 优化（token-tiled 注意力、indexer scorer 去 shared-memory bank-conflict、把 HC RMS norm 折进 FP16 投影输入）、小 Q4 张量并行批次加速、批处理 server 的 prefill quantum 可配、GLM 批处理 session 内存计入多卡放置。

2026-09 增量：vendored MMQ 移植到 HIP——cuda/mmq/vendors/hip.h（+308 行）让同一套 MMQ 源码在 ROCm 下编译，b1b4ea0 借此在 DGX Spark 上加速 GLM prefill；f9c459c 修 MMQ scratch 生命周期与尾部读取（见 §4c）。

2026-09 第二轮：解码复用对齐的 Q8 scratch（b42682f，make test-cuda-q8-scratch）；DSpark 在 CUDA 上保留全部五个部分前缀——部分接受不再回滚重放（f6c169f，ROCm 侧保留四个、第五个重放，394865c）；没胜出的 Q8 matvec 内核成对删除（4a14eb8，与 ROCm 侧清理同一次提交）。

4n. Metal 精确解码融合战役（Exact Decode Fusions）

2026-08 的一整轮 Metal 优化战役（4e401a1 → 0e8fb54，含 Ivan Fioravanti 的 exact MXFP4 加速合并 fd2d5e0）把 M5 Max 的 q2 decode 从 39.4 推到 ~45 t/s（+15%），并把同一套融合移植回 M1-M4（M3 Ultra +7.6%）。它的方法论比数字更值得学。

动机：decode 已经 ~97.5-98% GPU-busy——不是内核不够快，而是每个 token 的 dispatch 太多（M3 Ultra 每层 ~25 次、每次 ~3.9µs 的启动开销）。出路是把相邻小内核融合成一次 dispatch，同时绝不改变数值。

"精确"（exact）的合同：每个融合内核必须与旧的多 dispatch 路径产生逐位相同的全词表 logits 和相同 token 选择——不是"误差可接受"，是 bit-identical。验证靠两个新 A/B 基准（构建为 make metal-decode-schedule-bench / make metal-prefill-variant-bench，源码在 speed-bench/）：

decode 侧：prefill 两个 session，同时交替变体顺序和变体↔session 分配，除非每一行 full-vocab logits 逐位相同（且 --include-selection 时选中相同 token）否则 abort
prefill 侧：ABBA/BAAB 块序、每变体独立新 session、≥32 token 预热、拷贝前毒化 host logits 缓冲（防止"看起来对其实是没写"的假通过）、8192-token 前缀默认值，同样要求逐位相同

融合清单（kernel 名 ↔ 回滚环境变量，PRE_M5 前缀 = M1-M4 也默认启用）：

融合	内核	回滚开关（`DS4_METAL_DISABLE_...`）	收益
HC RMS norm+mix	`kernel_dsv4_hc_rms_norm_mix_f16*`	`_PRE_M5_HC_NORM_MIX_FUSE`	+0.1%
四路 F16 压缩投影 + 状态存储	`kernel_mul_mv_f16_f32_pair_4_impl<2>`	`_PRE_M5_COMPRESSOR_QUAD_STORE`	+0.2%
路由器 + 共享专家 gate/up	`kernel_dsv4_router_shared_gate_up_q8_0`	`_PRE_M5_ROUTER_SHARED_FUSE`	+0.3%
QKV norm + KV RoPE + FP8 存储	`kernel_dsv4_kv_rope_fp8_store_f32`	`_PRE_M5_QKV_NORM_KV_STORE_FUSE`	+0.5-0.6%
注意力逆 RoPE 折进 FA reduce	`kernel_flash_attn_ext_vec_reduce_rope`	`_PRE_M5_ATTN_INV_ROPE_FUSE`	+0.1-0.3%
ratio-4 压缩器打包	`kernel_dsv4_compressor_pack_ratio4_decode_ggml`	`_PRE_M5_COMPRESSOR_RATIO4_DECODE_PACK_FUSION`	+0.3-0.4%
q_a/kv pair + quad 压缩器存储	`kernel_dsv4_qkv_pair_quad_compressor_store_q8_0`	`_PRE_M5_QKV_PAIR_QUAD_FUSE`	+0.5%
emit 路径压缩器 finalize	`kernel_dsv4_comp_row_finalize_f32`（每层 11 次单行 dispatch → 1 次）	`_PRE_M5_COMP_FINALIZE_FUSE`	+0.4-0.6%
压缩器精确 softmax+product	`kernel_dsv4_compressor_exact_softmax_product_ratio4`	—	+0.1-0.2%
聚簇 HC norm/mix 生产者	`kernel_dsv4_hc_rms_norm_mix_f16_cluster2`（两个 8-simdgroup 簇共用一个 512 线程组，保留精确 1024 虚拟线程 RMS 归约树）	—	+0.7-0.9%
打包 split-K FlashAttention 归约	`kernel_dsv4_flash_attn_vec_packed32_reduce_rope_f16_dk512_dv512`	`_PRE_M5_FLASH_ATTN_PACKED32_REDUCE`	+3.95%（单笔最大）
HC 生产者/pre-norm 续接	`kernel_dsv4_hc_rms_norm_mix_f16_cluster2_pre_norm`（跨 threadgroup 栅栏完成协议）	—	+1.2-2.9%
并发全 FFN	`MTLDispatchTypeConcurrent` 同命令缓冲调度（共享 Q8 gate/up 与路由 IQ2XXS pair-SwiGLU 独立启动 + 显式资源屏障）	`_PRE_M5_PARALLEL_FULL_FFN`	+1.2%

其中打包 split-K FA 归约值得展开：32 个 split-K 工作组让 8 个物理 simdgroup 分时复用、在 threadgroup 内存里物化精确 F32 权重与 (S,M) 统计、经衬垫的 32×33 平面做合并 V 偏导、保留原 32-lane 归约树并融合逆 RoPE——一次性消除了每层 4.2 MB 的 device 端 partial 往返。注意它的门很窄：仅 M5、F16 512 宽 K/V、单 query、64 头、≤1024 key 的特定几何。

自动门控（0ad494e）：战役初期每个融合要 DS4_METAL_ENABLE_* 手动开；收官时删掉了 ~76 个 ENABLE 开关和 ~27 个废弃的 M3 时代 DISABLE 别名（统一改叫 PRE_M5），融合路径按设备 + 形状默认启用：

// ds4.c — 统一开关（M5 端口默认走 M1-M4 标准路径之外的新融合）
static bool metal_graph_ported_m5_decode_feature_enabled(
        const char *pre_m5_disable_env, const char *m5_disable_env) {
    if (m5_disable_env && getenv(m5_disable_env) != NULL) return false;
    const bool pre_m5 = ds4_gpu_device_is_pre_m5_apple_silicon();  // 设备名判定
    if (pre_m5 && (getenv("DS4_METAL_DISABLE_PRE_M5_DECODE_PORTS") != NULL ||
                   (pre_m5_disable_env && getenv(pre_m5_disable_env) != NULL)))
        return false;                       // 聚合/单项回滚开关
    return pre_m5 || ds4_gpu_device_is_m5_apple_silicon();
}

留下的 DS4_METAL_DISABLE_PRE_M5_* / DS4_METAL_DISABLE_M5_* 全部是回滚控制（A/B 基准的 --candidate-env 正是拿它们当对照组），不是 opt-in。5186e2c 又给非 Apple 构建补了 ds4_gpu_device_is_m5/pre_m5_apple_silicon() 的返回 0 内联 stub，让 CUDA/ROCm-only 树照常编译。

性能进程（官方 ds4-bench 对比契约：Promessi sposi、ctx 2048、32 贪心 token、预热权重；M5 Max 128GB，ds4f-q2）：

里程碑	decode 中位	稳态
基线（融合全关）	39.39	39.60
七融合进 M5（`0e125e6`）	40.54	—
round 3（pair+quad / emit-finalize / softmax+product）	40.96	41.21
round 4（pool / HC cluster / packed FA）	42.76	42.97
round 5（`0e8fb54`，"past 45"）	45.25	45.47

移植回 pre-M5（e6a6ab5，M3 Ultra）：官方交替对比 decode 41.08 → 44.18 中位（+7.55%）；聚合回滚 harness 40.65 → 44.06（+8.40%）。> README 上的 q2 全曲线表（speed-bench/m5_max.csv，每前沿 128 token 生成）仍是另一份契约的 ~39.4 t/s @2k——两表口径不同，别混读。

MXFP4 exact 加速（f75a065 / bcd1e57，M3 Ultra 15 轮：prefill +10.9%、decode +9.5%）：decode MoE 静态行程计数、prefill pair-tile half 域缩放（E8M0 字节 103-142 内 half 乘法精确）、KV RoPE 尾折进 FP8 存储、逆 RoPE 折进 FA reduce，加上 §MXFP4→half 查找表（decode 步 +4%）与 host 侧管线查找 memoization。

长上下文正确性修复（0658dac）——两个独立 bug：

全局 consumed 标志串味：注意力逆 RoPE 融合曾用进程级 ds4_gpu_decode_attn_rope_fuse_used() 判断"本层是否已折进 reduce 内核"。一次 gathered-attention 层置位后，紧跟的 indexed-attention 层把这个陈旧标志误认成自己的，跳过了它必需的独立逆 RoPE。修复：改用局部 attn_inv_rope_fuse_armed 逐层跟踪。
屏障类型不足：kernel_dsv4_kv_rope_fp8_store_f32 里 RoPE helper 写的是设备内存、却被其它 lane 读取，原来的 threadgroup_barrier(mem_threadgroup) 不够——改成 mem_device_and_threadgroup 栅栏。

stream512 top-k 移除（023614e）：kernel_topk_stream512 的原子 append/紧缩路径在宽恢复 prefill 行（长上下文、top_k=512、n_comp>1024、n_tokens≥32）上输出顺序不可复现——即 top-k 选择不确定。整条路径删除（含 DS4_METAL_ENABLE_STREAMING_TOPK512），永远走确定性的 argsort+merge 路径。这是"快但不确定 = 坏"的直接教材：宁可慢也要可复现。

负面结果同样入档（bfeb61d 等）：DS4_METAL_Q8_MV_ROWS=4 看着像 43.7 t/s，实为跳行损坏（相关内核编译期 NR0=2）；M5 上 NAX n_tok=1 matmul 不满足逐位精确（拒绝）；各种 split 调度 2/0、8/16… 全在 ±0.15% 内（4/0 最优）；MTP 草稿深度 1-4 验证成本超过接受收益（28.2/23.5 t/s）；关 Metal 4 让 prefill 直接减半（NAX 张量 matmul 是 prefill 主力）。且只信进程内/交替对比——相隔时间窗的机器漂移可达 ~10%。

4o. 视觉编码器 GPU 后端（2026-09）

多模态视觉的编码器在三个后端都有实现（语言模型侧的管线与架构见 Part 3 §7）：

Metal：metal/glm53_vision.metal + ds4_metal.m 桥接
CUDA：ds4_glm53_vision_gpu.cuh / ds4_deepseek4_vision_gpu.cuh——两族共享同一组 header-only 内核；767e517 让视觉权重在 DGX Spark 这类集成 CUDA 系统上直接映射（ds4_gpu_set_aux_model_map_range，零拷贝）
ROCm：rocm/ds4_rocm_deepseek4_vision.cuh（Vision-Exp，e976a54）；GLM 5.3 视觉走 ds4_rocm_glm.cuh

编码入口统一走 ds4_engine_vision_encode_image() 的 vision_kind 分发（DS4_VISION_GLM53 / DS4_VISION_DEEPSEEK4），产出 ds4_vision_embedding（token_count × 4096 浮点 + 网格维度 + 指纹）。DeepSeek 侧的 visual_router_bias 是张量级的路由干预：ds4_gpu_router_select_batch_visual_tensor 对 batch 里的视觉 token 施加专门偏置。QA 用 tests/vision-fixtures/（两族各带图集 + 用例）与 tests/test_glm53_vision_engine.c / test_deepseek4_vision_image.c 做快照回放。

4p. Metal Prefill 战役（MXFP4 路由 GEMM + 平铺位置表，2026-09）

§4n 的 decode 融合战役收官后，同一套"测量驱动 + 可回滚"方法论转向 prefill。四个独立抓手：

1. MXFP4 路由专家上 TensorOps（2240272）。路由专家的 prefill GEMM 不再逐 token matvec，而是实例化现成的 mul_mm_id（分组矩阵乘）TensorOps 管线——metal/moe.metal 新增 kernel_mul_mm_id_mxfp4_f32_mpp / _f16_mpp，由 ds4_metal.m 的 ds4_gpu_get_mul_mm_id_pipeline() 按激活精度选择。数学相同、累加顺序不同（矩阵核 vs 逐行），所以 frontier logits 不再逐位一致——这正是 §4n "exact 合同"只签给 decode 的原因：prefill 输出只作为后续计算的输入，不直接当结果比对。DS4_METAL_DISABLE_MXFP4_MM_ID_MPP=1 回退。

2. 位置表平铺（8cbdd76）。MLA 压缩器的 APE（位置）表展开此前按 token 分发单 threadgroup 编码器——一个 2048-token prefill 里 25K 个编码器 dispatch 中 21.5K 个是在干这个。新内核 kernel_cpy_tile_rows_f16_f32 / _f32_f32（metal/cpy.metal，host 侧 ds4_gpu_encode_cpy_ape_tile_rows）把整段表拷贝折叠成一次平铺行 dispatch。

3. 路由激活复用（2ea3839 / 5806d09）。同一专家批内复用 gate/up 激活，MXFP4 激活打包进两 rank TP prefill（DS4_METAL_DISABLE_ROUTED_MPP_PACKED）；配套 tiny-pair / routed-pair 的 SwiGLU 融合（各有 DS4_METAL_DISABLE_* 回滚）。§4j 双 Mac TP 的 prefill ~850 t/s @2048 有这两笔的直接贡献。

4. 双缓冲 dense 内核的正确性修复（8fcd61d）。TensorOps 双缓冲 dense 内核需要 8 KiB threadgroup 内存（2u * 64u * 32u * sizeof(uint16_t)），host 侧分配却按 4 KiB 给——undersized。回归 make test-metal-dense-mpp（tests/test_metal_dense_mpp.c，在 MTL_DEBUG_LAYER=1 下跑——Metal 调试层会替你抓住 threadgroup 超配）；另有 make test-metal-moe-prefill（tests/test_metal_moe_prefill.c，285 行）。

单机 M5 Max prefill（上游 docs/PERFORMANCE.md）：790 t/s @2048 ctx、573 @16384、399 @65536。

4q. DeepSeek V4.1 Flash Metal 后端（Engram 磁盘表 + 宽批 prefill + 原生会话解码，2026-09）

bd66c40（2026-09-12）把 V4.1 Flash 带进 Metal——不是给旧图加分支，而是新开一条图（ds41_gpu_graph）+ 一个新内核源文件 metal/dsv41.metal（kernel_dsv41_rope、kernel_dsv41_bf16_linear/_quantize（BF16/E8M0/E4M3 激活舍入）、kernel_dsv41_engram_add、kernel_dsv41_pool2、kernel_dsv41_candidate_blocks/_filter、kernel_dsv41_carry_copy、kernel_dsv41_indexer_pack/_scores_packed，因果版 argsort 变体 kernel_argsort_*_causal[_shuffle/_merge]，打包 MoE GEMM kernel_mul_mm_id_{iq2_xxs,q2_K}_mpp_packed_m32n128 + kernel_moe_packed_offsets 头部通道）。架构动机与 Engram 机制见 Part 3 §8 与 Engram 术语表；这里讲执行侧的四个故事。

1. 宽批 SSD prefill：算读重叠。V4.1 的层主序 prefill sweep（ds41_graph_prefill_sweep）把计算深度设为 1 的双缓冲做到极致：算第 il 层时就用 pthread 页面工作线程预读第 il+1 层（metal_graph_stream_prepare_start_if_needed(..., pread_only=true)，join 后才 map）。批量上限 DS41_PREFILL_CAP = 8192 token；层间要搬的"紧凑 carry"（BF16 残差、pre、ffn_split、selected_comp、位打包 block_mask）约 3 GiB，由 kernel_dsv41_carry_copy 无损换页。≥16384 token 的大 prefill 还会 mlock 前 20 层的路由专家（60 个 span，借用 decode 专家缓存预算，ds41_encoder_acquire/release），扫完后用最后映射的层回填 decode 专家缓存（ds41_prefill_seed）。Engram 行的预取独立成军：所有行号先一次性哈希进 prefill_ids，pthread 预取线程每 2048 token 发布一批（原子计数器），第 2 张表在第 2 层才开读——正好被编码器计算掩盖。

2. 原生会话解码 2–8 路。ds41_sessions_batch_supported 收 2 到 DS4_TP_BATCH_MAX_ROWS（=8）个 session，条件苛刻：各自有有效检查点、pos == checkpoint.len、非分布式、非 SSD streaming、非 quality/imatrix 采集，且 image_count == 0（原生批处理图没有逐行 Engram 抑制逻辑，图像位置只有逐行路径实现了 image_text_mask/ds41_image_at）。能原生批的关键是逐位对齐：批处理 decode 复用 §4n 的 tiny-pair 融合内核（IQ2_XXS/Q2_K 与 Q4_K 两套配方都要"在 BF16/router 边界前保持标量归约顺序"），ds4_gpu_dsv41_projection_rows 保持"与逐行 matvec 完全相同的算术"。TP 模式下 3–8 路才原生、2 路退有序兜底——两行摊不平批量交换与投影的固定开销；resident 单机从 2 路起原生。

3. 数值契约的边界。和 §4p 一样，"exact 合同"只签给 decode：批处理 decode 逐位对齐标量。但层主序 prefill 的分组 MPP GEMM 改变归约顺序——QA §17 记录了已知差异：Q4 batched prefill 在短官方续写测试上概率分数略降（resident 0.25241 vs SSD 0.24652），关掉层主序 prefill 能逐字节复现 SSD 分数；top-token 总数不变但上游明说"这不能证明等价"。另有 TP 注意力输出：每 rank 发未舍入的部分和、图侧求和后才做 BF16 舍入。RoPE 内核包在 #pragma float_control(precise, on) 里、频率由 CPU 预算——"频率舍入误差在长上下文累积成相位误差"。

4. 内存与容量。Q2 主权重 152 GiB（单机 SSD streaming / 双 128 GB Mac TP 每 rank ~81 GiB，两台都要完整 GGUF）；Q4 294 GiB（SSD streaming 或 512 GB M3 Ultra 常驻——QA 跑到 130,062 token）。Engram 表任何模式都不常驻。Metal 侧新增缓冲：打包 RHS 上限提到 512 MiB（V4.1 形状专用，带 offsets 头部通道）；每图 Engram 缓冲 engram_rows（6144 floats）+ engram_prefetch（prefill_cap×6144）；会话内存经 ds41_session_bytes/ds41_memory_admit 统一准入。配套修复：Q4 间接专家表（q4_expert_table 等全家）与有界 SSD 缓存互斥——间接表保留全部专家驻留集，装不进有界缓存；流式 C 调用路径没有外层 @autoreleasepool，blit 簿记会跨 token 累积保留 slab——ds4_gpu_tensor_copy、TP gate encode、专家缓存种子等入口都补了池。

QA §17 的教训也值得记：Q4 首次 SSD 跑把一台 M5 直接重启了（92 秒 watchdog 超时，123.4 GiB wired vs 计划 76.69 GiB——PRO/Q4 间接专家表把整层驻留了）——修复成只驻留选中专家的缓存；Q4 8 行 decode 曾超 2e-4 容差（0.1419），靠标量专家归约修回、容差不动；512 token 的"热 SSD 批处理阈值"提案因长 NLL 回归被否。实测（512 GiB M3 Ultra，Q4）：prefill 342–716 t/s、decode ~18.1–18.7 t/s。

发布后两个跟进（2026-09 中旬）：① decode 用上并发 Engram reader（6c00e2d）：ds4_engram_read_batch 本就并发扇出行读取，但 256 行的阈值只有批量 prefill 够得着——普通 decode 的 48 个串行 pread 全都排在 GPU 开工之前（GPU 干等，串行 107.9 µs/行 vs 并行 12.8）。修复把 decode 读路由进批量入口、阈值降到 8（Apple 限定，非 Apple 路径每批建线程不划算——077a257），M3 Ultra Q2 常驻 decode 19.16 → 21.06 t/s（+9.9%）、输出逐字节相同。行复用率只有 22.5%——赢的是并发不是缓存。② 路由 softplus 小 logit 精度（0719a0b）：fp32 里 log(1+exp(x)) 在 exp(x) < 1/32 时把小 exp 加进 1.0 会丢位，而路由门控喂的是专家 top-k 选择——小误差翻转路由。四个 DeepSeek 路由变换点改用四项 log1p 级数（多项式输入 clamp 保 SIMD 未选中车道有限），与共享精确参考比对不再放宽容差；配套 112 条官方 V4.1 续写的路由质量对照（默认 NLL 0.396840→0.395353）。

4r. DeepSeek V4.1 Flash CUDA 后端（有界专家缓存 + 双 Spark TP，2026-09）

a04f46f 把 V4.1 文本推理带上 CUDA（单 Spark Q2 SSD streaming + 双 Spark常驻 TP；视觉/DSpark/ROCm 不支持）。与 Metal 路径（§4q）的分工很有代表性——同一个模型，两种存储哲学：

	Metal SSD streaming	CUDA SSD streaming
权重映射	`MAP_SHARED` 共享 mmap，缺页即算	有界 device 专家缓存：专家按需搬进固定预算的显存缓存
prefill 重叠	页面工作线程预读下一层 mmap	恰好两个 8 MiB pinned staging 缓冲 + 一个读线程，读下一层专家进"预留槽"（无第二专家 arena）
读入时机	与计算流水	当前层计算时预读下一层；≥2K token 才开（更短的 prefill 全层读反而慢，`DS4_CUDA_DISABLE_SSD_PREFETCH` 回退）
预留槽可见性	—	拷贝全部完成才发布；未用上的预读保持冷（在需求热专家之前先被逐出）

要点：短 prefill 走精确路径（< 256 token 矩阵阈值以下用精确的 2–8 行 chunk，与标量逐位一致）；中等追加一次层扫（e9e1aaa：带可观尾 tile 的追加复用 carry 工作区、每层只拜访一次而不是两次，保留 2048 行算术分区；DS4_CUDA_DISABLE_SSD_MEDIUM_SWEEP 回退——3,241-token README prefill ~55 → ~98 t/s）；8 行 decode 批处理（mixed prefill+decode 行支持，更大混合走有序）。Engram 磁盘读与 Metal 同款（表永远在盘上）。

双 Spark 网络 TP 的传输与协议细节见 §4j 尾部（RoCEv2 host staging、协议 14）；性能数字见 §4b。QA 记录强调：CUDA 的 SSD 常驻策略是"缓存预算换 prefill 吞吐"——自动缓存下 decode 微降 ~5%，但整段请求时长显著缩短。

4s. Qwen3.8 Flash Next GPU 图与批处理战役（2026-09）

Qwen3.8 的 GPU 图（ds4.c 57,547 起，Metal 源 metal/qwen4.metal、CUDA 源 ds4_qwen4_cuda.cuh ~2.3K 行）是一条图服务 decode（T=1）与批量 prefill：每次 forward 一个 command batch，f32 瞬态按 cap_tokens 行定尺寸，同一套内核跑 1 行或 8192 行。内存契约值得背下来：

瞬态（DS4_QWEN4_SCRATCH_FIELDS）每次 forward 重算、可以共享；其余全是每会话私有——KV、索引器与块键缓存、GDN/PLE 递归历史、MTP 快照、rope 位置、logits。共享只在"同一时刻只跑一个 forward"下成立——批处理必须是"多行一次 forward"，不能是"多个 forward 共享 arena"（第二个会话会在 GPU 读走之前覆写第一个的瞬态）。

原生跨会话批处理战役（ba3b0c2 起一族提交）把 Qwen 从"每 session 整图跑一遍"（16 个 session = 16 倍权重流量）变成：dense 读权重的活（HC 混合、投影、MoE、输出头）对堆叠的行整批跑一次；递归状态（GDN scan）、注意力缓存、PLE 历史按 session 私有，经行表（ds4_gpu_qwen4_gdn_row / _attn_row：{state, history, 快照 GPU 地址, 首行, 行数}）寻址——从每周期 ~2,300 个短 dispatch 收敛到表寻址的批量内核。会话 prefill 瞬态并进共享 arena（每 session 内存 ÷3.6 @ctx4096）。配套：dense GEMM 窄批 regime 的 split-K 路径、dense/MoE/混合 GEMM 阈值从"T=1 或 T≥8192"改服役 4..64 行 decode 批、Q8 批 GEMM 以 16 行半批服务 9..32 行。C=16 @ctx4096 M5 Max：~50 → ~143 t/s 聚合。浮点归约顺序在批内核处改变（与 §4q 的数值契约同款立场：decode 不承诺逐位、只承诺行为一致）。

推测批处理（a504136→ab38bc6）把 Part 4 §5 的 MTP 扩到批：每 session 的 token 作行 0、待验证草稿搭车作行 1，预测器层对整批跑一遍（trunk 流复用批 R 的行位置、每 session 最后提交行 gather 后一次批量输出头 GEMV + 逐 session argmax）；拒绝恢复行快照。全部机制与策略（all-or-nothing 起草、EMA 定价、16 session 上限）见 Part 5 §6。

盘上 n-gram 的并发读（d0b7434 + 077a257）：PLE 表不缓存（engram 同款立场），每 token 16 行 × 320 B 的 pread 串行读要 1.3 ms——19 ms decode 步的 7%、推测周期的 10%，全卡在 GPU 开工前。≥2 行的请求在 Apple 上走 dispatch_apply_f 并发 reader（非 Apple 保持 256 行阈值 + pthread——避免 CUDA/ROCm 每批建线程）；C=1 普通 51.3 → 54.5 t/s、MTP 代码 69.9 → 78.9、散文 58.3 → 66.0，logits 逐位相同；M5 Max 暖盘 decode 12.62 → 13.39 t/s。

CUDA 侧（df6eedd 起一族）：Qwen 优化内核覆盖投影与专家（分组查询头共享的 prefill 注意力、跨注入 chunk 复用的 prefill 归一化、对 CPU 参考逐核回归），Spark 上与 Metal 同一套命令（无转换、无特性开关）；CUDA 目前按会话有序 decode（原生跨会话批处理是 Metal 专属）。TP/流水线/SSD expert streaming/ROCm 均未实现。

5. 增量吞吐量基准测试（ds4-bench）

传统基准测试报告全流程平均速度，ds4-bench 在不同 context frontier 处测量瞬时吞吐量：

每个 frontier（如 2048, 4096, 6144...）只计算新增 token 区间的 prefill 速度
prefill 后保存内存快照，做固定 128 token 贪心解码测 generation 速度
快照保存/恢复时间不计入测量窗口
输出 CSV：每个 frontier 的 prefill t/s、generation t/s、kvcache_bytes

// ds4.h — 快照 API
typedef void (*ds4_session_progress_fn)(void *ud, int progress);
void ds4_session_set_progress(ds4_session *s, ds4_session_progress_fn fn, void *ud);

void *ds4_session_save_snapshot(ds4_session *s);
int ds4_session_load_snapshot(ds4_session *s, void *snapshot);
void ds4_session_snapshot_free(void *snapshot);

测试数据使用 speed-bench/promessi_sposi.txt（23329 行意大利语公共领域文本）作为固定输入序列。bench/ 目录已重命名为 speed-bench/，新增 M2 Ultra (192 GB) 和 M4 Max 基准数据，以及 plot_speed.py CSV→SVG 可视化脚本。

2026-08 又用 ds4-bench（Promessi sposi 输入、2048-token 阶梯、128 贪心 token）重测了 M5 Max（Metal）与 DGX Spark GB10（CUDA）的 q2 全曲线，存于 speed-bench/m5_max.csv / gb10.csv：M5 Max 2k 上下文 prefill ~790 t/s、generation ~39 t/s（65k 上下文降到 ~399 / ~27.6）；GB10 prefill ~826 t/s、generation ~18 t/s（65k ~823 / ~13.8），CUDA prefill 几乎不随上下文衰减。

同月 M5 精确解码融合战役（§4n）用另一份对比契约（ctx 固定 2048、32 贪心 token、预热权重、交替副本取中位）把 M5 Max q2 decode 从 39.39 推到 45.25 中位 / 45.47 稳态，pre-M5（M3 Ultra）从 41.08 到 44.18。两个数字口径不同：全曲线表量的是"每个前沿的真实生成速度"，战役契约量的是"融合前后的受控对比"。

speed-bench/ 还新增两个 A/B 基准（战役的验证工具，也是此后 Metal 优化的标准方法论）：make metal-decode-schedule-bench 与 make metal-prefill-variant-bench——除非新旧路径的全词表 logits 逐位相同（decode 侧还要求 token 选择一致）否则 abort，--candidate-env 直接拿 DS4_METAL_DISABLE_* 回滚开关当对照组。

15 天学习总结

C 语言技能树

LLM 推理知识树

ds4.c 的设计哲学

垂直整合：一个文件做所有事，零外部依赖
专用化：只服务一个模型，硬编码所有参数
零拷贝：mmap + Metal/CUDA no-copy buffer，数据永远不离开映射区
热路径零分配：预分配 scratch buffer，生成循环中不 malloc
早失败：加载时严格验证，不匹配就退出
可扩展后端：通过 ds4_gpu.h 抽象接口，Metal/CUDA/ROCm/CPU 四后端可互换；融合算子作为可选后端钩子，新后端按需提供
KV cache 是一等公民：渲染文本前缀匹配 + 对齐前沿 + KTM 工具记忆，让长上下文推理跨重启存活
方向性引导：推理时低秩编辑激活空间，无需重训练即可控制输出风格
分布式推理：coordinator/worker 模型切片执行，PRO 模型可跨多台机器运行
PRO 模型支持：384 专家 Q4_K 量化路径，CPU/Metal/CUDA 三后端完整覆盖
SSD Streaming：路由专家缓存 + SSD 按需加载 + 路由热度淘汰（每 16 token 衰减、提示局部重置但缓存跨会话保温）+ 三路 prefill，让模型在有限 RAM 下也能运行
合作式中断：Ctrl+C 在安全检查点优雅停止，KV cache 状态不丢失
张量并行：Mac-to-Mac（RDMA over Thunderbolt）与 CUDA（--cuda-tensor-parallel）两种形态，两机同 token 同时切层内工作量，在同步门交换部分和，降低单 token 延迟
多 GPU 放置：--gpu-vram/--gpu-devices 统一参数 + 层打包器单调连续放置 + 跨设备 P2P 拷贝，把模型按预算分到多卡
会话批处理：--batched-session N 预分配 N 个常驻 KV session，就绪 decode 步打包进同一内核，原生批处理 + 有序精确兜底
多模型：不再 DeepSeek 专用——GLM 5.2 在 Metal/CUDA/ROCm 三后端运行，路由专家支持多种量化布局（详见 Part 3）

Part 6: GPU 加速 ​

1. GPU 后端 + 总结 ​

C 知识点 ​

1. Objective-C 互操作 ​

2. Metal 计算管线 ​

可选 buffer 绑定占位 ​

张量句柄的双重释放守卫 ​

3. Metal Shading Language ​

Metal Neural Acceleration (NAX) ​

Compressed KV in F16 ​

GPU 功耗节流 ​

Typed Pointer vs void Pointer ​

LLM 知识点 ​

1. GPU 加速原理 ​

2. Flash Attention ​

3. 零拷贝 MTLBuffer ​

模型视图与张量覆盖不变量 ​

4. 分布式推理（Distributed Inference） ​

4a. SSD Streaming（SSD 流式推理） ​

mlock pinning（让缓存专家驻留 RAM） ​

Metal 4 prefill 加速 ​

GLM 全量检查点 streaming 成熟化（2026-09） ​

4b. DGX Spark / GB10 后端（HBM 缓存） ​

Metal View Cap 修复 ​

模型缓存 FD 绑定 ​

2026-09：常驻会话 + 小 batch 对齐内核 ​

2026-09 中旬：V4.1 与 Qwen3.8 登陆 Spark ​

4c. Strix Halo / ROCm 后端 ​

4d. 帮助系统（Help System） ​

4e. 模型下载脚本更新 ​

4f. 方向性激活引导（Directional Steering） ​

4g. CUDA Q8→F16 缓存显存预算守卫 ​

4h. CUDA 长上下文修复 ​

MoE Down 路径清理 ​

CUDA 内存优化 ​

CUDA Managed KV Cache ​

Metal Q4 Model View 回退 ​

4i. CUDA Prefill 性能优化 ​

4j. 张量并行（Tensor Parallelism） ​

双 Mac TP 延迟工程战役（2026-09，4886c56 合入 macm5us/mac-tp-1-sep-2026） ​

4k. 多 GPU 放置与层打包（Multi-GPU Placement） ​

4l. GLM 5.2 / 5.3 后端支持 ​

4m. CUDA 性能（vendored MMQ + decode-island graphs + Blackwell） ​

4n. Metal 精确解码融合战役（Exact Decode Fusions） ​

4o. 视觉编码器 GPU 后端（2026-09） ​

4p. Metal Prefill 战役（MXFP4 路由 GEMM + 平铺位置表，2026-09） ​

4q. DeepSeek V4.1 Flash Metal 后端（Engram 磁盘表 + 宽批 prefill + 原生会话解码，2026-09） ​

4r. DeepSeek V4.1 Flash CUDA 后端（有界专家缓存 + 双 Spark TP，2026-09） ​

4s. Qwen3.8 Flash Next GPU 图与批处理战役（2026-09） ​

5. 增量吞吐量基准测试（ds4-bench） ​

5. 增量吞吐量基准测试（ds4-bench） ​

15 天学习总结 ​

C 语言技能树 ​

LLM 推理知识树 ​

ds4.c 的设计哲学 ​

反向链接

出链

Part 6: GPU 加速

1. GPU 后端 + 总结

C 知识点

1. Objective-C 互操作

2. Metal 计算管线

可选 buffer 绑定占位

张量句柄的双重释放守卫

3. Metal Shading Language

Metal Neural Acceleration (NAX)

Compressed KV in F16

GPU 功耗节流

Typed Pointer vs void Pointer

LLM 知识点

1. GPU 加速原理

2. Flash Attention

3. 零拷贝 MTLBuffer

模型视图与张量覆盖不变量

4. 分布式推理（Distributed Inference）

4a. SSD Streaming（SSD 流式推理）

mlock pinning（让缓存专家驻留 RAM）

Metal 4 prefill 加速

GLM 全量检查点 streaming 成熟化（2026-09）

4b. DGX Spark / GB10 后端（HBM 缓存）

Metal View Cap 修复

模型缓存 FD 绑定

2026-09：常驻会话 + 小 batch 对齐内核

2026-09 中旬：V4.1 与 Qwen3.8 登陆 Spark

4c. Strix Halo / ROCm 后端

4d. 帮助系统（Help System）

4e. 模型下载脚本更新

4f. 方向性激活引导（Directional Steering）

4g. CUDA Q8→F16 缓存显存预算守卫

4h. CUDA 长上下文修复

MoE Down 路径清理

CUDA 内存优化

CUDA Managed KV Cache

Metal Q4 Model View 回退

4i. CUDA Prefill 性能优化

4j. 张量并行（Tensor Parallelism）

双 Mac TP 延迟工程战役（2026-09，`4886c56` 合入 `macm5us/mac-tp-1-sep-2026`）

4k. 多 GPU 放置与层打包（Multi-GPU Placement）

4l. GLM 5.2 / 5.3 后端支持

4m. CUDA 性能（vendored MMQ + decode-island graphs + Blackwell）

4n. Metal 精确解码融合战役（Exact Decode Fusions）

4o. 视觉编码器 GPU 后端（2026-09）

4p. Metal Prefill 战役（MXFP4 路由 GEMM + 平铺位置表，2026-09）

4q. DeepSeek V4.1 Flash Metal 后端（Engram 磁盘表 + 宽批 prefill + 原生会话解码，2026-09）

4r. DeepSeek V4.1 Flash CUDA 后端（有界专家缓存 + 双 Spark TP，2026-09）

4s. Qwen3.8 Flash Next GPU 图与批处理战役（2026-09）

5. 增量吞吐量基准测试（ds4-bench）

5. 增量吞吐量基准测试（ds4-bench）

15 天学习总结

C 语言技能树

LLM 推理知识树

ds4.c 的设计哲学