Model-to-NPU Pipeline for Qualcomm Snapdragon

Documentation: English · Русский · Project State · Android APK

The world's first fully functional on-device Stable Diffusion XL (SDXL) pipeline running natively on the Qualcomm Hexagon NPU (Snapdragon 8 Elite / OnePlus 13) without cloud servers, without Termux, without Root, and without model splitting.

---

⚡ What's New in v0.6.0-core

1. 🚀 72.7% LPDDR5X Memory Bandwidth Roofline Achieved

• 7.13 TOPS sustained across all 362 layers of SDXL 2.57B (W8A16) on Hexagon V79 HTP.
• UNet single pass latency dropped to 880.3 ms (down from ~956 ms and ~2411 ms historically).
• 72.7% of uncompressed LPDDR5X DMA bandwidth ceiling reached (~640 ms theoretical memory bottleneck at ~58 GB/s). Compute cores (HMX tensor & 1024-bit HVX vector units) maintain significant headroom.
• 8-step generation UNet time dropped to 11.59 seconds!
• Total Cold Start generation (including model deserialization from UFS 4.0 flash) is ~20–21 seconds.
• Host overhead is virtually eliminated (< 0.7%): ARM64 NEON FP16 MLP FMA (3.4 ms), 64KB L1 double-buffer copy to RPCMEM (2.2 ms, ~49.2 GB/s), FastRPC transport (1.7 ms).
• NPU Pipeline Duty Cycle = 98.7%!

2. 🧩 Monolithic W8A16 UNet (No More Model Splitting)

• The UNet model is no longer split into separate encoder and decoder halves!
• A single unified context (unet.serialized.bin, ~2.44 GiB) runs directly in NPU memory.
• Eliminated 11 skip-connection memory buffering bottlenecks, process synchronization delays, and split-graph IPC overhead.

3. 📐 Dynamic Arbitrary Resolution Engine (0.26 MP to 2.36 MP)

• Arbitrary aspect ratios and resolutions from 512×512 up to 1536×1536 / 1344×1728 on the fly without recompiling models or switching contexts.
• Centered Spatial-CFG Sub-Canvas Framing: The active image window is placed at the optical center (512, 512) with spatial CFG masking (w_CFG = 3.5 active, transitioning to 1.0 uncond at borders).
• 0% extra NPU latency overhead compared to standard square.
• Continuous noise variance preserves 100% GroupNorm numerical stability (zero NaNs).
• Completely eliminates edge reflection artifacts and bilateral symmetry (Rorschach mirror seams).
• Integrated Catmull-Rom bicubic reconstruction + Contrast-Adaptive Sharpening (CAS, 59 ms).

4. 📱 Zero-Root, Zero-Termux, Zero-Python Architecture

• No Root required.
• No Termux or Python runtime required on the phone for inference.
• Standalone native C inference engine (qnn-multi-context-server) handles CLIP, Monolithic UNet, and VAE directly via QNN System & Backend APIs.
• Can be executed with a single command via ADB or embedded into an Android APK.

5. 📲 Native Android APK (v0.6.4)

• Standalone Zero-Root & Zero-Termux Android application (Releases) packing the monolithic C-engine, Progressive CFG (dCache), Dynamic Resolution, Dynamic LoRA selector, and strict industrial dark UI.

---

🖼️ Gallery

All gallery samples above are 1024×1024 outputs from the Lightning-merged SDXL path running directly on-device.

---

📱 Proof that it actually runs on-device

Earlier public screenshot — 273.6s total

v0.2.0 public marker — 100.8s total

v0.2.3 screenshot (Live Preview ON) — 78.0s total

v0.4.7 cold-start APK proof — 34.6s total Measured accelerator-visible time inside this run: ~16.25 s.

v0.6.4 Zero-Root APK proof — 20.38s total CLIP: 158 ms · UNet (8 steps, dCache): 11.28 s · VAE: 2.32 s.

On-Device Telemetry & Milestone Progression

Public screenshot lineage so far: 273.6 s → 100.8 s → 78.0 s → 34.6 s → 20.38 s (11.28 s UNet 8 steps)!

================================================================================
  OnePlus 13 (Snapdragon 8 Elite / Hexagon V79 HTP) — On-Device Execution Log
================================================================================
  [Init] QNN backend + system loaded in 142ms
  [Load] CLIP-L + CLIP-G loaded in 312ms
  [Load] UNet Monolithic W8A16 (2.44GB) loaded in 5821ms
  [Load] VAE Decoder FP16 loaded in 284ms
  [CLIP] Text encode finished: 281ms
  [Zero-Copy Denoise] Running autonomous 8-step monolithic UNet in NPU memory...
  [UNet 1/8] 880.3ms (CFG active)  | 6 HVX threads @ 7.13 TOPS
  [UNet 2/8] 879.8ms (CFG active)  | 6 HVX threads @ 7.13 TOPS
  [UNet 3/8] 878.9ms (CFG active)  | 6 HVX threads @ 7.13 TOPS
  [UNet 4/8] 876.5ms (CFG active)  | 6 HVX threads @ 7.13 TOPS
  [UNet 5/8] 874.2ms (CFG dCache)  | 6 HVX threads @ 7.13 TOPS
  [UNet 6/8] 874.1ms (CFG dCache)  | 6 HVX threads @ 7.13 TOPS
  [UNet 7/8] 874.0ms (CFG dCache)  | 6 HVX threads @ 7.13 TOPS
  [UNet 8/8] 874.5ms (CFG dCache)  | 6 HVX threads @ 7.13 TOPS
  [Zero-Copy Denoise] UNet total: 11,590ms (880.3ms/pass, Duty Cycle: 98.7%)
  [VAE] Decode: 2,204ms (FP16)
  [Save] PNG written to storage in 189ms
  Total Cold Generation: ~20.9s | Warm Generation: ~14.1s
================================================================================

---

📊 Performance Benchmarks (Snapdragon 8 Elite / OnePlus 13)

UNet Pass Latency Breakdown (1024×1024)

Stage	Hardware Unit	Latency	Share	Details
Host temb MLP	Oryon CPU (ARM NEON)	3.43 ms	0.4%	FP16->FP32 FMA (vld1q_f16, vfmaq_f32)
Host -> RPCMEM ION	System Memory Bus	2.20 ms	0.2%	64 KB L1 double copy (~49.2 GB/s)
Quant & Transpose	Oryon CPU	0.41 ms	0.0%	NEON vectorized quantization
FastRPC Transport	ARM <-> Hexagon Bus	1.72 ms	0.2%	ION buffer synchronization
Pure Hexagon V79 Accel	Hexagon V79 HTP (NPU)	880.31 ms	98.9%	6 HVX hardware threads unlocked
Total UNet Pass	—	888.07 ms	100%	NPU Duty Cycle: 98.7%

Historical Progression

Version	Total Time	UNet (8 steps)	CLIP	VAE	Architecture & Pipeline
v0.1.0	273.6 s	—	—	—	Initial public proof-of-concept
v0.1.3	104.4 s	91.5 s	2.0 s	9.0 s	Added mmap
v0.2.0	79.7 s	72.4 s	2.9 s	3.4 s	Sustained high performance mode
v0.2.5	75.6 s	66.6 s	2.8 s	3.0 s	Native accel helper, per-step qnn-net-run
v0.3.0	30.4 s	19.3 s	2.8 s	1.9 s	Persistent QNN server, split UNet (enc+dec)
v0.4.7	34.6 s	14.2 s	0.1 s	1.8 s	APK cold start marker (split UNet)
v0.6.0-core	20.9 s	11.59 s	0.28 s	2.20 s	Monolithic W8A16 UNet, 72.7% Memory Roofline, 6 HVX threads
v0.6.4 (APK)	20.38 s	11.28 s	0.16 s	2.32 s	True Zero-Root & Zero-Termux APK, Progressive CFG (dCache), Dynamic LoRA

Note: v0.6.4 total time (20.38 s) is a cold-start APK measurement including deserializing 2.44 GiB of model weights from UFS 4.0 flash storage. Subsequent warm generations run in ~14 s total!

Dynamic Resolution Benchmarks (Centered Spatial-CFG)

Target Resolution	Aspect Ratio	Megapixels	UNet (8 steps)	Wall Time (Cold)	Notes
1024 × 1024	1:1 Square	1.05 MP	11.59 s	21.06 s	Reference ground truth
832 × 1216	9:16 Portrait	1.01 MP	11.64 s	20.64 s	Perfect anatomy, zero edge seams
1024 × 768	4:3 Landscape	0.79 MP	11.91 s	20.74 s	High detail landscape
1344 × 1728	3:4 Hi-Res	2.32 MP	11.21 s	20.03 s	Ultra-crisp Catmull-Rom CAS

---

🛠️ Quick Start & Usage

Hardware & Software Requirements

• Target Device: Qualcomm Snapdragon 8 Elite (OnePlus 13 or similar) with Hexagon V79 HTP.
• RAM: 16 GB LPDDR5X (Peak memory footprint ~3.5 GB for monolithic pipeline).
• Storage: ~8 GB UFS 4.0 storage on device.
• Root: Not required!
• Termux: Not required!

1. Build the Native Server (Host PC)

Requires Android NDK (r26+) and Qualcomm QAIRT SDK (2.31+):

python scripts/build_qnn_multi_context_server.py --deploy

2. Push Models & Assets to Phone

adb push D:/platform-tools/sdxl_npu/unet.serialized.bin /data/local/tmp/sdxl_qnn/
adb push D:/platform-tools/sdxl_npu/vae_decoder.serialized.bin /data/local/tmp/sdxl_qnn/
adb push D:/platform-tools/sdxl_npu/clip_l.serialized.bin /data/local/tmp/sdxl_qnn/
adb push D:/platform-tools/sdxl_npu/clip_g.serialized.bin /data/local/tmp/sdxl_qnn/
adb push D:/platform-tools/sdxl_npu/tokenizer /data/local/tmp/sdxl_qnn/

3. Generate Image on Phone via ADB

adb shell "LD_LIBRARY_PATH=/data/local/tmp/sdxl_qnn/lib /data/local/tmp/sdxl_qnn/bin/qnn-multi-context-server \
  --backend libQnnHtp.so \
  --system libQnnSystem.so \
  --generate-dyn \
  --prompt 'masterpiece, 1girl, cyberpunk aesthetic, neon city, highly detailed' \
  --width 832 --height 1216 --steps 8 --cfg 3.5 --seed 42 \
  --output /sdcard/Download/output.png"

---

📁 Repository Structure

• NPU/qnn_multi_context_server.c — The complete standalone C inference engine (NEON, RPCMEM, 6 HVX threads, Spatial-CFG).
• scripts/build_qnn_multi_context_server.py — Host build script compiling with Android NDK Clang.
• phone_generate.py — Standalone Python entrypoint for debugging and evaluation.
• PROJECT_STATE.md — Complete engineering record, memory roofline calculations, and architectural notes.
• APK/ — Android Studio project for the standalone Zero-Root mobile app (v0.6.4).
• SDXL/ — SDXL conversion, calibration, and ONNX graph manipulation tools.
• WAN 2.1 1.3B/ — WAN research workspace and video diffusion tools.

---

📜 License

This project is licensed under the PolyForm Noncommercial License 1.0.0.
You are free to use, modify, and distribute this software for non-commercial and research purposes.