Reverse engineered Qualcomm NPU compiler to uncover undocumented behaviour

[Darshan]

This was my first reverse‑engineering attempt on a proprietary software piece. I work in Edge AI, and I was about to throw in the towel because there's barely any documentation on NPUs - everyone from Qualcomm to MediaTek to Apple keeps it hidden.

I discovered some genuinely new stuff:

1. The compiler can silently downgrade your model's weight precision without telling you.

2. Memory placement uses HiGHS (an open‑source linear programming solver), not some heuristic.

3. The same model on two chips with identical reported VTCM (where NPUs store weights) can show a 33× difference in DDR (regular RAM) traffic. Why? I have no clue; Qualcomm never published VTCM capacity for any of their SoCs.

4. There's an undocumented internal simulator called Hextimate that prices ops without the hardware.

[The Whole Writeup](https://datavorous.github.io/writing/qairt/)

Also, the VTCM memory placement works like a 3‑D bin‑packing - each block gets a tuple of three u32 integers. That shocked me because everyone assumed it was just a linear memory lane. The allocator (they literally call it `FancyAllocator`) even uses recursive backtracking.

I've also written about [NPUs before](https://datavorous.github.io/writing/npu/). Maybe I'm the only one bridging the gap between marketing fluff and academic literature on NPUs on the internet XD

Hope this helps folks using the Qualcomm DevKit (like me!) to get past the itch.

Edit: added full forms

[Raj]

Absolute wizardry! You're basically writing all the NPU literature on the internet by yourself.

[Komal]

I didn't get what you meant.

[Basavaraj]

Yeah, cool stuff. I tried to dive deep into NPUs once during college, but man, the documentation was almost zero.

Back then I wanted to build a stock‑exchange engine from scratch in C++. Even in this AI era it took ages because there's hardly any docs for that kind of high‑frequency‑trading stuff.

[Hari]

Did you use AI for the reverse engineering?

[Shivendra]

One tip: when you write about a niche, mention the full form of abbreviations at least once. It makes the content a bit more readable and understandable.

[Arisha]

Did you use multiple devices?

[Norman]

Thanks a lot, mate. I've tried running the Gemma‑4 E2B model on my phone before because it had 12 GB RAM and the mid‑range processor claimed an "AI processing unit", but it failed. I fell for their marketing hype, so it's great that you're testing their claims.

[Meera]

This was a really interesting read. The memory‑allocation part is new to me - instead of a 1‑D view it's almost like a 3‑D matrix. Shows I still have a lot to learn. I'll get back to my ESP32s for now.

[Avni]

I've worked on Qualcomm's NPUs before and their docs are a nightmare. Even when we asked engineers inside Qualcomm who were the POC, we kept getting "we don't know" answers. What you've done is really valuable.

[Pooja]

I've been working with low‑level ML systems like Blackwell optimisation, experimenting with different DSLs and kernels, so this caught my interest.

I appreciate your write‑up - their docs are a black hole. For anyone wanting to dig deeper, check the Qualcomm site OP mentioned to see what SNPE, QAIRT, etc., stand for; a quick scan helps you know the landscape.

To me it looks like an inference engine that does system‑level optimisations. The blog feels fairly complete.

Any ML‑systems discussion needs both the architecture and the DSL. Qualcomm isn't forthcoming about the architecture. That 3‑D matrix could have been explained if we knew their hierarchy (like CUDA core → warps → warp groups → SMs). That's why the hexagons stay a black box. Nice that you discovered they use systolic arrays.

That could also explain the memory layout and whether they use access‑pattern tricks like swizzling.

For OP, enable the web‑search tool and give Claude Code another try with those .dlc files. Maybe run a separate Claude Code call to decompose the NPU architecture. It's possible they have specialised hardware like tcgen05 on Blackwell, or separate scalar and tensor units like on Google TPUs.

Qualcomm should be more open about this. People are keen to optimise inference kernels - think of the optimized kernels for Apple Silicon on MLX - so exposing NPUs would boost adoption.

They do provide some C++ and Java tooling to serve compiled models as apps, but again, check the website to see what I'm referring to.