Using AMD's hawkpoint NPUs to run LLMs!! its a llama.cpp plugin not a standalone runtime to ensure supporting most models, but it has a few cool kernels targeted at the the xdna1 arch bcs as it turns out, the support for xdna1 NPUs is like really dead everyones like xdna2 yay, but i hv a xdna1 CPU so ima use every bit of it ig lol
Run GGUF quantized models directly on AMD XDNA1 NPUs (Ryzen AI 7040 / 8040 series) with llama.cpp
its a fork or bettter said plugin of llama.cpp specially made to run models using amd NPUs
Using AMD’s hawkpoint NPUs to run LLMs!! its a llama.cpp plugin not a standalone runtime to ensure supporting most models, but it has a few cool kernels targeted at the the xdna1 arch bcs as it turns out, the support for xdna1 NPUs is like really dead everyones like xdna2 yay, but i hv a xdna1 CPU so ima use every bit of it ig lol
7h 47m 9s logged
Finished multi-model crossover validation across three scales (0.5B, 3B, 27B):
// Adaptive Heterogeneous Scheduling
if (N <= 256) {
// Tiny matrices (<60 KB) run on CPU SIMD to avoid driver dispatch overhead
cpu_engine->run(weights, weight_bytes, N, K, act, out);
} else {
// Medium/Large matrices stream to 16 AIE2 compute tiles
xdna_engine->run(weights, weight_bytes, N, K, act, out);
}
Shipped v1.0.0 with standalone binaries and clean architecture docs!
12h 48m 59s logged
Architected the fix for the 27B memory collapse: a bounded double-buffered staging ring in src/xdna-gemv-engine.cpp.
// Allocate exactly two 64 MB host-only staging buffers (128 MB total pinned memory)
static constexpr size_t STAGING_BO_SIZE = 64 * 1024 * 1024;
bo_staging[0] = std::make_unique<xrt::bo>(*device, STAGING_BO_SIZE, xrt::bo::flags::host_only, 0);
bo_staging[1] = std::make_unique<xrt::bo>(*device, STAGING_BO_SIZE, xrt::bo::flags::host_only, 0);
// Double-buffered stream from mmap page cache
size_t stage_idx = current_op % 2;
std::memcpy(bo_staging_map[stage_idx], mmap_weight_ptr, weight_bytes);
kernel_dpu->operator()(..., *bo_staging[stage_idx], ...);
flags::host_only: Limits pinned kernel memory strictly to $2 \times 64\text{ MB} = \mathbf{128\text{ MB}}$.mmap_weight_ptr: Weights stay in standard page cache (MAP_SHARED), freely managed by Linux VM.Result: pgmajfault and swap dropped to 0. 27B decode speed jumped from 0.15 tok/s to 1.72 tok/s—an 11.5× end-to-end recovery.
12h 13m 38s logged
Tried loading Qwen3.8-27B-Q4_0 (15.2 GB). Everything collapsed: token speed plummeted to 0.15 tok/s (6.67 seconds/token).
Ran /proc/meminfo and /proc/vmstat diffs during decode to find the bottleneck:
/proc/meminfo during 27B crash:
MemTotal: 31,842 MB
Unevictable: 12,640 MB <-- 369 persistent XRT BOs locked in RAM!
pgmajfault: 520,118 <-- 520k disk page reads PER TOKEN
pswpout: 18,400 <-- active swap storm
xrt::bo objects for all 369 model tensors pinned 12.6 GB into Unevictable kernel memory.Persistent BO allocation for large models is fundamentally broken on consumer RAM.
9h 43m 3s logged
Built tools/bench-aie-gemv-standalone.cpp to measure pure hardware DMA bandwidth without graph scheduler overhead.
auto t0 = std::chrono::high_resolution_clock::now();
bo_weights->sync(XCL_BO_SYNC_BO_TO_DEVICE);
kernel_run.wait();
auto t1 = std::chrono::high_resolution_clock::now();
double gb_s = (bytes / (dt_ns * 1e-9)) / 1e9;
sync(TO_DEVICE): Flushes dirty CPU lines and triggers DMA transport.kernel_run.wait(): Waits for AIE vector compute completion.Tested synthetic matrix sizes out to 1 GB ($K=5120$):
Linear fit: $T(\text{bytes}) \approx 0.15\text{ ms} + \frac{\text{bytes}}{38.5\text{ GB/s}}$. The ~150 µs startup overhead gets amortized as matrix size grows.
11h 12m 31s logged
Standard GGUF interleaves nibbles row-by-row. But AIE2 vector tiles need weights split into 4 independent planar column channels aligned to 64-byte boundaries.
Wrote xdna-q4-prepack.cpp to deinterleave GGUF Q4 weights:
// 4-Column Planar Deinterleaving
for (size_t col = 0; col < 4; ++col) {
uint8_t* dst_col = packed_out + col * col_bytes;
for (size_t row = 0; row < N; ++row) {
const uint8_t* src_nibbles = src_weights + row * row_stride + col * (K / 4) / 2;
deinterleave_and_align64(src_nibbles, dst_col + row * 64, K / 4);
}
}
col * (K / 4): Strips matrix columns across the 4 independent NPU memory tiles.64-byte alignment: Direct DMA streaming into tile L1 SRAM without unaligned stalls.aie::unpack converts uint4 to BF16, scales, and computes fused vector MACs.First interactive generation on physical NPU! Cosine similarity against CPU FP32 oracle: 0.9999.
13h 5m 32s logged
Before writing custom AIE kernels, we needed an unyielding golden oracle to check numerical correctness and see how fast the host CPU is.
Wrote an unrolled AVX2/FMA Q4_0 vector engine in src/cpu-q4-gemv-engine.cpp:
const __m256i raw = _mm256_loadu_si256(reinterpret_cast<const __m256i*>(w_ptr));
const __m256i lo_nibbles = _mm256_and_si256(raw, _mm256_set1_epi8(0x0F));
const __m256 w_f32 = _mm256_cvtepi32_ps(_mm256_sub_epi32(lo_nibbles, _mm256_set1_epi32(8)));
acc = _mm256_fmadd_ps(_mm256_mul_ps(w_f32, scale_vec), act_vec, acc);
_mm256_and_si256(..., 0x0F): Masks out the low 4-bit nibbles from 32 packed weights._mm256_sub_epi32(..., 8): Subtracts the Q4_0 symmetric offset (8)._mm256_fmadd_ps: Fused multiply-accumulate with FP32 activations in single-cycle throughput.Benchmarked Qwen2.5-0.5B on CPU: 106.5 tok/s (9.38 ms/tok) on 8 threads. Modern Zen 4 AVX2 is blazing fast on small models, so the NPU has a high bar to clear.
8h 1m 44s logged
why have random shi like gate-1? bcs xdna1 is the worst supported from amd, idk their docs are incomplete im like testing and figuring stuff out
Gate 1 milestone: getting real AIE execution contexts spinning via AMD XRT.
Initial prototype had a huge bug: we were creating an xrt::hw_context and reloading the .xclbin bitstream inside the operator compute loop. Profiling showed this added ~14 ms of pure driver overhead per tensor dispatch—which murdered performance before we even started doing math.
Fixed it by making the context a persistent backend singleton:
// Initialize context & load bitstream ONCE at startup
device = std::make_unique<xrt::device>(device_index);
xrt::uuid uuid = device->register_xclbin(xclbin_img);
hw_ctx = std::make_unique<xrt::hw_context>(*device, uuid);
xrt::module mod(elf_img);
kernel_dpu = std::make_unique<xrt::kernel>(xrt::ext::kernel(*hw_ctx, mod, "DPU"));
device->register_xclbin: Flashes the DPU bitstream once into the NPU fabric.hw_ctx: Persistent hardware context reused across all tensor iterations.Dispatch overhead instantly dropped from 14 ms down to ~38 µs.
12h 0m 11s logged
So the plan is simple: get GGUF models running straight on the Ryzen AI NPU without touching bloated ONNX or Vitis runtimes.
First step was checking if we could even talk to /dev/accel/accel0 directly without root. Wrote a quick probe calling the kernel’s amdxdna DRM driver ioctls:
struct amdxdna_drm_get_info info{};
info.param = DRM_AMDXDNA_GET_INFO_TOPOLOGY;
ioctl(fd, DRM_IOCTL_AMDXDNA_GET_INFO, &info);
DRM_IOCTL_AMDXDNA_GET_INFO: Queries the hardware directly from the kernel.info.param = TOPOLOGY: Confirmed our Phoenix chip has 4 columns × 4 compute rows = 16 AIE2 tiles.Instead of dragging in a 10 GB SDK, we wrote our own include/xdna/amdxdna_uapi.h directly matching the kernel driver structs. Zero bloat, instant builds.
