LocalAI/backend/cpp/llama-cpp-localai-paged/docs/ACCELERATOR_PORTING_SCOPE.md

Accelerator-porting scope: bringing the paged backend's portable benefits to Metal / SYCL / Vulkan (+ a ROCm note)

Source-only analysis (no GPU, no build) of which llama-cpp-localai-paged benefits are portable off the CUDA family, and what each port costs per accelerator. This is the umbrella doc; it BUILDS ON, and does not repeat, UPSTREAM_LAYER2_SCOPE.md (the GDN/SSM fusion kernel scope) - that doc remains the authoritative reference for benefit #1 below.

The backend ships CUDA-only today (README sections 4c, 8): off-CUDA the fusions gate off (patch 0030) and NVFP4 falls back to dequant, so it is neutral-to-slightly-negative there and non-CUDA users run the stock llama-cpp. "Porting the benefits" is the upstream-contribution track that would make these wins real on the other accelerators. Methodology for the work itself is in .agents/vllm-parity-methodology.md.

We have no Metal / SYCL / Vulkan / ROCm hardware here, so every port is gated by test-backend-ops (backendX-vs-CPU) on the target hardware - the same gate discipline the existing layer-2 doc sets out.

---

0. The four benefits and their portability class

#	Benefit (patches)	Portable off CUDA?	Where scoped
1	GDN/SSM decode fusions (0018-0022, 0028) - in-place state write-back, fused recurrent-state gather, conv-state in-place fusion, o_proj MMQ reshape, occupancy retune	YES - per-backend KERNEL work	UPSTREAM_LAYER2_SCOPE.md (consolidated in section 1 here)
2	Paged KV in-kernel block-table flash-attn read (0009-0011)	YES - per-backend KERNEL work	Section 2 here (the new analysis)
3	Decode-first prefill scheduler (0013/0016)	YES - FREE, host-side, zero kernel work	Section 3 here
4	NVFP4 FP4-MMA + its decode levers (0017/0023/0025)	NO (Blackwell FP4-MMA) - out of scope; two analogues flagged	Section 4 here

The two kernel-bearing tracks (#1 and #2) share an identical port SHAPE - they touch the same decode kernel(s), the same supports_op, the same dispatch guard, and sequence the same way (ops-first PR, then one PR per backend). They should be bundled into one per-backend PR, not pursued as two separate efforts; section 5 sequences them together. Tracks #3 (free) and #4 (out of scope) are independent.

---

1. Benefit #1 - GDN/SSM decode fusions (consolidated; full scope is the layer-2 doc)

Do not re-derive this here. UPSTREAM_LAYER2_SCOPE.md already establishes, and this doc adopts wholesale:

• The base GGML_OP_GATED_DELTA_NET + GGML_OP_SSM_CONV + GGML_OP_SSM_SCAN kernels already exist on Metal, Vulkan AND SYCL, so the Qwen3.6 hybrids RUN on all three today via the upstream non-fused path. Layer-2 is the decode SPEEDUP, not "make it run." (NB: the README section 4c no longer carries the stale "no Vulkan kernel" line that the layer-2 doc section 0 was written to correct - that correction has since been folded into the README, so treat layer-2 section 0 as historical context, not a live correction.)
• The four fusion ops (A in-place state 0018, B fused state gather 0019, C conv-update in-place 0021, D conv-tap gather 0028) reuse the existing op enums with extra src[] discriminators; only OP C is a genuinely new kernel, the rest redirect the read source / write target of the EXISTING kernel. The builders, CPU reference kernels, model graph and test-backend-ops cases are SHARED and already done.
• Per-backend net-new work, effort and gotchas: SYCL easiest (near-verbatim CUDA mirror, ~250-350 LOC, no shader-gen), Metal medium (~350-500 LOC, fixed 32 simdgroup = simplest bit-exactness), Vulkan hardest (~450-650 LOC + shaders-gen + descriptor growth + per-vendor subgroup validation).
• Bit-exactness is per-backend BY CONSTRUCTION (the fusions redirect addresses, not the f32 reduce order); gated by test-backend-ops (backendX-vs-CPU).
• Upstream path: ops-first PR (incl. the capability-driven replacement for patch 0030's backend-name allow-list), then one PR per backend.

The value/effort ranking from that doc (Metal 1st, SYCL 2nd, Vulkan 3rd) is adopted unchanged here and, as section 5 shows, coincides with benefit #2's ranking

• which is why the two bundle cleanly per backend.

---

2. Benefit #2 - paged KV in-kernel block-table flash-attn read (NEW scope)

2.0 What it is, and why it is the lever that makes paged KV non-negative off-CUDA

On CUDA, patches 0009-0011 replaced the per-step host-side K/V gather (patch 0003) with an in-kernel paged read. ggml_flash_attn_ext gained an optional src[5] = an I32 block table [n_view, n_stream] in token-POSITION order; the fattn vec/tile kernel maps logical KV index j to physical cell block_table[seq*ne11 + j] and reads K0 + cell*nb11 / V0 + cell*nb21 in place, so the get_rows of K and V (the bulk of the gather) is gone. A null block table is the stock contiguous read, byte-identical. Position ordering keeps the online-softmax reduction order identical to stock, so it is bit-exact (CPU/batch1) by construction.

The crucial point for portability: the entire host side is already backend-agnostic. The block-table fill (llama_kv_cache::get_block_table), the K/V views, the mask compaction, the input_block_table graph input, and the ggml.c / ggml.h builder (ggml_flash_attn_ext_set_block_table) all live in src/ and ggml/... shared code. The ONLY per-backend work is, in each backend's flash-attn kernel: (a) thread one extra source through to the kernel, and (b) do the indexed read at the K/V load sites. The CPU reference already does it (patch 0009, ops.cpp).

Off-CUDA today the paged path falls back to the host-side gather (patch 0003), which the README section 4c measured as neutral-to-slightly-negative on the M4 (~0-3% slower decode / ~2-8% slower prefill vs stock's contiguous read - pure overhead, because the in-kernel read that recovers the gather cost is CUDA-only). Porting the block-table read is exactly what flips paged KV from "neutral-to-negative" to "neutral-to-positive" off CUDA - it removes the gather overhead so paged KV's memory-management and prefix-sharing wins come for free instead of at a decode tax. (The big decode multipliers on the hybrids are still the benefit-#1 GDN fusions; this benefit is what makes the paged allocator pay its own way off CUDA.)

2.1 The cross-cutting finding (applies to all three backends)

The indexed per-cell read only fits the vec / scalar decode kernel. Every backend's FAST attention path - CUDA mma, Metal simdgroup_load MM, Vulkan coopmat2, SYCL tile - loads K/V as contiguous tiles (8-cell simdgroup_load, coopMatLoadTensorNV over a linear stride, shared-memory tile loads) that cannot express an arbitrary per-cell gather without a staging pre-pass. This is exactly why the CUDA port (0009-0010) wired ONLY the vec kernel and added a dispatch guard (if (dst->src[5]) force vec).

So each port mirrors that: route any FA op carrying a block table onto the vec / scalar kernel; leave the fast MM path contiguous-only, and keep the null-table contiguous read on the fast path untouched. The decode shape (1 query token/stream) naturally lands on or near the vec/scalar kernel on all three, so this is a small routing change, not a rewrite of the fast path.

2.2 SYCL - EASIEST (near line-for-line CUDA mirror)

• Exists today: ggml-sycl/fattn-vec.hpp is a DPCT-style near-verbatim mirror of CUDA fattn-vec.cuh; the kernel signature ends in the same nb11..nb33 cluster the CUDA patch appends const int* block_table to (fattn-vec.hpp:65-76). Args are passed by SYCL lambda value-capture - no descriptor/binding/push- constant bookkeeping at all (strictly easier than CUDA). supports_op (fattn.cpp -> ggml_sycl_get_best_fattn_kernel) needs no change to ACCEPT src[5].
• Port shape (value: medium / effort: LOW): append const int* block_table to the kernel + fattn_kernel_t typedef + lauch_kernel/launch_fattn (sourcing dst->src[5]->data); 3 read-site substitutions (K at line 318, V at 389 and 410): K0 + block_table[seq*ne11 + k_VKQ_0 + i_KQ]*nb11.
• Two SYCL-specific gotchas:
  1. Pointer pre-advance. The vec kernel advances K/V by k_VKQ_0 OUTSIDE the inner read (fattn-vec.hpp:293-300), so i_KQ/k are tile-local. The port must keep an UN-advanced base K0/V0, drop the per-iteration K +=/V += on the paged path, and reconstruct the absolute cell. Get this wrong and you read the wrong cells with NO compile error.
  2. Dispatch guard is bigger than CUDA's. f16-GQA decode routes to the TILE kernel, not vec (fattn.cpp:198-208 fall-through). Add if (dst->src[5]) return BEST_FATTN_KERNEL_VEC; near the top of ggml_sycl_get_best_fattn_kernel. The shared fattn_kernel_t typedef means the tile kernel must gain a matching ignored block_table param (or split the typedef) - a trivial chore.
• Bit-exact: sub-group width (16) is fixed and the indexed read does not touch lane assignment, loop bounds, or the XOR-reduction stride - reduction order is invariant, so the paged vec path is byte-identical to SYCL's own contiguous vec path. Gate: test-backend-ops FLASH_ATTN_EXT (with a block-table case) on Intel GPU.

2.3 Metal - EASY-MEDIUM (decode already routes to the vec kernel)

• Exists today: decode (1 query token/stream, GQA) dispatches to kernel_flash_attn_ext_vec (ggml-metal-ops.cpp ..._use_vec: ne01 < 20). Metal IS a true vec-equivalent (not a single unified FA kernel), and the vec kernel's quantized K/V branches ALREADY compute a per-cell base address (k + ((ic + NE*cc + ty)*nb11), ggml-metal.metal:6934 / V at :7045) - so a per-cell indexed read is unambiguously admissible. supports_op (ggml-metal-device.m FLASH_ATTN_EXT) inspects no src count, so src[5] is accepted as-is.
• Port shape (value: HIGH / effort: EASY-MEDIUM): append a device const char * block_table param after dst (buffer index 8 for vec)
  • a kargs field + a has_block_table function-constant; reuse the existing "bind dummy when null" idiom for a missing table; substitute the cell index with block_table[seq*ne11 + cell] at the K reads (lines 6919/6934) and V reads (7032/7045) - a localized rewrite of ~2 loops (the fast path must adopt the per-cell base form the quantized branch already uses).
• Gotcha: the non-vec MM kernel is a HARD blocker - simdgroup_load(..., NS10, ...) reads 8 physically-CONTIGUOUS KV cells as one matrix tile (lines 6160 / 6339-6363); an arbitrary gather can't be a single strided matrix load. Mitigate exactly as CUDA did: force any block-table op onto the vec kernel in ..._use_vec (ggml-metal-ops.cpp:2517); leave the MM path contiguous-only. Also watch a NAME COLLISION: kernel_flash_attn_ext_blk is an existing mask-skip optimization, NOT a paged block table.
• Bit-exact: fixed 32-wide simdgroup + address-only redirect = byte-identical to Metal's own vec contiguous path. Gate: test-backend-ops on Apple Silicon.

2.4 Vulkan - MEDIUM (the fast NVIDIA decode path cannot do it)

• Exists today: three FA shaders - flash_attn.comp (scalar/vec), flash_attn_cm1.comp (coopmat1, stages K/V through shared memory), flash_attn_cm2.comp (coopmat2, the fast NVIDIA path). FA uses 7 descriptor bindings (0-6); supports_op (ggml-vulkan.cpp FLASH_ATTN_EXT) checks specific srcs only, no count check; but src[5] is not even threaded today - ggml_vk_flash_attn stops at src[4] (ggml-vulkan.cpp:14537), so wiring it through is part of the work.
• Port shape (value: HIGHEST breadth / effort: MEDIUM): add binding 7 in the shader(s), bump 7->8 in the three ggml_vk_create_pipeline calls (:3997, :4033, :4070) and the two dispatch subbuffer lists (passing a dummy when null), and wrap the indexed read in one phys_kv() helper applied at the ~4 K + 2 V load sites (flash_attn.comp; the logical index is the same (j*Bc + ...) expression at every site).
• Two gotchas, one structural:
  1. Push constants are FULL. vk_flash_attn_push_constants is exactly 128 bytes with a static_assert(... <= 128) (the Vulkan guaranteed minimum) - no room for a new field. Signal "block-table enabled" via the existing Flags spec constant (flash_attn_base.glsl, constant_id=10, already bit-packed) - add a BLOCK_TABLE_ENABLE bit. The per-seq stride is already p.KV; the seq index is derivable in init_indices().
  2. coopmat2 (the fast NVIDIA GQA-decode path) is INCOMPATIBLE. Its K/V load is a hardware coopMatLoadTensorNV over a LINEAR stride (flash_attn_cm2.comp:307-313/377-383); the decode callback only dequantizes, it cannot remap the physical address. The indexed read drops cleanly into scalar (which non-GQA decode already uses) and cm1 (which stages through shmem - remap the staging loop), but not cm2. With a block table present, NVIDIA GQA decode falls back to scalar/cm1 (slower than cm2, still correct); the null-table path keeps using cm2 unchanged. AMD/Intel (no cm2) are fully covered by scalar/cm1.
• Net positive? Yes. Non-GQA decode already runs scalar (paged read ~free); AMD/Intel covered; only NVIDIA GQA decode trades cm2 for scalar/cm1 when a table is supplied, and paged KV's payoff is allocator/memory + prefix-sharing, not raw FA throughput, so the trade is contained and the fast contiguous path is untouched.
• Bit-exact: the read is a per-thread scalar load, subgroup-size agnostic (already abstracted via the SubGroupSize spec constant); position ordering keeps the reduction order identical, so byte-identical to the backend's own scalar/cm1 contiguous path. Build burden is low - these are EXISTING shader variants recompiling (no new string_to_spv shape), so no shaders-gen matrix growth. Gate: test-backend-ops per vendor (AMD + Intel + NVIDIA).

2.5 Benefit-#2 ranking and the shared dispatch/supports_op pattern

backend	value	author effort	structural risk	rank
SYCL	medium (Intel GPU)	LOW (line-for-line; no bindings)	low (pointer pre-advance; force-vec guard)	easiest
Metal	HIGH (largest non-CUDA base)	EASY-MEDIUM (decode = vec already)	medium (MM blocker -> force vec)	mid
Vulkan	HIGHEST breadth (AMD+Intel+NVIDIA)	MEDIUM (7->8 bindings; Flags bit)	medium (cm2 can't; full push-const)	hardest

Common to all three (mirrors CUDA 0009-0010): (1) supports_op needs no change to ACCEPT src[5]; (2) a dispatch guard forces any block-table op onto the vec/scalar kernel; (3) the fast MM/coopmat2 path stays contiguous-only and the null-table read on it is byte-identical to stock.

---

3. Benefit #3 - decode-first prefill scheduler (FREE portable win, confirmed)

Patches 0013 (static LLAMA_PREFILL_BUDGET) and 0016 (dynamic decode-first max(n_ubatch, T-D)) are pure host-side scheduler policy inside update_slots() with zero libllama / zero ggml-backend changes (README sections 2, 3). They change only the count of prefill tokens admitted per step; they touch no kernel, no supports_op, no device code. They are therefore already backend-portable with no per-accelerator work - they run identically on Metal, SYCL, Vulkan, ROCm, CPU. Byte-identical when off (default-off / short prefill == upstream -b chunking).

This is the cheapest portable benefit: it needs no port at all, only the decision to leave it enabled in the (currently CUDA-only) build, or to upstream the policy. The only reason it is not "live everywhere" today is that the backend ships CUDA-only; the code itself is accelerator-neutral. If the scheduler levers are upstreamed independently of the kernels, they help any llama.cpp build on any accelerator at once - the lowest-effort, broadest-reach contribution of the whole series.

---

4. Benefit #4 - NVFP4 FP4-MMA (NOT portable) + two backend-agnostic analogues

The NVFP4 decode track is Blackwell-specific and out of scope for accelerator porting: Metal, SYCL, Vulkan and ROCm/AMD lack native FP4-MMA (Metal supports_op already excludes NVFP4 from MUL_MAT/MUL_MAT_ID/GET_ROWS; on non-Blackwell the FP4 path dequants). Patch 0017 (dense FP4-GEMM occupancy tune) ships only as the parity gate + default-off instrumentation even on CUDA, so there is nothing to port.

Two of the NVFP4 decode levers, however, have backend-agnostic analogues worth a note (do not over-claim - these are observations, not scoped ports):

• 0023 (NVFP4 activation-quantize de-dup) - the IDEA generalizes, the patch does not. The MoE broadcast up/gate projections re-quantize the same token activation once per expert; 0023 quantizes the unique activations once and byte-copies them into the expert-gathered layout. Any backend whose MoE path requantizes a shared activation per-expert (e.g. a Q8 activation-quant before an integer-dot MoE GEMM) could dedup the same way. It is NOT NVFP4-specific in PRINCIPLE - but it IS the one quant-specific patch in the series (README section 6), so a port is a per-backend MoE-quant investigation, not a lift-and-shift. Low priority.
• 0025 (MoE decode re-graph / LLAMA_MOE_FORCE_GRAPHS) - keeping the graph/ capture path on across the grouped-MMQ MoE decode step is a CUDA-graphs concept. Metal/Vulkan/SYCL have their own command-buffer/graph reuse machinery; the generalizable finding is "the grouped MoE decode step has no host sync, so it is safe to keep in a captured/replayed command buffer." Whether each backend's graph layer already covers this is a per-backend question. The methodology note (README dev notes: graph/stream coverage was a FLAT lever beyond 0025 on CUDA) is the more durable takeaway - do not expect a large graph-coverage win on any backend.

Neither analogue is on the critical path; both are recorded so the next person does not mistake them for free ports.

---

5. Combined sequencing and top recommendations

Benefits #1 (GDN fusions) and #2 (block-table FA read) share the port shape (vec/scalar decode kernel + supports_op/dispatch guard + ops-first-then-per-backend PR) and rank in the SAME order per backend. So sequence them TOGETHER, per backend, behind one shared ops-first PR:

1. PR #1 - OPS (largely done, upstreamable as-is): the ggml.h/ggml.c builders, the CPU reference kernels, the CUDA kernels, the test-backend-ops cases (GDN fusions AND a FLASH_ATTN_EXT block-table case), and the capability-driven gate replacing patch 0030's backend-name allow-list (make supports_op + the dispatch guard authoritative, so routing falls out of the normal scheduler fallback and no backend name is hard-coded). Independently mergeable.
2. PR #2 - Metal: GDN fusion kernels (layer-2 doc) + block-table read into kernel_flash_attn_ext_vec + the force-vec routing guard. Gate on Apple Silicon.
3. PR #3 - SYCL: the near-verbatim CUDA mirror of both tracks + the force-vec guard. Gate on Intel GPU.
4. PR #4 - Vulkan: GDN fusion shaders + the scalar/cm1 block-table read (cm2 stays contiguous, falls back when a table is present) + the Flags spec-constant bit + the 7->8 binding bump. Gate per vendor.

Do NOT bundle the backends into one PR (each needs its own hardware for test-backend-ops; reviewers are backend-specialized; a regression in one must not block the others).

Top recommendations

1. Metal first, both benefits together. Largest non-CUDA LocalAI base; the decode shape already routes to the Metal vec kernel (block-table read is EASY-MEDIUM there) and the base GDN/conv kernels already exist (fusions are MEDIUM); fixed 32-wide simdgroup makes bit-exactness the simplest of the three. Highest value at moderate effort.
2. SYCL second as the cheap mechanical follow-on. Both tracks are near line-for-line CUDA mirrors with no binding/shader-gen bookkeeping, so it is low-cost insurance even though the Intel-GPU audience is smaller. Budget the effort on the two SYCL gotchas (pointer pre-advance; the force-vec guard since f16-GQA decode routes to tile), not on plumbing.
3. Vulkan last as the high-breadth capstone. Reaches AMD + Intel + NVIDIA, but carries the most host glue and the coopmat2 limitation (NVIDIA GQA decode trades the fast path for scalar/cm1 only when a table is present). Do it once the pattern is proven on Metal + SYCL.

A cheaper variant (from the layer-2 doc, reaffirmed): ship Metal + SYCL together right after the ops PR and treat Vulkan as a separate later effort.

---

6. ROCm note

ROCm is in the CUDA family, not a separate port: patch 0030's allow-list already admits "CUDA"/"ROCm"/"MUSA", and the CUDA kernels compile for HIP, so benefits #1 and #2 are largely already-built or near-free on ROCm rather than a from-scratch accelerator port. Two caveats:

• FP4-MMA (benefit #4) stays NVIDIA-Blackwell-only - AMD has no native FP4-MMA, so the NVFP4 path dequants on ROCm exactly as elsewhere.
• The block-table read's force-vec routing matters on AMD too. The AMD fast FA path is the wmma/mma kernel (fattn-wmma-f16), which - like CUDA mma, Metal MM and Vulkan cm2 - ignores the block table; the CUDA dispatch guard already forces a block-table op onto the vec kernel, so ROCm inherits correct routing, but the perf trade (vec vs wmma for AMD GQA decode with a table present) should be measured on AMD hardware before claiming a win. The GDN fusions, being plain CUDA-C, port to HIP with the rest of the CUDA path.

Net: ROCm is a "validate, don't re-port" follow-up - confirm the HIP build picks up the fusions + the force-vec block-table routing and gate it with test-backend-ops on an AMD GPU. It is genuinely separate from, and lighter than, the Metal / SYCL / Vulkan ports.

---

7. Summary

• Benefit #3 (decode-first scheduler) is free and already portable - host-side policy, zero kernel work; it only needs to be left enabled / upstreamed.
• Benefits #1 (GDN fusions) and #2 (block-table FA read) are the real ports - both are vec/scalar-decode-kernel + supports_op/dispatch-guard changes, both rank Metal-then-SYCL-then-Vulkan, and they bundle into one per-backend PR behind a shared ops-first PR.
• Benefit #2 is the lever that makes paged KV non-negative off CUDA - it removes the host-gather overhead the README measured as neutral-to-slightly-negative on the Mac. Feasibility: SYCL EASY, Metal EASY-MEDIUM, Vulkan MEDIUM. The universal constraint is that only the vec/scalar kernel admits the indexed read; the fast MM/coopmat2 path is contiguous-only, so route block-table ops onto vec (as CUDA already does) and leave the fast path's null-table read byte-identical.
• Benefit #4 (NVFP4 FP4-MMA) is out of scope (Blackwell only); 0023's de-dup and 0025's graph-coverage have backend-agnostic ideas but no lift-and-shift port.
• ROCm rides the CUDA path (validate, don't re-port); FP4-MMA stays Blackwell-only.
• Everything is bit-exact per-backend BY CONSTRUCTION (position-ordered table + address-only redirect = identical reduction order), gated by test-backend-ops (backendX-vs-CPU) on the target hardware, which we do not have here.