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 中的流程：

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 会报校验错误：

// 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/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：

1. Indexer Scores NAX — 压缩 KV 注意力的索引评分，将评分计算卸载到 Neural Accelerator
2. 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 反映正确的只读访问语义：

// 改前：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，可压缩为更低精度减少网络开销）。

# 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）。

命令行选项：

选项	说明
--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);

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）：

1. ROCm 安装补全：Ubuntu 26.04 的 librocwmma-dev 缺 rocwmma/internal/ 头文件，需手动 clone rocWMMA 7.1.0 补齐；若 ROCm 装在 /usr 而工具链期望 /opt/rocm，要建符号链接。

2. 显存扩展（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）。

3. 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）。

更安全的 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 这类新后端能选择性提供融合内核而不必复制所有分支逻辑。

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

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 使回答更简洁

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（~2.2K 行）。它有两种部署形态，共用同一套 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 不能与 SSD streaming、分布式模式、MTP drafting、CPU 后端同时启用（ds4_tp_validate_engine_options()）。GLM 5.2 在 CUDA 上走普通层放置（非 TP）；DGX Spark 是单卡，不能用 --cuda-tensor-parallel。

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 后端支持

DwarfStar 不再是 DeepSeek 专用——现在也支持 GLM 5.2（详见 Part 3 §GLM 5.2）。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
• 限制：不支持方向性引导、--power < 100、显式 --prefill-chunk、外部 --mtp 文件；MTP 块在主 GGUF 内部，用 --glm-mtp 启用实验性贪心推测
• SSD streaming：GLM 也支持（Metal + ROCm），路径保留能装下的最大全层前缀常驻，余量给动态专家缓存
• 两机张量并行：要求 ownership-aware 的 IQ2_XXS 或 Q2_K 路由布局；路由 Q4 的 GLM 必须在评估前拒绝

ROCm 侧为 GLM 做了大量调优（rocm/ds4_rocm_glm.cuh ~4.2K 行）：selected/causal 注意力 GEMM、wave value projection、slice decode、routed IQ2 down 等，多数是 opt-in 实验路径，验证有效后才转为默认。

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 内存计入多卡放置。

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 几乎不随上下文衰减。

---

15 天学习总结

C 语言技能树

LLM 推理知识树

ds4.c 的设计哲学

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