llama.cpp: CDNA sparse attention support, revised kernel selection

1. Removing the AMD execution barrier in shared CUDA code

The change was merged into llama.cpp on October 29. It removes conditions blocking sparse attention in HIP builds from the backend shared by CUDA and HIP, AMD's GPU execution and development environment. It extends the existing CUDA path to AMD CDNA while retaining its selection conditions.

Previously, HIP excluded sparse attention code from compilation, and the function gathering required positions also aborted execution. HIP now performs both this preparation and sparse computation. On CDNA, the rules prioritize sparse matrix kernels over the existing workload-size rules when sparse conditions are met.

2. Reducing computation at the cost of gathering positions first

Sparse attention gathers accessible key positions into a list from a mask that specifies which positions to include, then processes the selected keys and values. Rather than computing masked-out positions, it handles only needed positions. This trades reduced attention computation for the upfront cost of building the position list.

This implementation stores position selections to match the size of the GPU execution group. It uses a 64-bit mask to accommodate CDNA's 64 parallel execution lanes and adjusts the counting of selected positions accordingly.

3. Choosing sparse kernels to fit the workload

According to code comments, CDNA sparse support is compiled for the variant with DKQ, the query/key dimension, set to 576 and ncols2, the kernel's column grouping setting, set to 16. It first checks whether this variant is available, then checks the actual workload conditions. The author left support for other architectures, including RDNA, for later work.

At runtime, workloads qualify as sparse-path candidates when a mask exists and the bias added to results and the value cap setting are both 0. On CDNA, the key length must be at least 4096 and at least twice the maximum number of keys in the sparse list. The software thus checks workload size and shape to decide the execution path.

4. Synthetic test throughput and variability

pp512 and tg64 identify test categories, while d16384 and d32768 indicate test depth. The two columns below follow the measurement order in the source. The ± values beside throughput are also unchanged.

Table: Two measurement results from the qwen4exp synthetic test

Conditions: Synthetic test by the PR author. qwen4exp A3B Q4_K - Medium, 103.68 GiB, 176.94 B, ROCm, 50 repetitions. Loaded across 2x mi100s, 1x r9700, 1x rx7900xtx. Includes a separate patch to enable the sparse path.

Unit: t/s (tokens processed per second)

In the first measurement, increasing depth from d16384 to d32768 lowers pp512 throughput from 760.52 to 633.18 t/s and tg64 from 35.81 to 30.78 t/s. The pp512 ± value also rises from 55.29 to 99.63, indicating lower throughput and greater variability in the deeper test for this measurement.

The author described the performance improvement in this test as clear. They also argued that the results understate the gains because roughly half the model was loaded on an architecture unaffected by this change. The performance table does not identify before and after, so the two columns alone cannot establish an improvement rate.

5. Accounting for older ROCm builds

A compatibility definition uses the existing __ballot instead of __ballot_sync to collect per-lane selections on versions below ROCm 7.0. The author noted that before ROCm 7.0 the __ballot_sync built-in was hidden behind HIP_ENABLE_WARP_SYNC_BUILTINS. The code added in the source is shown below.

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

The author said __ballot can simply be used in these cases for this kernel. The compatibility macro itself, however, is defined in the shared HIP header ggml/src/ggml-cuda/vendors/hip.h for HIP_VERSION < 70000000, with no kernel-specific restriction. He also flagged an outstanding problem: __ballot_sync is used in another file, where the mask is actively used on wave64 devices, which makes that use incorrect for wave64 on ROCm 6.3-6.4. He said he will follow up on it, and no such follow-up is included in this PR, so the compatibility described here does not resolve that remaining correctness limitation.

The core change enables a path that computes only required positions on CDNA and connects it to kernel selection based on workload conditions.