System Integration & Automation Engineer

Erik Paperniuk

I work on backend systems, automation tools and applied machine learning. Recently that has included a lot of GPU kernels for running language models locally on Apple Silicon. Below are some of the projects, with the code and the numbers.

Portrait of Erik Paperniuk

Open source

Projects

The first two are forks: the engines are Inco's Splash and antirez's ds4, and my part is making them run well on M1 and M2. Everything here is public.

Fork of Inco's Splash 109Metal · C++ · Apple7/8 GPUs

Splash for M1 and M2

Upstream Splash needs an M3 or newer. Its kernels use bfloat16, which the M1 has no hardware for, and Apple's MPP matrix library, which reaches about a third of the GPU's peak there. I rewrote those parts to use formats the M1 runs natively.

  • 2.1× decode on npanj's five-prompt benchmark, 18.9 to 39.1 tok/s on the same Mac
  • 7.4 TFLOPS prefill matmul, up from 2.4 to 3.5 through MPP
  • 4-bit weights turned into exact half values with a bit trick, products and sums stay in fp32
  • Found a split-K reduction reading through a non-coherent cache that made the fastest decode path print garbage on M1
  • Verification kernel unpacks each weight once for all parallel requests: 18 to 65 tok/s with four requests
  • 35B-A3B at 144 tok/s, and the Neural Engine prefill from upstream 1.3 works on M1/M2 too
  • simdgroup_matrix
  • Metal
  • speculative decoding
  • GGUF / MLX 4-bit
Fork of antirez/ds4 16C · Metal · Apple Silicon, 64 GB and up

dstar: Qwen3.8-Flash-Next on a Mac

A launcher and a set of Metal kernels that run ISTA-DASLab's small quants of a 126B mixture-of-experts model (6.7B active) with the full 262K context, MTP and vision, behind an OpenAI and Anthropic compatible server.

  • 44 tok/s at 4K, 35.4 at 398K tokens of context, mostly the same code on every Apple GPU
  • Metal decoders for every quant type in the files, each checked bit-exact against the CPU
  • A 95 GiB n-gram table is never loaded: the 16 rows a token needs are read from the SSD
  • Indexer scorer rewritten for the M1: decode at 128K went from 23.8 to 33.8 tok/s, byte-identical output
  • Server keeps state at the last turn marker, so a retried agent turn on 31K tokens takes 0.3 s instead of 109 s
  • DeepSeek V4 Flash (81 GiB) streamed from the SSD on a 64 GB Mac at 12 tok/s
  • C
  • Metal
  • MTP
  • MoE
  • OpenCode / Claude Code compatible
Installer 20LM Studio · llama.cpp

LM Studio + Ternary Bonsai 2

Makes LM Studio load Ternary Bonsai 2 GGUF models (PQ2_0 and PTQ1_0) through Prism's llama.cpp fork, with a reasoning-effort selector. Linux and Windows, double-click install.

Patch + guide 6oMLX · mlx-vlm

oMLX + Prism Hadamard packs

Makes oMLX load Prism Hadamard Qwen3.5 packs (Ternary Bonsai 2), plus a guide to adding an unrecognised architecture to mlx-vlm.

Research

Deepfake detection

Other projects

Automation and applied ML

Benchmarks

Local inference numbers

Everything is measured on one M1 Max with 64 GB. Decode is limited by memory bandwidth and prefill by matrix throughput, so the work is getting each close to what the chip can do without changing what the model says.

Speed while the context fills up

One chat grown to its full context on an M1 Max with 64 GB, reasoning xhigh, each turn adds about 30K tokens of C source. MTP speculative decoding is on for the Qwen runs.

    Decode speed in tokens per second against context size
    ContextQ2_0IQ3_XXSDeepSeek V4 Flash
    4K44.035.212.1
    128K43.034.29.1
    256K38.831.7 (at 259K)not run
    398K35.4not runnot run

    44tok/s

    Qwen3.8-Flash-Next decode at 4K context, Q2_0, M1 Max

    398Ktokens

    context where it still decodes at 35.4 tok/s

    2.1×

    Splash on the same M1 Max with my kernels, 18.9 to 39.1 tok/s

    7.4TFLOPS

    my prefill kernel on a 10.4 TFLOPS GPU, Apple's MPP gets 2.4 to 3.5

    0.70J/token

    after 5 minutes under load at 73 °C, no throttling

    Same math, formats the M1 has

    The M1 has no bfloat16. A 4-bit weight becomes an exact half value by putting it in the mantissa of 1024.0 and subtracting 1024. Activations widen to fp32 by appending zero bits. Every product is exact, only the order of the additions changes.

    // q is 0..15
half w = as_type<half>(ushort(q | 0x6400))
       - 1024.0h;          // exactly q

    How close to the hardware

    Measured on the M1 Max, not taken from the spec sheet. Plenty is still left, a lot of it lost to around 1000 small dispatches per token.

    Decode memory traffic150 of 400 GB/s
    Flash-Next prefill matmul4 of 10 TFLOPS
    Splash prefill kernel7.4 of 10.4 TFLOPS

    Checking the output did not change

    Every kernel is tested against a reference path before it ships. Benchmarks use greedy sampling, a fresh server per build and no prompt cache, and the stock and new builds run in ABBA order.

    • Perplexity (NLL), before / after0.45044 / 0.45045
    • Top-1 token agreement99.92%
    • Mean KL divergence0.00013 nats
    • Quant decoders vs CPUbit-exact

    Writing

    Notes with the numbers

    Long write-ups of the projects above, with charts and the commands to reproduce them. Most started as posts on r/LocalLLaMA.

    All posts

    About

    Erik Paperniuk

    I'm a system integration and automation engineer. The public side of my work is mostly two things: GPU kernels for running models locally, and tools that automate something tedious, like contact extraction or Moodle setup.

    The inference work started because the engines I wanted to use skipped my own laptop. I measure everything on that one M1 Max and publish the raw numbers, so if you have an M2, M3 or M4 and want to compare, I'd like to hear from you.