Skip to content

b12208

Pre-release
Pre-release

Choose a tag to compare

View all tags
@github-actions github-actions released this 08 Nov 15:47
b12208
57d0482
This commit was created on GitHub.com and signed with GitHub’s verified signature.
GPG key ID: B5690EEEBB952194
Verified
Learn about vigilant mode.

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

Website:

Attestations:

macOS/iOS:

Linux:

Android:

Windows:

openEuler:

  • DISABLED
  • openEuler x86 (310p)
  • openEuler x86 (910b, ACL Graph)
  • openEuler aarch64 (310p)
  • openEuler aarch64 (910b, ACL Graph)

UI:

Assets 40
Loading