设计决策
本文档只保留仍然直接影响 CUDA 实现或公开集成面的核心 ADR,删除了与外部流程框架绑定的历史治理层。
ADR-1:head_dim 仅支持 32 / 64 / 128
背景
Kernel 的 tiling、共享内存占用和寄存器压力都依赖编译期常量。
决策
仅支持 32、64、128 三个 head_dim,并通过显式模板实例化进行分发。
影响
- 共享内存与启动参数更可预测
- 测试面更集中
- 非标准
head_dim需要由上层做 padding 或适配
ADR-2:FP16 路径内部统一使用 FP32 累加
背景
Attention 分数累加和 online softmax 在半精度下容易出现精度损失与数值不稳定。
决策
GMEM 边界可使用 half,但内部累加、归一化和关键中间值统一使用 float。
影响
- 长序列下更稳定
- 内部资源压力略增
- FP16/FP32 两条路径的行为更一致
ADR-3:Python 集成停留在 C ABI 层
背景
仓库强调轻量、可嵌入和低依赖,不希望为绑定层引入大型运行时。
决策
保留 C ABI,并以 ctypes 风格的使用方式作为官方文档路径,而不是引入 pybind11 一类框架。
影响
- 依赖面更小
- 更容易被多语言复用
- 运行时类型检查不如高层绑定框架友好
ADR-4:核心线程块大小固定为 128
背景
当前 occupancy、warp 分工和 shared memory tiling 都围绕 128 线程块做了收敛。
决策
继续保持 blockDim.x = 128,并与显式启动配置校验配套。
影响
- 调度与性能特征更稳定
- 维护成本更低
- 若未来调整,需要重新调优并完整回归验证
ADR-5:仓库文档优先,移除外部流程框架
背景
仓库曾叠加多层 AI / 流程控制框架,导致说明重复、约束冲突、维护成本上升。
决策
把维护规则收敛到仓库原生文档:README、CONTRIBUTING、GitHub Actions 与 VitePress 页面;移除外部控制层。
影响
- 文档与代码更容易保持一致
- 删除了冗余流程负担
- 后续维护者只需阅读仓库本身即可理解项目