megakernel

LLM speculative inference server for heterogeneous hardware & consumer GPUs

An agent harness that compiles a model into one provably-correct, self-retargeting CUDA megakernel and self-tunes it past cuBLAS at batch-1 LLM decode, paper: https://arxiv.org/abs/2606.09682

Single-launch batch-1 decode engine for PrismML Bonsai 27B (ternary and 1-bit) on NVIDIA GPUs. 1.76x the vendor llama.cpp fork on H100, same outputs.

A light, transparent, and modular inference & quantization engine for studying LLMs.

TileMega is a compiler framework for generating persistent megakernels for LLM inference, using ISL-based symbolic task-dependence analysis and CuTe/CUTLASS-backed code generation.

Persistent-kernel CPU→GPU command bus. 64-byte packets, 0.5 µs enqueue, ~96 ns dispatch — no cudaLaunchKernel. C++/CUDA, 2 files.

Vendor-portable GPU decoders for quantum LDPC codes — Triton min-sum BP & Relay-BP on NVIDIA (CUDA), AMD (ROCm), and Apple-silicon (Metal), consuming any stim DetectorErrorModel.

Ampere megakernel for RTX 3090 (sm_86 / GA102). Fork of cohere-ai/cohere-megakernel: WGMMA+TMA → mma.sync.m16n8k16 + cp.async + ldmatrix. Beat vLLM on 3090, not match H100.