llama.cpp：支持 CDNA 稀疏注意力，调整内核选择

1. 在 CUDA 共享代码中解除 AMD 执行限制

此次更改于 10月29日合并至 llama.cpp。核心是在 CUDA 与面向 AMD GPU 的运行及开发环境 HIP 共享的后端中，移除阻止 HIP 构建使用稀疏注意力路径的条件。在保留现有 CUDA 路径选择条件的同时，将适用范围扩展至 AMD CDNA。

此前，HIP 的编译范围不包含稀疏注意力代码，收集所需位置的函数也会中止执行。更改后，HIP 也能连续完成这项准备工作和稀疏计算。内核选择规则也有所调整：CDNA 满足稀疏条件时，优先采用稀疏矩阵运算内核，而非现有的按任务大小选择内核的规则。

2. 以预先收集所需位置的开销换取计算量减少

稀疏注意力根据标示参与计算位置的掩码，将可引用的键位置收集成列表，再处理选中的键和值。它不计算被屏蔽的位置，而是仅处理所需位置，以预先生成位置列表的准备开销，换取注意力运算量的减少。

此次实现按照 GPU 中协同执行单元的大小存储位置选择结果。为容纳 CDNA GPU 中分担并行计算的 64 个执行通道，代码使用 64 位掩码，并相应调整了所选位置数量的统计处理。

3. 根据任务形态选择稀疏内核

据代码注释，CDNA 上编译稀疏支持的对象，是表示查询和键维度的 DKQ 为 576、内核列分组设置 ncols2 为 16 的变体。代码先确认能否选择该变体，再检查实际任务条件。作者将 RDNA 等其他架构的支持留待后续实现。

运行时，存在掩码、且加到计算结果上的偏置与数值上限设置均为 0 的任务，会被视为稀疏路径候选。CDNA 的键长度必须至少为 4096，且不小于稀疏列表可容纳的最大键数量的两倍。软件通过检查任务大小和形态来决定执行路径。

4. 合成测试的吞吐量与波动幅度

pp512·tg64 是测试类别，d16384·d32768 是测试深度标记。下方两列遵循原文中的测量顺序。吞吐量旁的 ± 数值也与原文一致。

表：qwen4exp 合成测试的两组测量结果

条件：PR 作者的合成测试。qwen4exp A3B Q4_K - Medium，103.68 GiB，176.94 B，ROCm，重复 50 次。分布式加载于 2x mi100s、1x r9700、1x rx7900xtx。包含用于启用稀疏路径的独立补丁。

单位：t/s（每秒处理的 token 数）

第一次测量中，深度从 d16384 增至 d32768 时，pp512 吞吐量从 760.52 降至 633.18 t/s，tg64 则从 35.81 降至 30.78 t/s。pp512 的 ± 值也从 55.29 增至 99.63，表明在这次测量中，更深的测试吞吐量更低、波动幅度更大。

作者认为此次测试的性能提升明显。他还指出，约一半模型加载在不受此次更改影响的架构上，因此认为测量结果低估了提升效果。性能表未标明应用更改前后的对应关系，难以仅凭两列数据计算提升率。

5. 兼顾旧版 ROCm 构建

代码新增了兼容定义，使低于 ROCm 7.0 的版本使用现有的 __ballot，而非用于收集各通道选择结果的 __ballot_sync。作者提到，在 ROCm 7.0 之前，__ballot_sync 内建函数被隐藏在 HIP_ENABLE_WARP_SYNC_BUILTINS 之后。原文新增代码如下。

#if HIP_VERSION < 70000000
#define __ballot_sync(mask, var) __ballot(var)
#endif //HIP_VERSION < 70000000

作者表示，在这些情况下该内核直接使用 __ballot 即可。不过，这一兼容宏本身并没有内核范围的限定，而是定义在共享 HIP 头文件 ggml/src/ggml-cuda/vendors/hip.h 中，条件为 HIP_VERSION < 70000000。同时他提醒了一个尚未解决的问题：另一个文件中也在使用 __ballot_sync，而那里在 wave64 设备上会实际使用掩码，因此该用法在 ROCm 6.3-6.4 的 wave64 上并不正确。作者表示会就这一用法作后续处理，但该后续提交并未包含在本 PR 中，因此上述兼容处理并未解决这个仍然存在的正确性问题。

此次更改的核心，是在 CDNA 上开通仅计算所需位置的路径，并根据任务条件选择相应内核。