llama.cpp b12208：Metal与CUDA的radix Top-K并列选取改为确定性，M5 Ultra上提速约 20–45%
llama.cpp发布b12208版本，ggml的radix Top-K不再用原子计数器收集选中索引，并列候选的排序和入选子集不再随线程调度变化；提交说明称Metal上输出跨运行字节一致，M5 Ultra上top-k耗时降低约 20–45%。
AI解读：llama.cpp b12208的主要变更是ggml的radix Top-K内核：原先用原子计数器收集被选中的索引，并列候选的顺序、以及并列数超过剩余名额时最终入选的子集，都取决于线程调度。
Metal后端改为确定性方案：在连续线程块内统计高于阈值和等于阈值的元素，用simdgroup scan加一次对radix直方图的串行扫描排名，按列顺序只写一次每个保留元素；提交说明称输出跨运行字节一致。
性能方面，提交说明称Metal上该改动比原子版本更快，strided原子循环改为连续访问，M5 Ultra上top-k时间降低约 20–45%，并列密集的形状在几个百分点以内。
后续提交将Metal紧凑化改为每个warp处理行内连续块、lane以 32 步长交错访问，以恢复被标量加载拖慢的M2 Ultra上大ncols形状性能；并列溢出时回退到逐步跨lane排名，保留列序最小的并列索引。
测试方面，Metal与ROCm后端启用了确定性Top-K自检；ROCm因无cub::DeviceTopK，每行走radix或bitonic路径，覆盖CUDA跳过的 1 行和 2 行并列用例。CUDA侧删除了gather前无效的radix Top-K计数器重置内核。
发布页同时列出macOS、iOS、Linux、Android、Windows各后端构建包，包括CUDA 12/13、Vulkan、ROCm 10.0、OpenVINO、SYCL以及Snapdragon CPU/Adreno GPU/Hexagon NPU等，openEuler构建标注为DISABLED。
llama.cpp发布b12208版本，其中ggml的radix Top-K实现改为确定性并列选取，涉及CUDA与Metal路径。提交说明称，这一改动让Metal上的输出跨运行字节一致，并在M5 Ultra上把top-k耗时降低约 20–45%。
问题：原子计数器让并列结果取决于线程调度
原radix top-k内核用原子计数器收集被选中的索引。提交说明指出，这使阈值处并列候选的顺序，以及并列数多于剩余名额时最终入选的子集，取决于线程调度。
新方案在连续线程块内统计高于阈值和等于阈值的元素数量，用simdgroup scan和一次串行扫描radix直方图来排名，并按列顺序对每个保留元素只写一次。提交说明称输出现在跨运行字节一致。
Metal性能数据与后续调整
提交说明称该确定性版本也快于原子版本：strided原子循环变为连续访问，M5 Ultra上top-k时间降低约 20–45%，并列密集的形状在几个百分点以内。
后续提交指出，确定性紧凑化按线程连续块做标量顺序加载，M2 Ultra的内存系统难以重叠，大ncols形状的radix Top-K时间回退约 1.5–2.4 倍。改为每个warp取行内连续块、lane以 32 步长交错访问，使访问完全合并；lane按lane-major排名写出保留元素，只有并列溢出时才回退到逐步跨lane排名，以保留列序最小的并列索引。
- Metal后端启用确定性top-k自检。
- ROCm也启用该自检：HIP无cub::DeviceTopK，每行走radix或bitonic路径，覆盖CUDA跳过的 1 行、2 行并列用例。
- CUDA侧删除了gather前的一个无效radix top-k计数器重置内核，因top_k_radix_init已对每块行清零greater_count。
范围与发布包
提交记录显示改动范围经过收窄：中间版本曾关闭Metal/ROCm路径并对Vulkan跳过，最终Metal与ROCm的确定性检查落地。提交说明还提到SYCL结果看起来也非确定，因此采用allowlist方式。
发布页列出各平台构建包，包括macOS arm64/x64、iOS XCFramework、Ubuntu（CPU、Vulkan、CUDA 12/13、ROCm 10.0、OpenVINO、SYCL FP32/FP16）、Android、Windows（CPU、CUDA 12/13、Vulkan、OpenVINO、SYCL、ROCm 10.0）及Snapdragon的CPU、Adreno GPU、Hexagon NPU组合。openEuler构建标注为DISABLED。