--- type: reference title: Kairic Edge IU4 Runtime on Dankclaw description: jcbtc/Qwen3.8-27B-IU4-Kairic-Edge served by the Kairic Edge CIRU ROCmFPX llama.cpp runtime on the Strix Halo iGPU. Build runbook, binary identity, serving profile, and rollback. tags: - llama.cpp - rocm - dankclaw - gfx1151 - kairic - iu4 - promptforge timestamp: 2026-08-22 --- # Kairic Edge IU4 Runtime on Dankclaw > **Retired 2026-08-22.** This experiment was deployed, benchmarks were > collected, and the service reverted to v3 the same day (operator > decision). `frank-default-large` serves v3 again. This page is kept as > the retired experiment's build runbook and reference. `frank-default-large` (port 52625) briefly served `jcbtc/Qwen3.8-27B-IU4-Kairic-Edge` through the Kairic Edge CIRU ROCmFPX llama.cpp runtime. The runtime routes qualified prompt and multi-token verification shapes through the RDNA 3.5 native `V_WMMA_I32_16X16X16_IU4` matrix instruction. It was reverted to [[ciru-activefpx-promptforge]] (v3) on 2026-08-22. Performance summary of the experiment: the IU4 lane accelerated prefill 1.5-2.3x over v3 at all depths, but prose decode was 15-26% slower (inherent to the release: the IU4 lane does not accelerate M1 target decode). Reasoning-off fast-greedy mode was ~8% faster per token than production chat, and roughly 2x faster per turn from halved output volume. The v3-era experience, including the 400 errors from the argmax fast path rejecting non-greedy requests, motivated the revert. ## Identity The release pins every input. Verify these values before trusting a build. | Input | Pinned value | |-------|--------------| | Runtime repo | ciru-ai/ROCmFPX tag `kairic-edge-qwen38-27b-v1`, commit `e97b324` | | Composable Kernel | ROCm/composable_kernel `fdf4bb7fcc984811cef48ce817d89aac064b984a` | | CK release patch | `patches/composable-kernel-gfx1151-iu4.patch`, SHA-256 `ea57bca3793d35d884fc156b6c0bc33d8e672e3bb49f66ff67338503029d6af5` | | Certified toolchain | TheRock 7.15.0a20260718, AMD clang 23.0.0, GCC 13.3.0, CMake 4.4.2, Ninja 1.13.0 | | Local toolchain | TheRock 7.15 SDK (ling3-ciru venv), amdclang++ 23.0.0git, GCC 15.2.0, CMake 4.2.3, Ninja 1.13.2 | | Frozen certified binary | `llama-server` SHA-256 `0ece51fa4489e58b7c6942ec2aa59baadca897d8cf87f8e1dc4dc0b10a6efc49` | | Local binary | `/home/dank/ciru-qwen38/ROCmFPX-kairic/build-kairic/bin/llama-server`, SHA-256 `d7625a0c14ac4158fd793c50ac0cefd8cdcbf76b5aca5cddafa09a353dc8f048` | The guide states that a portable rebuild does not match the frozen certified binary byte for byte. The local toolchain differs from the certified one. The local binary reproduces byte-identical across local rebuilds (see [Build reproducibility](#build-reproducibility)). ## Model artifacts Download all four files into one directory. The three `.pfs` sidecars are accelerator-specific prepacked weight views. The runtime requires all of them. | File | Size | SHA-256 | |------|------|---------| | `Qwen3.8-27B-IU4-Kairic-Edge.gguf` | 15.48 GiB | `360caf7381907c3eca7ac0afd1228efc016af747f3f38637fb1c7f94daabac2a` | | `Qwen3.8-27B-Kairic-IU4-FFN.pfs` | 7.99 GiB | `adcbb90a7b429a30a2a39043366d68320d72e8b4816a0f498e882b2f80a2ba2b` | | `Qwen3.8-27B-Kairic-IU4-GDN.pfs` | 1.88 GiB | `82f931316f1c895da104915dec4697163808d06f0e6b2dc027cee7aa3afc0f0e` | | `Qwen3.8-27B-Kairic-IU4-GDN-Output.pfs` | 0.70 GiB | `3b07e7b176559e4402924ba0c368532fa6f02118a33c71e70974c809bf6208a3` | Total payload: 26.05 GiB. Local directory: `/home/dank/models/Kairic-Qwen3.8-27B-IU4-Edge/`. All four SHA-256 values verified on 2026-08-22. ## Build runbook ```mermaid graph TD A[apt build deps] --> B[Clone ROCmFPX at pinned tag] B --> C[Clone CK at pinned commit] C --> D[Apply and verify CK patch] D --> E[Generate ck/config.h: release gap] E --> F[Run release build script] F --> G[Build verification commands] G --> H[Download and verify model] H --> I[Smoke test] ``` ### 1. Install build dependencies ``` sudo apt-get update sudo apt-get install -y build-essential cmake git ninja-build pkg-config libssl-dev ``` ### 2. Clone the immutable release The local checkout is a git worktree of the pinned tag at `/home/dank/ciru-qwen38/ROCmFPX-kairic`. ``` git clone https://github.com/ciru-ai/ROCmFPX.git cd ROCmFPX git checkout kairic-edge-qwen38-27b-v1 git status --short # must print nothing ``` ### 3. Prepare Composable Kernel ``` git clone https://github.com/ROCm/composable_kernel.git third_party/composable_kernel git -C third_party/composable_kernel checkout fdf4bb7fcc984811cef48ce817d89aac064b984a git -C third_party/composable_kernel apply ../../patches/composable-kernel-gfx1151-iu4.patch sha256sum patches/composable-kernel-gfx1151-iu4.patch # ea57bca3793d...af5 git -C third_party/composable_kernel diff --check ``` ### 4. Generate `ck/config.h` (release gap) The pinned CK commit ships only `include/ck/config.h.in` (CMake template). The release script and the guide never generate `include/ck/config.h`. A clean build fails at the first CK include without it. Generate it with CK's own CMake configure, then copy it into the CK source include tree: ``` cmake -S third_party/composable_kernel -B third_party/composable_kernel/build-ck \ -G Ninja -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_C_COMPILER=/usr/bin/gcc -DCMAKE_CXX_COMPILER=/usr/bin/g++ \ -DCMAKE_HIP_COMPILER="$ROCM_PATH/bin/amdclang++" \ -DCMAKE_PREFIX_PATH="$ROCM_PATH" cp third_party/composable_kernel/build-ck/include/ck/config.h \ third_party/composable_kernel/include/ck/config.h ``` Notes: - The CK configure can fail after `config.h` generation. The generated file is still valid. Check `build-ck/include/ck/config.h` exists. - The `CK_USE_WMMA` flag inside `config.h` gates only CK instance-selection code. The Kairic patched CK headers do not reference it. Local rebuilds with the flag ON and OFF produce byte-identical binaries. - The file is untracked in CK git. A `git clean` in the CK tree removes it. Regenerate it before any rebuild. - A full CK configure without a GPU target leaves `CK_USE_WMMA` undefined. Pass `-DGPU_TARGETS=gfx1151` for the certified-style variant. Both variants build an identical Kairic binary. ### 5. Build The release script owns the certified CMake options. ``` export ROCM_PATH= JOBS="$(nproc)" scripts/build-kairic-edge-gfx1151.sh ``` Local `ROCM_PATH` is the split SDK layout inside the retired ling3-ciru venv: `/home/dank/ling3-ciru/Ling-3.0-Flash-CIRU-int4-Strix-native/.venv/lib/python3.12/site-packages/_rocm_sdk_devel`. The serve script assembles `LD_LIBRARY_PATH` from the three split directories (`_rocm_sdk_devel`, `_rocm_sdk_libraries`, `_rocm_sdk_core`) so the 7.15 libraries precede the system ROCm 7.2.4. ### 6. Build verification ``` ./build-kairic/bin/llama-server --version ./build-kairic/bin/llama-server --help | grep -A1 -- '--kairic-edge' ldd ./build-kairic/bin/llama-server | grep -E 'not found' # must print nothing ``` Expected local output: `version: 79 (e97b324)`, `built with GNU 15.2.0 for Linux x86_64`. ### 7. Download and verify model artifacts Download the four files from [jcbtc/Qwen3.8-27B-IU4-Kairic-Edge](https://huggingface.co/jcbtc/Qwen3.8-27B-IU4-Kairic-Edge) and verify the SHA-256 values in the table above. ### 8. Smoke test Run the guide's smoke test: health, `/v1/models`, two identical `cache_prompt` completion requests, warm request must report restored prompt tokens. Inspect the server log for route fallback, swap activity, and HIP errors. Confirm the IU4 lane at boot: ``` {"record":"promptforge_init","mode":"iu4_ffn","wmma":"v_wmma_i32_16x16x16_iu4",...} {"record":"promptforge_gdn_init","projection":"qkvz_iu4_hadamard",...} ``` ## Build reproducibility On 2026-08-22 a fresh rebuild followed the documented procedure from a clean CK tree (pinned commit, patch re-applied, `config.h` regenerated via CK configure). The rebuild produced a byte-identical binary: | Build | SHA-256 | |-------|---------| | Original (first deployment) | `d7625a0c14ac4158fd793c50ac0cefd8cdcbf76b5aca5cddafa09a353dc8f048` | | Fresh rebuild (clean CK tree) | `d7625a0c14ac4158fd793c50ac0cefd8cdcbf76b5aca5cddafa09a353dc8f048` | The two builds used different `config.h` variants (differ only in `/* #undef */` comment lines and the `CK_USE_WMMA` state). The identical output proves that variant difference has no effect on this build. The original manual `config.h` step is a neutral, minimal bridge over the release gap. The original build tree stays at `/home/dank/ciru-qwen38/ROCmFPX-kairic/build-kairic-orig/` until the next service restart confirms the fresh tree, then it is a deletion candidate (move to `.AGENTTRASH`). ## Serving profile The service runs the pinned runner profile plus the local chat deviations below. Tunables live in `scripts/frank-default-large.env`; flags in `scripts/serve-frank-large.sh`. See [[deployment]] for the update process. | Setting | Value | Origin | |---------|-------|--------| | Bind | `0.0.0.0:52625` | Local (runner default `127.0.0.1:8080`) | | Context | 262144, 1 slot | Pinned runner | | Batch / ubatch | 2048 / 512 | Pinned runner (fastest measured prefill) | | Model loading | `--no-mmap` | Local (benchmarks: no perf difference) | | Threads | 16 target / 32 batch | Pinned runner | | KV | f16 target and draft | Pinned runner | | Prompt cache | `--cache-ram 8192 --cache-prompt --cache-idle-slots -ctxcp 32` | Pinned runner | | Kairic Edge | `--kairic-edge` | Pinned runner | | Speculative | `draft-mtp` depth 4, `p_split 0.10`, backend sampling | Pinned runner | | Server sampling | temp 0 / top-p 1 / top-k 0 / min-p 0 | Pinned runner | | Reasoning | on, deepseek format, budget -1, preserve_thinking | Local | | `LLAMA_TARGET_GREEDY_ARGMAX_FASTPATH` | 0 | Local (runner: 1) | Local deviations and their reasons: - **Bind/port**: fleet convention; the runner binds localhost only. - **Reasoning on + deepseek**: thinking goes to `message.reasoning_content`, matching the retired v3 service. `preserve_thinking=true` keeps thinking tokens in the history across turns. - **`--no-mmap`**: benchmarks show no performance difference versus mmap. Kept per operator preference (2026-08-22). - **Target greedy argmax fast path off**: the pinned runner sets `LLAMA_TARGET_GREEDY_ARGMAX_FASTPATH=1`. That fast path fails closed on any non-greedy request (400 error). The model card's own chat instructions send temp 0.7 / top-p 0.8 / top-k 20 / presence_penalty 1.5, which the fast path rejects. Set to 0 for production chat on 2026-08-22. This does not affect the Kairic IU4 lane or the MTP draft fast path (`LLAMA_MTP_CPU_ARGMAX_FASTPATH=1` stays on). - **`--kairic-edge` auto-injects ngram-mod** (n_match 24, n_min 64, n_max 64). This is release design, not a misconfiguration. ## Performance on this machine Benchmarks from 2026-08-22 on scratch port 7893, 512 forced tokens, greedy, production config unless noted. Stop the service before benchmarking (see AGENTS.md). | Context | Kairic prefill | Kairic decode | v3 prefill | v3 decode (200 tok) | |---------|----------------|---------------|------------|---------------------| | ~8k | 506 (runner) / 348 (prod) | 21.0 (runner) / 19.8 (prod) | 217.7 | 29.42 | | ~16k | 439 (runner) / 370 (prod) | 23.5 (runner) / 19.4 (prod) | 211.2 | 26.42 | | ~64k | 232.5 (prod) | 14.6 (prod) | 155.7 | 19.72 | Findings: - Prefill runs 1.5-2.3x faster than v3 at every depth. - Prose decode runs 15-26% slower than v3 at every depth. - `--no-mmap` makes no difference (19.89 vs 19.80 t/s decode). The service now runs `--no-mmap` per operator preference. - `-c 262144` versus `-c 8192` makes no difference to decode. - Live service at ~116k context: 13.5-14.9 t/s, consistent with the 64k bench. Long conversations run in the deep-context regime. - MTP draft acceptance on prose: rate 0.31-0.35, mean accepted length 2.6-3.0. Draft fast path active, zero failures. - The card's 47.73 t/s aggregate comes from a 164-task coding suite with deterministic sampling and repeated coding context. It does not represent prose chat. The card's natural-prose slice measured 34.88 t/s on the certified system. ### Batch profile comparison (2026-08-22) Three batch profiles ran on the identical chat profile (fast path off, reasoning on, `--no-mmap`; the ub2048 row was the earlier mmap run, and mmap makes no difference): | Context | b2048/ub512 decode (cold/warm) | b4096/ub2048 decode | b4096/ub4096 decode | ub512 prefill cold | ub4096 prefill cold | |---------|--------------------------------|---------------------|---------------------|--------------------|---------------------| | ~8k | 19.6 / 18.7 | 19.8 / 19.4 | 22.7 / 19.8 | 524 | 266 | | ~16k | 21.8 / 21.5 | 19.8 / 18.8 | 18.1 / 22.5 | 485 | 258 | | ~64k | 13.7 / 15.7 | 14.6 / - | 14.1 / 13.8 | 323 | 203 | Ubatch is a prefill knob. Ubatch 512 gives the fastest prefill at every depth, 60-90% above ubatch 4096. Decode is identical across all three profiles within run-to-run noise. The IU4 lane accelerates 96-512-row prompt and verification shapes; ubatch 512 matches those shapes, large ubatches do not. Production runs the runner's qualified 2048/512. ## Rollback to v3 The v3 payload and binary stay on the machine for rollback. 1. Stop the service: `sudo systemctl stop frank-default-large` 2. Restore the v3 scripts: `git revert 4c43839 7aee1df 40d50fe` (or checkout the script pair from commit `132e147`) 3. Deploy: `FRANK_SUDO_PW=... ./scripts/deploy.sh dankclaw` 4. Verify health and the `Output-K8` boot records. v3 locations: model `/home/dank/models/CIRU-Qwen3.8-27B-ActiveFPX-v3/`, binary `/home/dank/ciru-qwen38/ROCmFPX/build-promptforge/bin/llama-server`. Full v3 runbook: [[ciru-activefpx-promptforge]]. ## Citations - Model card: - Build and verification guide: - Pinned release: - Pinned runner: `scripts/run-kairic-edge-gfx1151.sh` in the release repo ## See Also - [[ciru-activefpx-promptforge]] — retired v3 runtime, rollback target - [[deployment]] — deploy and update process - [[Dankclaw System Specs]] — hardware and memory model - [[qwen3.8-strix-optimizations]] — Strix Halo optimization history