Devlog 08 — Aug 31: The Crossover Spectrum & v1.0 Ship

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);
}

Measured Crossover Scoreboard

• Qwen2.5-0.5B (403 MB): CPU favored (106.5 t/s vs 72.6 t/s XDNA hybrid) due to small 60 KB payloads.
• Qwen2.5-3B (1.86 GB): XDNA wins decisively: 16.5 tok/s vs 9.8 tok/s CPU (+68.8% / 1.69× speedup).
• Qwen3.8-27B (15.2 GB): XDNA sustains 1.72 tok/s streaming 12.6 GiB smoothly on a 32 GB laptop.
• Numerical Quality: Verified 100% Top-1 token match across all 3 scales.

Shipped v1.0.0 with standalone binaries and clean architecture docs!

Devlog 07: Bounded 128MB Staging Ring & The 11.5x Recovery

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.

Devlog 06 The 27B Disaster & 12.6 GB Unevictable BO Crisis

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

• The bug: Allocating persistent xrt::bo objects for all 369 model tensors pinned 12.6 GB into Unevictable kernel memory.
• The result: On a 32 GB laptop with OS and KV cache, Linux was starved of clean page cache. Every layer forced synchronous swap-outs and major page faults to disk.

Persistent BO allocation for large models is fundamentally broken on consumer RAM.

Fixed a few Dense path issues, because the MoE path somehow made the dense path completely usable, fixed a few small one line bugs i got to know from the last ship’s reviews, read to for ship 2!!

Devlog 05: Standalone GEMV Ceiling: Sustaining 38.5 GB/s

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$):

• 16 MB: 31.87 GB/s (0.53 ms)
• 128 MB: 34.27 GB/s (3.92 ms)
• 1024 MB: 38.57 GB/s (27.84 ms)

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.

Devlog 04: 4-Column Planar Prepacking & First Real AIE Run

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.
• Inside AIE: 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.

Hitting 100hrs and final ship

Didnt get time to work on this idea, since quite the while.

And it has js about reached the state i can probably get it to MoE arch works, dense models work, cpu/gpu streaming works well, context/kv cache is also dynamically quanted lm_head quants are also there now

The motivation really died out because of this project not getting reviewed its been over a month, leaving this as is, tho its not perfect, but it js about has enough features that i wanted

Here are the final results after a ton of benchmarking-

• Phi-4: ThinTensor 28.21 tok/s vs llama.cpp 17.50 tok/s — 1.61× faster (+61.2%). Quality: 0.999949 cosine vs 0.997281.

• OpenReasoning-Nemotron-14B: ThinTensor 23.79 tok/s vs llama.cpp 19.79 tok/s — 1.20× faster (+20.2%). Quality: 0.999533 vs 0.967140.

• Llama 3.2 11B: ThinTensor 41.02 tok/s vs llama.cpp 39.82 tok/s — 1.03× faster (+3.0%). Quality: 0.995095 vs 0.991938.

• Qwen3.5-9B: ThinTensor 43.28 tok/s vs llama.cpp 42.83 tok/s — 1.01× faster (+1.1%). Quality: 0.973334 vs 0.894908.

• GPT-OSS-20B: ThinTensor 50.34 tok/s vs llama.cpp 40.36 tok/s — 1.25× faster (+24.7%). Quality: 0.983299 vs 0.792350.

*all are apples to apples btw

Overall: ThinTensor beats llama.cpp Q4 throughput on 5/5 tested architectures, with a geometric-mean speedup of approximately 1.203×.

Probably gonna take a break, and learn before comming back to this, after the ship-2 on stardance!

Made it ready to ship!

pushed a linux binary on the releases page!
and wrote a good readme.md and made the repo structure better

Devlog 03: The CPU Baseline (AVX2 / FMA Golden Oracle)

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.

TUIEmu devlog #1 - two consoles, one terminal, zero GPU

the cores i did NOT write

writing a NES or Game Boy core from scratch is months of work. i am not
doing months of work in a weekend. so TUIEmu runs two vendored cores and
the whole job was the glue around them:

• PeakRacing/nes for the NES - hooks for memory-mapped ROM loading,
  ARGB8888 framebuffer output, sound off, every mapper family enabled
• deltabeard/peanut-gb for the Game Boy - clean, tiny, fast to bind

each core needs a port layer: boot the console once, hand it a
framebuffer, feed it a joypad every frame, run one frame, read pixels
back. that’s the whole engine loop in main.c, and honestly the fastest
way to “have two emulators” is to steal the two hardest parts and write
the fun glue yourself.

the render engine - this is where the real work happened

the cores give us a raw framebuffer (256x240 NES, 160x144 GB). the job:
get that onto a terminal without a GPU. the version history is basically
a pyramid of terminal graphics tricks:

1. half-blocks - every cell becomes a ▀ with the top half in fg
   color and the bottom in bg. 1x2 pixels per cell. the baseline.
2. quadrants - a 2x2 truecolor pixel per cell using block-element
   glyphs (▖▗▘▝▀▄▌▐█). every 1/2-color pattern maps to an exact glyph,
   and messy 3-4 color cells fall back to a diagonal or full block. no
   black holes: empty sub-squares inherit their neighbor’s color.
3. kitty graphics protocol - real pixels, base64-uploaded per frame
   and replaced in place. the terminal does the scaling.
4. DEC sixel - the OG 80s-era raster format. palette capped at 256
   sorted colors, 6-row bands, run-length slices, one DCS per frame.

the launch screen fires a capability probe: kitty q=1/q=2 plus the DA1
query (“what raster do you speak, terminal?”) and each mode quietly
auto-skips if the terminal can’t do it. mode keys cycle
half-block -> quadrant -> kitty -> sixel, everything fit-to-screen and
zoomable, nothing ever cropped or distorted.

the infinite hang

this project has one legendary bug. when all three probe replies landed
in a single read burst, the DA1 parser spun forever - the param scan
ended with q += strcspn(q, ";0123456789"), the digits are in the stop
set, so it advanced zero bytes, forever. the emulator just froze
silently. fixed with an explicit digit-walk loop. terminal reply inputs
now also can’t masquerade as keys (a stray kitty “OK” used to hit zoom).

input, the boring-but-vital part

keys -> gamepad: Q=A, E=B, arrows/WASD = directions, Enter = Start,
Tab = Select. directions are latched continuously while held - the core
reads the pad every frame and gets “pressed” every frame. no stutter.

heres contra running

Devlog 02: Hardware Smoke Tests & The 14ms Trap

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.

Devlog 01: Talking to the Silicon (DRM & Phoenix Discovery)

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.

Fixed vercel hosting hell, the whole thing js didnt want to work for soon long, small fix not much, not writing a huge devlog, because the ships already done, and probably wont be working on this again moving forward

devlog 3: debugging hell and nuking the UI

bruh the last 10 hours were an absolute nightmare. the web UI preview was completely broken. you’d hit play, and it would just flash and disappear. i was losing my mind trying to figure out why the terminal was instantly clearing itself.

turns out, the core engine’s TerminalSession was doing its job too well. when the animation finished, it sent the ANSI sequence to exit the alternate screen buffer, which is exactly what you want in a real terminal. but in xterm.js, exiting the alternate screen just leaves you with a blank black box!

had to wire up a --keep-screen flag all the way through the CLI architecture to intercept that sequence when running inside the web preview.

# src/clearfx/engine/terminal.py
def __exit__(self, exc_type, exc_val, exc_tb):
    if not self.keep_screen:
        # restore alternate screen
        self.write(b"\033[?1049l")

then there was the resizing bug. xterm.js was fitting to the container, but the PTY backend was hardcoded to 80x24. so it looked super misaligned. had to intercept term.cols and term.rows in React and pass them through the WebSocket JSON payload.

wsRef.current.send(JSON.stringify({ 
  type: 'run', 
  code, 
  width: termRef.current.cols, 
  height: termRef.current.rows 
}));

once the bugs were dead, i looked at the web UI and realized it looked like generic SaaS slop. huge glowing pills, purple gradients everywhere… it was dog shit.

so i nuked index.css entirely. built a custom design system inspired by factory.ai. pure dark operations console vibes. strictly monochrome, 8px spacing grid, geist monospace fonts for technical labels, and tiny orange/green signal dots instead of giant colored backgrounds.

/* replaced the generic slop with this */
:root {
  --bg-canvas: #101010;
  --bg-surface: #1D1A18;
  --border-hairline: #34312F;
  --signal-orange: #EE6018;
  --font-mono: 'Geist Mono', monospace;
}

the hero section now has a literal UI-in-UI terminal preview embedded in a custom .product-frame. the whole app feels like a precision engineering tool now, not a marketing page.

we finally shipped a safe, strictly declarative, gorgeous terminal engine.

devlog 2: nuking the backend and building a safe package format

so after the engine was running, i needed a marketplace. initially i spun up a whole fastAPI backend, but maintaining that sounded like a nightmare. so i ripped it all out and moved the entire catalog to firebase.

but here’s the catch: if i’m letting randos upload python animations to a public marketplace, they could literally just inject os.system("rm -rf /") inside their design.py and wipe everyone’s machines when they run clear. security was a massive issue.

to fix this, i built a compiler that parses the python animation locally, extracts all the elements and keyframes, and bakes them into a purely declarative design.json.

def pack(self, directory: str | Path) -> Path:
    # we literally run their code in an isolated builder
    # and dump the resulting state. no logic is saved.
    design_data = {
        "elements": builder._elements,
        "keyframes": builder._keyframes,
        "config": builder._config
    }
    
    with zipfile.ZipFile(out_path, 'w') as zf:
        zf.writestr("manifest.toml", toml_content)
        zf.writestr("design.json", json.dumps(design_data))

this means .clearfx packages are just dumb JSON files. the CLI’s DesignInterpreter reads the JSON and feeds it back into the engine safely. zero arbitrary code execution.

then i had to build the web studio so people could preview their animations in the browser. getting a live terminal running in react is a pain. i used xterm.js on the frontend, and built a websocket server in python using pty.fork() to spin up a fake terminal session.

@app.websocket("/ws/studio")
async def websocket_studio(websocket: WebSocket, width: int = 80, height: int = 24):
    pid, fd = pty.fork()
    if pid == 0:
        # child process - sets up the fake terminal size
        winsize = struct.pack("HHHH", height, width, 0, 0)
        fcntl.ioctl(sys.stdout.fileno(), termios.TIOCSWINSZ, winsize)
        os.execvpe("python", ["python", "-m", "clearfx.cli.main", "preview", temp_dir], env)

so when you click “execute” in the web UI, it literally sends the python code over websockets, saves it to a temp dir, forks a new PTY, runs the engine natively, and pipes the raw ANSI output back to xterm.js. it’s basically magic. but the resize bugs were driving me insane.

devlog 1: building a terminal engine from scratch because clear is boring

bro i got so tired of just typing clear and seeing the terminal instantly snap to empty. it’s so boring. so i decided to build a full double-buffered animation engine in python just to replace the clear command.

the idea was to make something that hooks into your .bashrc or .zshrc and intercepts clear, plays a sick terminal animation, and then actually clears the screen.

first thing i realized: you can’t just print strings to the terminal and expect it to look like a smooth 60fps animation. the flickering is insane if you do that. so i had to build a double-buffered FrameBuffer system from scratch.

class FrameBuffer:
    def diff(self) -> list[tuple[int, int, Cell]]:
        """returns only the cells that changed since last frame."""
        changes = []
        for i in range(len(self.front.cells)):
            if self.front.cells[i] != self.back.cells[i]:
                x = i % self.width
                y = i // self.width
                changes.append((x, y, self.front.cells[i]))
        return changes

this is the secret sauce. instead of redrawing 80x24 characters every frame, we only send ANSI escape codes for the exact pixels (cells) that changed.

then i built a DiffRenderer that takes those changes and constructs the smallest possible ANSI sequence to move the cursor and update the colors.

# snippet from renderer.py
for x, y, cell in changes:
    # only move cursor if we aren't already there
    if not (y == cy and x == cx):
        out.append(f"\033[{y+1};{x+1}H")
    
    # write the actual character with truecolor
    out.append(f"\033[38;2;{cell.fg[0]};{cell.fg[1]};{cell.fg[2]}m{cell.char}")

after getting the engine running, i needed a way to safely inject it into the shell without destroying people’s configs. built a ShellIntegration class that detects zsh/bash/fish, backs up the config, and safely appends a wrapper.

# >>> clearfx managed block >>>
if command -v clearfx &>/dev/null; then
  clear() {
    command clearfx play --clear-after
  }
fi
# <<< clearfx managed block <<<

i spent like 10 hours just getting the core engine to hit 60fps consistently. python is slow, so minimizing string concatenations and using sys.stdout.buffer.write instead of print() was mandatory. next up was actually building out the animation formats and getting a marketplace running so people could share these.