serverAppConfigs/serverAppConfigs-docs/LLM/kairic-edge.md

---
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-large` serves v3 again. This page is kept as
> the retired experiment's build runbook and reference.

`frank-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=<path to the 7.15 SDK _rocm_sdk_devel directory>
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-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-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: <https://huggingface.co/jcbtc/Qwen3.8-27B-IU4-Kairic-Edge>
- Build and verification guide: <https://github.com/ciru-ai/ROCmFPX/blob/kairic-edge-qwen38-27b-v1/docs/kairic-edge-gfx1151.md>
- Pinned release: <https://github.com/ciru-ai/ROCmFPX/tree/kairic-edge-qwen38-27b-v1>
- 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