hexagon: support for multi-NPU devices (IQ9, IQ10) and fully asynchronous backend

[max-krasnyansky]

Overview

Kind of a big overhaul of the Hexagon backend to enable support for mulit-NPU devices, and to make it fully asynchronous: async graph-compute, async events, async tensor-copy, cross-device fences -- the whole enchilada!

The PR also includes everything required to properly support and enable multi-NPU tensor-split mode (ie integration with the meta backend) on devices like the IQ9 and IQ10 series.
Tested on the IQ9 EVK.
See examples below.

The asynchronous implementation ended up requiring many other updates besides just making graph_compute and cpy_tensor_async not block. To make it fully usable we need to use non-host buffers by default (ie to force async copies and repack), which in turn requires support for more types in SET/GET_ROWS, etc. The CPY op needed to be optimized to not slow things down, and to support fused FENCE op to allow for truly async copies.
Hence this PR includes all those updates.

The multi-NPU support requires new device naming scheme. The original naming HTP0,...HTP4 is still supported and is enabled if the environment variable GGML_HEXAGON_NDEV is defined. The new scheme is HTP0:0, HTP1:0 where the first number is physical NPU id and the second number is the virtual id of the session.

To simplify the migration and support for the devices I included new ./scripts/snapdragon/build.py and ./scripts/snapdragon/run.py and update the documenation to cover the new scripts, devices, etc.

Additional information

Here is an example gemma-4-E4B (q4_0 QAT model) on IQ9 EVK in tensor split mode

./scripts/snapdragon/run.py --target ubuntu:maxk@192.168.1.87 --device HTP0:0,HTP1:0 -- llama-completion \
     -m gemma-4-E4B_q4_0-it.gguf -f sample_prompt_1024.txt --jinja -st --split-mode tensor --ctx-size 8192

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: Yes, Antigravity for a ton of refactoring, testing, debugging and python script updates

[CreaV]

Hi @max-krasnyansky — I've been running this branch on an SA8797 (automotive, v81) and hit a limitation in the device→domain mapping. Patch below; it applies cleanly on qualcomm:hexagon-async and I think it belongs with this PR rather than on its own, since it only matters once multi-NPU exists.

The problem

get_domain_id() computes CDSP_DOMAIN_ID + physical_idx and htpdrv_get_domain() looks that up in the SDK's static supported_domains[]. On SA8797 that only reaches two of four NSPs.

The board's FastRPC runtime resolves these (queried with FASTRPC_GET_EFFECTIVE_DOMAIN_ID):

name	effective domain id
cdsp = nsp1000	1600
cdsp1 = nsp1001	1700
nsp1002	1800
nsp1003	1900

cdsp0, cdsp2, cdsp3 do not resolve at all. So devices 2 and 3 become ids 5 and 6, which aren't in supported_domains[] — that table comes from incs/domain_default.h and stops at CDSP1 = 4 — and allocate() throws unable to get domain struct for CDSP (domain_id 6) before touching a driver.

incs/domain.h picks incs/domain_auto.h under _AUTO, which presumably has the automotive entries, but that header isn't in the public 6.4.0.2 or 6.6.0.0 SDK releases, so the table can't be completed from what ships.

The change

Ask the runtime for the id by domain name instead of computing it. Candidate names are probed in order — existing cdsp/cdspN first, then nsp100N — so nothing is required from the user and parts that already worked are unaffected. GGML_HEXAGON_DOMAIN_PREFIX is there for a third naming scheme, and a runtime without the query falls back to the current arithmetic. A missing supported_domains[] entry stops being fatal, since it only feeds the URI fallback for when FASTRPC_GET_URI fails (that path is now guarded).

Results

llama-batched-bench, one process per NSP, no env vars set:

• before: HTP2/HTP3 could not open at all
• after: each NSP alone matches HTP0's throughput, and all four concurrently give 3.82x aggregate prefill / 3.60x aggregate decode, per-NSP decode down ~10% from DDR contention

I don't have an IQ9/IQ10 to check this doesn't regress there — the probe order should make it a no-op since cdsp/cdspN are tried first, but it would be worth a run on your EVK.

Two other things I ran into on this board, separate from this patch, in case they're useful:

• -sm tensor across the two/four domains aborts in GGML_ASSERT(ggml_backend_buffer_is_meta(meta_buf)) on master, and on this branch with multi buffers are not supported by the meta backend (ggml-backend-meta: multi buffers are unsupported leading to vulkan segfault #22197). The small--c workaround in that issue doesn't help here.
• The LM head can't reach the DSP on tied-embedding models: ggml_hexagon_supported_get_rows only accepts F32, so a quantized token_embd is pinned to a CPU buffer and the output matmul follows it. Giving the GGUF a separate output.weight and raising the ne[1] > 32768 limit moved it over and took steady-state CPU from ~974% to ~10% of twelve cores, with perplexity unchanged (spread 0.012 against ±0.31 per-estimate uncertainty). Happy to open that separately if it's interesting.

Patch

diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
index dadf3bd..898d36c 100644
--- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp
+++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
@@ -13,6 +13,7 @@
 #include <cstddef>
 #include <stdexcept>
 #include <string>
+#include <vector>
 #include <sstream>
 #include <iomanip>
 #include <unordered_set>
@@ -77,15 +78,80 @@ struct ggml_hexagon_device_config {
 
 static ggml_hexagon_device_config opt_device_configs[GGML_HEXAGON_MAX_SESSIONS];
 
-static int get_domain_id(int physical_idx) {
-    return CDSP_DOMAIN_ID + physical_idx;
+// Resolve a device index to a FastRPC domain by asking the runtime.
+//
+// The names and ids are not the same on every SoC. Mobile parts expose cdsp and
+// cdsp1 at CDSP_DOMAIN_ID and CDSP_DOMAIN_ID+1, which is what the arithmetic
+// below assumes. Automotive parts do not: SA8797 has four NSPs named
+// nsp1000..nsp1003 whose effective domain ids are 1600, 1700, 1800 and 1900, so
+// both the naming and the CDSP_DOMAIN_ID+i formula stop working after the
+// second device.
+//
+// The Hexagon SDK acknowledges the split - incs/domain.h picks
+// incs/domain_auto.h over incs/domain_default.h under _AUTO - but domain_auto.h
+// is not part of the public 6.4.0.2 or 6.6.0.0 releases, so the static
+// supported_domains[] table available to us stops at CDSP1 = 4 and
+// htpdrv_get_domain() fails for anything beyond it.
+//
+// Asking FASTRPC_GET_EFFECTIVE_DOMAIN_ID for the id of a given domain *name*
+// avoids both problems and needs no table. Candidate names are probed in order
+// and the first that resolves wins, so no configuration is required; the
+// GGML_HEXAGON_DOMAIN_PREFIX override exists for SoCs that use some third
+// naming scheme.
+static bool query_domain_id(const std::string & name, int * out_id) {
+    remote_rpc_effective_domain_id_t d;
+    memset(&d, 0, sizeof(d));
+    d.domain_name     = const_cast<char *>(name.c_str());
+    d.domain_name_len = name.size();
+    d.session_id      = 0;
+
+    if (remote_session_control(FASTRPC_GET_EFFECTIVE_DOMAIN_ID, (void *) &d, sizeof(d)) != AEE_SUCCESS) {
+        return false;
+    }
+    if (out_id) {
+        *out_id = (int) d.effective_domain_id;
+    }
+    return true;
 }
 
-static std::string get_domain_name(int physical_idx) {
+// Names to try for device i, most specific first.
+static std::vector<std::string> domain_name_candidates(int physical_idx) {
+    const char * prefix = getenv("GGML_HEXAGON_DOMAIN_PREFIX");
+    if (prefix && prefix[0] != '\0') {
+        return { std::string(prefix) + std::to_string(physical_idx) };
+    }
+
+    std::vector<std::string> names;
     if (physical_idx == 0) {
-        return CDSP_DOMAIN_NAME;
+        names.push_back(CDSP_DOMAIN_NAME);
+    } else {
+        names.push_back(std::string("cdsp") + std::to_string(physical_idx));
+    }
+    names.push_back(std::string("nsp100") + std::to_string(physical_idx));
+    return names;
+}
+
+static std::string get_domain_name(int physical_idx) {
+    const std::vector<std::string> names = domain_name_candidates(physical_idx);
+    for (const std::string & name : names) {
+        if (query_domain_id(name, nullptr)) {
+            return name;
+        }
+    }
+    return names.front();
+}
+
+static int get_domain_id(int physical_idx) {
+    int id = 0;
+    for (const std::string & name : domain_name_candidates(physical_idx)) {
+        if (query_domain_id(name, &id)) {
+            return id;
+        }
     }
-    return std::string("cdsp") + std::to_string(physical_idx);
+
+    // Runtimes that do not implement the query keep the legacy numbering, which
+    // is correct on the SoCs that only have cdsp and cdsp1.
+    return CDSP_DOMAIN_ID + physical_idx;
 }
 
 static int    opt_arch    = 0; // autodetect
@@ -2122,10 +2188,12 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {
 
     GGML_LOG_DEBUG("ggml-hex: %s allocating new session\n", this->name.c_str());
 
+    // Only used to build a fallback URI if FASTRPC_GET_URI fails below, so a
+    // domain the SDK's static table does not know about is not fatal.
     domain * my_domain = htpdrv_get_domain(this->domain_id);
     if (my_domain == NULL) {
-        GGML_LOG_ERROR("ggml-hex: unable to get domain struct for CDSP (domain_id %d)\n", this->domain_id);
-        throw std::runtime_error("ggml-hex: failed to get CDSP domain (see log for details)");
+        GGML_LOG_DEBUG("ggml-hex: no static domain entry for domain_id %d; relying on FASTRPC_GET_URI\n",
+                       this->domain_id);
     }
 
     std::string dom_name = get_domain_name(phys_idx);
@@ -2168,6 +2236,12 @@ void ggml_hexagon_session::allocate(int dev_id) noexcept(false) {
 
         int err = remote_session_control(FASTRPC_GET_URI, (void *) &u, sizeof(u));
         if (err != AEE_SUCCESS) {
+            if (my_domain == NULL) {
+                GGML_LOG_ERROR("ggml-hex: FASTRPC_GET_URI failed for %s (error 0x%x) and there is no static entry for domain_id %d to fall back to\n",
+                               dom_name.c_str(), err, this->domain_id);
+                throw std::runtime_error("ggml-hex: cannot build a session URI (see log for details)");
+            }
+
             // fallback to single session uris
             int htp_URI_domain_len = strlen(htp_uri) + MAX_DOMAIN_NAMELEN;

---

AI usage disclosure: this patch and comment were written by Claude (Anthropic) working from my analysis and hardware; I reviewed both and ran all the measurements quoted above on my board.

[max-krasnyansky]

Hi @max-krasnyansky — I've been running this branch on an SA8797 (automotive, v81) and hit a limitation in the device→domain mapping. Patch below; it applies cleanly on qualcomm:hexagon-async and I think it belongs with this PR rather than on its own, since it only matters once multi-NPU exists.

@CreaV sorry for the delay. Yes, we flagged this issue internally as well. The plan is to add runtime query in the next iteration. For now I updated the hardcoded domain IDs that should for your SA8797. The reason I didn't add the query thing is because that remote_... call is not available on Android (latest gen might have it), WoS and Ubuntu targets (IQ8/Ventuno, IQ9, etc). Please give the latest version a try and we'll add the optional query in the next round.

[CreaV]

Thanks for the update and for explaining the API availability constraint. I tested the latest head (24e22fa) on our SA8797 Android 16 guest (v81), using a clean rebuild of both the host libraries and the DSP skel.

The new static mapping gets HTP0:0 and HTP1:0 working, but HTP2:0 and HTP3:0 still fail before driver access:

ggml-hex: unable to get domain struct for CDSP (domain_id 18)
ggml-hex: unable to get domain struct for CDSP (domain_id 19)

With the public Hexagon SDK 6.6.0.0, supported_domains[] still has no entries for 18/19, so changing get_domain_id() alone is not enough on this board. Reapplying the runtime-query patch on top of 24e22fa makes all four physical NSPs enumerate successfully, and both 2-way and 4-way tensor-split benchmarks complete.

One configuration detail I also noticed: GGML_HEXAGON_DEVICES=4 now creates four virtual sessions on physical HTP0, so for the four physical SA8797 NSPs I used GGML_HEXAGON_DEVICES=HTP0:0,HTP1:0,HTP2:0,HTP3:0.

Thanks again for the quick turnaround and all the work on this backend. I'll be happy to test the optional runtime-query version when it is ready.

[max-krasnyansky]

@lhez for review/ack (I know it's huge but it got a bit stuck behind meta changes and got better and better :))

[jeff-zhouwg]

line 1278: GGML_ASSERT(offset == 0);
line 1280: repack_tiled_q4_0(data, tensor, offset, size);

under what scenario would offset != 0? is this added parameter offset intended for multi‑NPU cases?

[max-krasnyansky]

get_tensor_2d (used with tensor-split) uses offsets. And the set_tensor... path with host_buffer enabled does partial sets, our new shadow buffer hides that though.

[max-krasnyansky]

Thanks for the update and for explaining the API availability constraint. I tested the latest head (24e22fa) on our SA8797 Android 16 guest (v81), using a clean rebuild of both the host libraries and the DSP skel.

The new static mapping gets HTP0:0 and HTP1:0 working, but HTP2:0 and HTP3:0 still fail before driver access:

ggml-hex: unable to get domain struct for CDSP (domain_id 18)
ggml-hex: unable to get domain struct for CDSP (domain_id 19)

With the public Hexagon SDK 6.6.0.0, supported_domains[] still has no entries for 18/19, so changing get_domain_id() alone is not enough on this board. Reapplying the runtime-query patch on top of 24e22fa makes all four physical NSPs enumerate successfully, and both 2-way and 4-way tensor-split benchmarks complete.

Ah. Bummer. Ok. Let's fix it in the next round then. I forgot that we need more infra in there.

One configuration detail I also noticed: GGML_HEXAGON_DEVICES=4 now creates four virtual sessions on physical HTP0, so for the four physical SA8797 NSPs I used GGML_HEXAGON_DEVICES=HTP0:0,HTP1:0,HTP2:0,HTP3:0.

Yep. That's intentional. As part of testing the ggml-meta backed fixes I was playing with Metal and CUDA and they use GGML_..._DEVICE=N to create virtual devices on a single GPU, and name them as CUDA0, CUDA1, ... so I figured I'll make Hexagon behave the same to be consistent. The majority of the Snapdragon devices right now are single NPU so that's the common case. And for multi-NPU we can let uses specify the exact breakdown of phy/virt devices.

Thanks again for the quick turnaround and all the work on this backend. I'll be happy to test the optional runtime-query version when it is ready.

Thanks

[max-krasnyansky]

@ggml-org/maintainers can I get the second approval.
Sorry for noise. I added merge-ready label but it'd be good to merge soon since we have a bunch of follow ups based on this.

[taronaeo]

Providing another approval :)

[CISC]

@max-krasnyansky Tests don't work:
https://github.com/ggml-org/llama.cpp/actions/runs/33031193441/job/98385718584
https://github.com/ggml-org/llama.cpp/actions/runs/33121088119/job/98690786119

[max-krasnyansky]

@max-krasnyansky Tests don't work: https://github.com/ggml-org/llama.cpp/actions/runs/33031193441/job/98385718584 https://github.com/ggml-org/llama.cpp/actions/runs/33121088119/job/98690786119

Oops. My bad. I noticed the failure but at first glance it looks like one of those intermittent things (ie runner failed to call QDC API). But now I see that it's my script changes that broke it. Will fix asap. Thanks for the ping.

[GaTmaNnes]

Thanks for putting this together, @max-krasnyansky. Based on our root-cause analysis from the K3-K5 campaign, here is how our findings align with what this PR actually fixes regarding our current bottleneck:1. Where PR #26501 helps (Actionable & Validated)lm_head Offloading (Real Impact): Moving output.weight to HTP (via Q8_0 GET_ROWS support) is the single biggest win here. In our profiling, this single-handedly dropped CPU usage from ~974% down to ~10% on SA8797 by eliminating heavy CPU-side operations in the decode graph.Fully Async Backend (Moderate Impact): Making graph_compute, events, and tensor copies non-blocking is a great architectural improvement. While it doesn't reduce the orchestration overhead itself, it successfully allows overlapping DSP execution with AP-side work (~10-20% expected gain).2. The Remaining Bottleneck (Why we still need more)Despite these fixes, our profiling shows we are still bounded by the GGML scheduler generating ~730 micro-graphs per token. This leads to a massive RPC overhead ratio:$T_{\text{DSP}} \approx 66\text{ms}$ (stable, HTP hardware is fast)$T_{\text{AP/RPC}} \approx 65\text{ms}$ (orchestration overhead per split)Because interleaved CPU ops (like rms_norm and samplers) force the scheduler to split the graph at every CPU $\leftrightarrow$ HTP backend switch, the 730 RPC calls per token remain untouched by async execution alone. To fully bridge the gap to QAIRT-level performance (~17 tok/s compared to our current ~7.3 tok/s limit), we are looking at either:Moving remaining CPU ops entirely onto the HTP backend, orModifying the scheduler/implementing a topological sort pass to merge adjacent backend switches (similar to what we explored in our pass-4.5 attempt).Bottom line: This PR provides essential building blocks—especially the lm_head offloading, which we consider production-critical. However, it's part 1 of solving the throughput wall; graph-splitting overhead will need a follow-up strategy.