b12208

<details open> ggml : update CUDA/Metal radix Top-K to make tied selections deterministic (#31084) * ggml : fix for the radix TOP_K thread timing bug * make the key an ordered 64-bit pair * a rewrite without shift-skip logic * fix casts * warp-contiguous layout * reduce the scope removing vulkan+metal changes * turn it off for Metal/ROCm, note about cub::DeviceTopK * fix the hip build * skip Vulkan as well * popc lambda from fattn * do block_cols math in int64_t * make it allowlist bc SYCL looks non deterministic too * ggml : make Metal radix Top-K tied selections deterministic The radix top-k kernel gathered the selected indices with atomic counters, so the order of the tied candidates at the threshold (and, with more ties than remaining slots, the subset itself) depended on the thread scheduling. Replace the two atomic compaction loops with a deterministic scheme: count the elements above and tied with the threshold in contiguous per-thread chunks, rank them with a simdgroup scan and a serial scan over the (now unused) radix histogram, and write each kept element exactly once in column order. The output is now byte-identical across runs. This is also faster than the atomic version: the strided atomic loops become contiguous passes, reducing top-k times by ~20-45% on M5 Ultra, with tie-heavy shapes within a few percent. Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB) * tests : enable deterministic top-k check on Metal, add radix-path tie cases Metal now produces deterministic tied selections for the radix top-k path, so the determinism self-check runs for the MTL backend as well. Add eval cases covering the radix path (ncols > 2048) with more ties at the threshold than remaining slots, and tie-heavy perf cases to monitor the cost of the deterministic tie handling. Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB) * ggml : use coalesced accesses in the Metal radix Top-K compaction The deterministic compaction walked contiguous per-thread chunks with scalar sequential loads, which the memory system of the M2 Ultra can not overlap well: the radix Top-K times regressed by ~1.5-2.4x on the large-ncols shapes. Give each warp a contiguous chunk of the row instead and let the lanes interleave within it with stride 32, so every access is fully coalesced. The lanes write their kept elements at lane-major ranks, which is deterministic, and only a tie overflow falls back to a per-step cross-lane ranking to keep the lowest tied indices in column order. Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB) * tests : remove the extra top-k tie and perf cases The radix-path eval and tie-heavy perf cases were added to validate and benchmark the Metal deterministic tie handling. Keep only the enabling of the determinism self-check for the MTL backend. Assisted-by: dsh:llama.cpp/ggml-org/GLM-5.3-Flash-GGUF (Mac Studio, M5 Ultra, 256GB) * metal: scan the radix bins and warp totals with a simdgroup The radix top-k selected the threshold bucket with a serial scan of the 256 bins by thread 0 in each of the 4 passes, and ranked the warps of the compaction with two serial scans by thread 0. The first simdgroup now does both: each lane sums 8 bins, an exclusive prefix sum and a ballot find the lane that reaches the k-th value, and the warp totals get their exclusive bases from simd_prefix_exclusive_sum. * tests: enable the deterministic top-k check on ROCm HIP has no cub::DeviceTopK, so every row goes through the radix or bitonic path, which break ties in the same way in each run. The check covers all tie cases on ROCm, including the 1 and 2 row ones that CUDA skips. * cuda: remove the redundant radix top-k counter reset Address review from ORippler. top_k_radix_init already zeroes greater_count for each chunk of rows and top_k_radix_select keeps it, so the reset kernel before the gather is one launch with no effect. --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> Co-authored-by: Pascal <admin@serveurperso.com> </details> **Website:** - <https://llama.app> **Attestations:** - <https://github.com/ggml-org/llama.cpp/attestations/65072510> **macOS/iOS:** - [macOS Apple Silicon (arm64)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-macos-arm64.tar.gz) - macOS Apple Silicon (arm64, KleidiAI enabled) [DISABLED](https://github.com/ggml-org/llama.cpp/pull/23780) - [macOS Intel (x64)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-macos-x64.tar.gz) - [iOS XCFramework](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-xcframework.zip) **Linux:** - [Ubuntu x64 (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-x64.tar.gz) - [Ubuntu arm64 (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-arm64.tar.gz) - [Ubuntu 24.04 s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-s390x.tar.gz) - [Ubuntu 26.04 s390x (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-26.04-s390x.tar.gz) - [Ubuntu 26.04 s390x (zDNN)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-zdnn-s390x.tar.gz) - [Ubuntu x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-vulkan-x64.tar.gz) - [Ubuntu arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-vulkan-arm64.tar.gz) - [Ubuntu x64 (CUDA 12)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-cuda-12.8-x64.tar.gz) - [CUDA 12.8 libraries](https://github.com/ggml-org/llama.cpp/releases/download/b12208/cudart-llama-b12208-bin-ubuntu-cuda-12.8-x64.tar.gz) - [Ubuntu x64 (CUDA 13)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-cuda-13.5-x64.tar.gz) - [CUDA 13.5 libraries](https://github.com/ggml-org/llama.cpp/releases/download/b12208/cudart-llama-b12208-bin-ubuntu-cuda-13.5-x64.tar.gz) - [Ubuntu arm64 (CUDA 13)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-cuda-13.5-arm64.tar.gz) - [CUDA 13.5 libraries](https://github.com/ggml-org/llama.cpp/releases/download/b12208/cudart-llama-b12208-bin-ubuntu-cuda-13.5-arm64.tar.gz) - [Ubuntu x64 (ROCm 10.0)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-rocm-10.0-x64.tar.gz) - [Ubuntu x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-openvino-2026.4.1-x64.tar.gz) - [Ubuntu x64 (SYCL FP32)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-sycl-fp32-x64.tar.gz) - [Ubuntu x64 (SYCL FP16)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-ubuntu-sycl-fp16-x64.tar.gz) - [Linux arm64 (Snapdragon: CPU, Adreno GPU, Hexagon NPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-linux-arm64-snapdragon.tar.gz) - [setup guide](https://github.com/ggml-org/llama.cpp/blob/master/docs/backend/snapdragon/linux.md) **Android:** - [Android arm64 (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-android-arm64.tar.gz) - [Android arm64 (Snapdragon: CPU, Adreno GPU, Hexagon NPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-android-arm64-snapdragon.tar.gz) - [setup guide](https://github.com/ggml-org/llama.cpp/blob/master/docs/backend/snapdragon/README.md) **Windows:** - [Windows x64 (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-cpu-x64.zip) - [Windows arm64 (CPU)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-cpu-arm64.zip) - [Windows arm64 (OpenCL Adreno)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-opencl-adreno-arm64.zip) - [Windows x64 (CUDA 12)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-cuda-12.4-x64.zip) - [CUDA 12.4 DLLs](https://github.com/ggml-org/llama.cpp/releases/download/b12208/cudart-llama-bin-win-cuda-12.4-x64.zip) - [Windows x64 (CUDA 13)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-cuda-13.5-x64.zip) - [CUDA 13.5 DLLs](https://github.com/ggml-org/llama.cpp/releases/download/b12208/cudart-llama-bin-win-cuda-13.5-x64.zip) - [Windows arm64 (CUDA 13)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-cuda-13.5-arm64.zip) - [CUDA 13.5 DLLs](https://github.com/ggml-org/llama.cpp/releases/download/b12208/cudart-llama-bin-win-cuda-13.5-arm64.zip) - [Windows x64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-vulkan-x64.zip) - [Windows arm64 (Vulkan)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-vulkan-arm64.zip) - [Windows x64 (OpenVINO)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-openvino-2026.4.1-x64.zip) - [Windows x64 (SYCL)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-sycl-x64.zip) - [Windows x64 (ROCm 10.0)](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-bin-win-rocm-10.0-x64.zip) **openEuler:** - [DISABLED](https://github.com/ggml-org/llama.cpp/pull/23705) - openEuler x86 (310p) - openEuler x86 (910b, ACL Graph) - openEuler aarch64 (310p) - openEuler aarch64 (910b, ACL Graph) **UI:** - [UI](https://github.com/ggml-org/llama.cpp/releases/download/b12208/llama-b12208-ui.tar.gz)