ouroboros-kernelsmith-minicpm5-1b-GGUF

OUROBOROS Kernelsmith (MiniCPM5-1B) - GGUF

A 1-billion-parameter model that writes fused Triton GPU kernels. This is the quantized GGUF build, so you can run it on your own machine with llama.cpp and let it draft kernels offline.

It started as OpenBMB's MiniCPM5-1B and was fine-tuned on a loop that has no human labels in it: supervised training on kernels that a referee had already verified, then reinforcement learning where the only reward is that same referee's verdict (does the kernel compile, is it correct against PyTorch, is it faster). The model you are downloading is the small one. The source LoRA adapter is YMRohit/ouroboros-kernelsmith-minicpm5-1b, and the 27B sibling for heavier work is YMRohit/ouroboros-kernelsmith-qwen3.6-27b.

This GGUF is the local/offline path: it runs through llama.cpp and does not need a Modal call. The same project also has a Modal-backed path: the 27B run was trained and served on Modal H200s, and the public evidence is linked from the 27B model card, the corpus RESULTS.md, and the Space RUN.md backend endpoint section.

A 1B model writing GPU code is going to miss sometimes. That is the whole point of the project: you do not trust the model, you trust the referee. Draft a few kernels, throw away the ones that fail, keep the one that passes.

File

Install

pip install llama-cpp-python huggingface_hub

For GPU offload (much faster), install a CUDA build instead. Pick the index folder that matches your CUDA (the repo publishes cu121, cu122, cu123, cu124); for CUDA 12.4:

pip install llama-cpp-python \
  --extra-index-url https://abetlen.github.io/llama-cpp-python/whl/cu124

Download the GGUF:

huggingface-cli download YMRohit/ouroboros-kernelsmith-minicpm5-1b-GGUF \
  minicpm5-1b-kernelsmith-q8_0.gguf --local-dir .

The prompt format matters

MiniCPM5 is a reasoning model that uses ChatML. If you prompt it the normal chat way it will spend its budget thinking out loud and a 1B does not have the budget to spare. So we suppress the reasoning trace by closing an empty <think> block before the answer, and we ask for exactly one fenced code block. Get this wrong and the model rambles instead of writing a kernel.

The exact string we feed the model is:

<|im_start|>system
{system}<|im_end|>
<|im_start|>user
{user}<|im_end|>
<|im_start|>assistant
<think>

</think>

System prompt:

You are an expert GPU kernel engineer. Write a single correct, fast Triton kernel. Output ONLY one fenced python code block defining run(*inputs) and its @triton.jit kernel. Accumulate reductions in float32. No prose.

The user message describes the op (a short spec and the function signature) and hands the model one valid kernel for a different op as a style guide. This example matters more than anything else in the prompt. The model was trained with a row-wise reduction kernel (rmsnorm, one row per program) as the guide for almost every op, so that is the structure it learned to copy. If you hand it an elementwise kernel instead, it will happily write elementwise code for a reduction op like softmax and return the wrong answer. So: use the rmsnorm kernel below as your style guide for anything that reduces over the last dimension, which is most of what this model is good at.

Run it

This mirrors what the live Space does. Note create_completion on the raw rendered string, not create_chat_completion (which would turn the thinking trace back on).

from llama_cpp import Llama

SYSTEM = (
    "You are an expert GPU kernel engineer. Write a single correct, fast Triton kernel. "
    "Output ONLY one fenced python code block defining `run(*inputs)` and its @triton.jit "
    "kernel. Accumulate reductions in float32. No prose."
)

# A row-wise reduction kernel (rmsnorm). This is the style guide the model was trained with, and
# it is what teaches the model the one-row-per-program structure. Use it for any reduction op.
STYLE = '''@triton.jit
def _rmsnorm_kernel(x_ptr, w_ptr, y_ptr, stride, N, eps, BLOCK: tl.constexpr):
    row = tl.program_id(0)
    x_ptr += row * stride
    y_ptr += row * stride
    acc = tl.zeros([BLOCK], dtype=tl.float32)
    for off in range(0, N, BLOCK):
        cols = off + tl.arange(0, BLOCK)
        x = tl.load(x_ptr + cols, mask=cols < N, other=0.0).to(tl.float32)
        acc += x * x
    rms = tl.rsqrt(tl.sum(acc) / N + eps)
    for off in range(0, N, BLOCK):
        cols = off + tl.arange(0, BLOCK)
        mask = cols < N
        x = tl.load(x_ptr + cols, mask=mask, other=0.0).to(tl.float32)
        w = tl.load(w_ptr + cols, mask=mask, other=0.0).to(tl.float32)
        tl.store(y_ptr + cols, (x * rms * w), mask=mask)

def run(x, w):
    M, N = x.shape
    y = torch.empty_like(x)
    _rmsnorm_kernel[(M,)](x, w, y, x.stride(0), N, 1e-6, BLOCK=1024)
    return y
'''

# Describe the op you want. Keep it short and concrete. Note the style guide is a DIFFERENT op
# (rmsnorm) than the one you are asking for (softmax) - that is intentional.
USER = (
    "Op `softmax`: numerically stable softmax over the last dim (subtract the row max).\n"
    "Signature:\n  run(x: Tensor[M, N]) -> Tensor[M, N]\n\n"
    "Here is a valid Triton kernel for a DIFFERENT op (`rmsnorm`) as a style guide:\n"
    f"```python\n{STYLE}\n```\n"
)

def render(system, user):
    return (
        f"<|im_start|>system\n{system}<|im_end|>\n"
        f"<|im_start|>user\n{user}<|im_end|>\n"
        "<|im_start|>assistant\n<think>\n\n</think>\n\n"
    )

llm = Llama(
    model_path="minicpm5-1b-kernelsmith-q8_0.gguf",
    n_ctx=4096,
    n_gpu_layers=-1,   # -1 offloads everything to the GPU; use 0 for CPU-only
    verbose=False,
)

out = llm.create_completion(
    render(SYSTEM, USER),
    max_tokens=768,
    temperature=0.7,
    top_p=0.97,
    stop=["<|im_end|>", "<|im_start|>"],
)
print(out["choices"][0]["text"])

You get back a fenced Python block with a @triton.jit kernel and a run(...) entry point. If you want better odds, sample two or three at temperature=0.7 and keep the first one that passes.

Do not skip the referee

The model is the easy half. The kernel it writes is a guess until something checks it. Before you run a generated kernel anywhere real, put it through a verifier that:

1. compiles it,
2. checks it against the PyTorch reference on awkward inputs (allclose, not eyeballing), and
3. times it with CUDA events.

The exact harness we used is public and is the same referee that trained the model. It also blocks the obvious ways to cheat a benchmark (memoizing the output, mutating the input). Grab it here:

• Referee + live demo: build-small-hackathon/ouroboros-kernel-mint
• Verified kernels + how it was built: YMRohit/ouroboros-kernel-corpus

What it is good at, and what it is not

It is good at memory-bound fusion ops: normalizations, activations, gated MLP halves, softmax variants, and combinations of those (rmsnorm fused with an activation, add+layernorm, and so on). On those it reliably writes kernels that beat PyTorch eager, and the trained kernels hold a solid margin against torch.compile max-autotune on an H200.

It is not writing FlashAttention or matmul kernels, and it is not inventing new algorithms. These are scheduling wins on the kind of small fused ops that show up all over a transformer's forward pass. That is a narrow target on purpose, because a narrow target is one a 1B model can actually hit and a referee can actually check.

License and provenance

The fine-tuning, the quantization, and our code are MIT. The base weights are OpenBMB's MiniCPM5-1B and the base model's own license terms apply to them. No human-labeled kernel data was used at any point; every training signal came from the referee.

Built for the Hugging Face Build Small hackathon.