Misc. bug: Vulkan ARGSORT ne=[2048,1,1,1] only sorts half of the array on some devices

快速结论：这个报错通常出现在支持较大 workgroup（例如 2048 invocations）的 Vulkan 设备上，运行 llama.cpp 的 ARGSORT 后端测试时，数组只被排序了一半。优先排查 ggml_vk_argsort() 中 use_small 路径是否在 ne0=2048 时被错误启用。

适用环境：问题在 Linux aarch64、llama.cpp 0.5.0-dev（build 11188，commit e85e15cf6）上报告，使用 Vulkan 后端，设备为 Adreno A750（maxComputeWorkGroupInvocations=2048）。桌面 RDNA3（RX 7900 XTX、Ryzen 7 7800X3D iGPU，workgroup 上限 1024）无法复现。

最快修复方案：暂无确认的一步修复方案。Issue 被关闭时没有合入修复补丁；讨论中给出的建议是修改 ggml_vk_argsort()，让 use_small 使用未 clamp 的索引，即要求 ncols_pad_log2 == pipeline_idx，但这只是建议，尚未验证。

注意事项：上述改动属于源码级修改，需要重新编译 llama.cpp，且必须用目标 Vulkan 设备回归测试。受影响的窗口是 ne0 在 (1024, 2048] 之间、且设备声明最大 2048 invocations 的场景；workgroup 上限为 1024 的设备会走 large 路径，不受影响。

问题场景

在 llama.cpp 的 Vulkan 后端测试中运行 test-backend-ops，单独测试 ARGSORT 算子：ne=[2048,1,1,1]、数据类型 f32、order=0。在 Adreno A750 这类支持 2048 个 workgroup invocations 的设备上，测试失败，结果只有数组的一半被正确排序。

报错原文

[ARGSORT] ERR = 0.999998560 > 0.000000100   ARGSORT(type=f32,ne=[2048,1,1,1],order=0): FAIL
[ARGSORT] ERR = 0.999998601 > 0.000000100   ARGSORT(type=f32,ne=[2048,1,1,1],order=0): FAIL

原因分析

可能原因与 ggml_vk_argsort() 中的 pipeline 选择逻辑有关。评论分析指出：当 ne0 = 2048（log2 = 11）、设备 max_workgroup_size_log2 也为 11 时，pipeline_idx 会被 clamp 到 10，对应 BLOCK_SIZE 为 1024 的 pipeline_argsort_f32[10]。由于该 pipeline 存在，use_small 被置为 true，但该 pipeline 只覆盖 1024 列；同时跨 workgroup 的合并循环位于 if (!use_small) 之后而被跳过，最终在同一行上派发 2 个 1024 宽的 workgroup，导致前半部分被排序两次、后半部分未排序。

环境排查

确认 llama.cpp 版本与 commit：报告中使用 0.5.0-dev（build 11188，commit e85e15cf6）。
确认操作系统：Linux aarch64（报告环境）。
确认 Vulkan 设备及 maxComputeWorkGroupInvocations：Adreno A750 为 2048；桌面 RDNA3 为 1024 时无法复现。
确认 maxComputeWorkGroupSize：Adreno A750 为 [1024, 1024, 1024]。
确认测试的 ne 是否落在 (1024, 2048] 区间，这是复现窗口。

解决步骤

在目标 Vulkan 设备上复现：运行 ./test-backend-ops -o 'ARGSORT(type=f32,ne=[2048,1,1,1],order=0)'，确认 ARGSORT 测试 FAIL。
查看 ggml_vk_argsort() 中 pipeline_idx、num_argsort_pipelines、use_small 的判断逻辑，确认 pipeline_idx 是否因为 max_workgroup_size_log2 被 clamp 到 10。
按讨论中的建议，调整 use_small 的门控条件：改用未 clamp 的索引判断，要求 ncols_pad_log2 == pipeline_idx 时才允许走 small 路径，避免被 clamp 后的 pipeline 误入 small 路径。
重新编译 llama.cpp 后，在同一设备上重跑上述 ARGSORT 测试。
如果无法修改源码，可临时避免在该设备上触发 ne0 落在 (1024, 2048] 的 ARGSORT small 路径，或改用 workgroup 上限为 1024 的设备/驱动配置。

验证方法

重新运行 ./test-backend-ops -o 'ARGSORT(type=f32,ne=[2048,1,1,1],order=0)'，确认 ARGSORT 测试由 FAIL 变为 PASS，且不再出现 [ARGSORT] ERR = ... > ... 的报错。评论中桌面 RDNA3 设备（workgroup 上限 1024）在同一命令下均 PASS，可作为对照。