Skip to content
AI optimization copilot

AI that makes your models faster
— and can't lie about it.

Not a chatbot guessing at code. Kernalin's copilot studies where your GPU time actually goes, writes faster versions of your model, tests each one, and hands you only the version that's provably correct and genuinely faster. Everything that cheats or breaks, it throws away.

Request early access Explore the platform
Available
Not a chatbot

It writes kernels. Then it proves them.

A general chat model will happily hand you a plausible kernel that is subtly wrong. The copilot instead runs a closed loop: telemetry → hypotheses → variants → benchmark → verification. Every step is grounded in measurement, and nothing leaves the loop until it survives the gauntlet.
How verification works The profiler it reads
optimization copilot · softmax.py
variant 3 · triton kept 
@triton.jit
def softmax(x_ptr, y_ptr, n):
-  m = tl.max(x, 0)
-  e = tl.exp(x - m)
+  m = tl.max(x, 0, keep_dims=True)
+  e = tl.exp(x - m)   # fused, online
+  s = tl.sum(e, 0, keep_dims=True)
   tl.store(y_ptr, e / s)
Why this change

Online softmax fuses the max/exp/sum reduction into one pass, cutting a full DRAM round-trip. Grounded in the profiler's memory-bound diagnosis.

1 Workload
2 Profiler telemetry
3 Hypotheses
4 Kernel variants
5 Benchmark
6 Verification
7 Best verified impl 1.49×
The loop

From workload to best verified implementation.

Seven stages, run automatically. The copilot never ships a variant it hasn't measured and verified — the last box is the only one you act on.
01

Workload

The kernel or module you want faster, under its real shapes and dtypes.

02

Profiler telemetry

The copilot reads timing and profiles (and hardware counters when available) — it optimizes from evidence, not guesswork.

03

Hypotheses

It forms ranked, confidence-scored hypotheses grounded in the named bottleneck.

04

Kernel variants

Each hypothesis becomes a concrete rewrite — Triton today, CUDA C++ in preview — not a suggestion.

05

Benchmark

Variants race under distribution-aware measurement, clock-locked when the host permits.

06

Verification

Every survivor clears the correctness gate, anti-cheat, and the statistical verdict.

07 verified

Best verified impl

You get the fastest implementation that provably still computes the right thing.

kernalin · optimize loop best so far 2.84 ms
Propose
Verify
Measure
Decide
keep the winner · revert everything else · loop
iterationverdict
  • 01 tile 64×64 kept −12.3%
  • 02 vectorize loads revert
  • 03 fuse epilogue kept −8.4%
  • 04 enable fast-math revert
  • 05 stage in shared kept −16.2%
  • 06 unroll ×8 revert
3 kept · 3 reverted · verified winner 1.91 ms −32.7%

Watch it live: the reverts are the point — a variant dies the moment it fails correctness, reaches for numerics it wasn't granted, or lands inside the noise. Best-so-far only ever ratchets down.

What it recommends

Techniques grounded in a 6-tier playbook.

The copilot doesn't reach for tricks at random. It works a tiered playbook — from low-risk rewrites to algorithmic reformulations — and only escalates when the evidence supports it.

Kernel fusion

Collapse elementwise chains and reductions into one pass to kill DRAM round-trips.

Mixed precision

Run in fp16/bf16 where it is safe and accumulate where it is not.

Tensor-core targeting

Reshape math onto tensor cores when the shapes and dtypes cooperate.

Memory improvements

Shared-memory tiling and better reuse to lift a memory-bound kernel off DRAM.

Vectorized loads

Widen loads and stores so each transaction moves more useful bytes.

Coalescing

Reorder access so a warp touches contiguous memory in a single sector.

The playbook

1 Tier 1 · Low-risk rewrites — vectorized loads, coalescing, launch-config tuning.
2 Tier 2 · Fusion — elementwise and reduction chains collapsed into one pass.
3 Tier 3 · Precision — mixed-precision compute with safe accumulation.
4 Tier 4 · Tensor-core — math reshaped onto tensor cores where shapes allow.
5 Tier 5 · Memory hierarchy — shared-memory tiling and reuse restructuring.
6 Tier 6 · Algorithmic — online/streaming reformulations of the kernel itself.

Illustrative tiering. The copilot escalates only when profiler evidence justifies the risk.

Verification-gated

Every variant is verification-gated.

A faster kernel isn't useful if it's wrong. Whatever technique the copilot applies, the output runs the same gauntlet as any other candidate — correctness gate, input-differential and operation-coverage anti-cheat, statistical benchmark. A variant that fails is discarded, not shipped. On an H100 in September 2026, Claude's first rewrite of a Triton kernel passed and was 3.48× faster; three later rewrites were correct but not faster, so they were reverted.

Typical AI optimizer

Unproven
1 Generate
2 Benchmark
3 Ship "48% faster" could be wrong

A headline number with nothing behind it. If the kernel dropped an op, ignored an input, or memorized the reference output, you ship a bug at speed.

Kernalin

Verified
Generate
Correctness gate
Input-differential
Op-coverage
Statistical benchmark
Verified speedup deserves to exist

Every speedup survives correctness, anti-cheat, and statistical checks before it counts. The number you see is the number you can ship.

See the full verification gauntlet
Preview
In preview

PyTorch → Triton, automatically.

Source-to-source conversion: hand Kernalin a PyTorch op and get back a Triton implementation that is run through the same 5-stage correctness gate on your GPU. It works today for elementwise ops such as GELU and SiLU-multiply; wider coverage is in progress. The correctness contract doesn't change; only the starting point does.
Today

Verification-gated kernel rewrites in Triton; CUDA C++ in preview. PyTorch → Triton conversion for elementwise ops.

Next

PyTorch → Triton conversion for more of a module: reductions and fused chains.

Early access is open

Optimization you can act on.

Put a verification-gated copilot on your hottest kernels. Join the early-access program.

Request early access Explore the platform