tessel · a tile language and compiler
one kernel, three chips.
Designed, built and measured by Sankur KunduOpen source (Apache-2.0) · page updated 15 Oct 2026
Kernels are written against tiles, small blocks of values, and the compiler decides which thread holds which element. The same kernel compiles to CUDA for NVIDIA tensor cores, to Metal for Apple GPUs and to Pallas for TPUs, and is fast enough to run an LLM at vLLM's speed.
Layouts decide everything
Every tile gets a linear layout, as in Triton, and the code generator lowers it to the instructions each chip wants.
CUDA for tensor cores
ldmatrix, mma.sync, XOR-swizzled shared memory and software-pipelined loads, tuned on the GPU it runs on and compiled at run time with NVRTC.
Metal and Pallas
The same kernels compile to the Metal Shading Language and to Pallas, where range analysis must prove every block access in bounds or the kernel is refused.
Checked four ways
CI runs every kernel against a reference interpreter on an emulator of the CUDA execution model, compiles it with NVRTC, runs Pallas's TPU interpret mode and compiles with Apple's Metal compiler.
Token for token with vLLM
An LLM engine whose every GPU operation is a tessel kernel: fused QKV projection, rotary embeddings, a paged KV cache, grouped-query attention, SwiGLU, continuous batching, packed prefill and CUDA graphs.
TinyLlama-1.1B, 32 requests at once on a Tesla T4 (session 2; session 1: 1858 vs 1764). With 8 requests tessel reaches 0.94–0.99× vLLM. Generations identical to vLLM's.
At the speed of the vendor library
A Kaggle Tesla T4, 4 rounds of 100 launches each; every output checked against a float64 reference.
| Case | tessel | Fastest other | tessel's speed |
|---|---|---|---|
| GEMM 2048³ fp16 | 0.985 ms | cuBLAS 0.941 ms | 1.00× |
| GEMM 4096³ fp16 | 7.381 ms | cuBLAS 7.357 ms | 1.04× |
| Causal attention, S = 512 | 0.146 ms | SDPA 0.276 ms | 1.90× |
| Causal attention, S = 4096 | 4.416 ms | SDPA 6.362 ms | 1.44× |
tessel's speed is the median, over 4 rounds, of its speed relative to the fastest other engine timed in the same round (above 1 is faster). Times are medians of each round's medians, so the columns need not divide exactly: the T4 throttles as it heats. cuBLAS is torch.matmul with fp32 accumulation, as tessel.
A compiler, not a wrapper
The ideas behind Triton, rebuilt from the language up and taken to three kinds of hardware.
Linear layouts, from scratch
A typed tile IR where every tile has a linear layout, lowered to tensor-core instructions.
Vendor-level GEMM
Matches cuBLAS on fp16 GEMM after XOR swizzling and software pipelining.
Portable by construction
One kernel compiled to CUDA, Metal and Pallas, with bounds proved for TPUs.
An LLM engine on top
Serves TinyLlama at 1.05× vLLM's throughput, with identical generations.
tessel has not yet been timed on a TPU or an Apple GPU: Metal and Pallas output is compiled and checked in CI, but not benchmarked. With 8 concurrent requests the LLM engine reaches 0.94–0.99× vLLM.
@software{kundu2026tessel,
author = {Kundu, Sankur},
title = {tessel: a tile language and compiler for CUDA, Metal and Pallas},
year = {2026},
url = {https://github.com/Shakhtar-Sankur/tessel},
license = {Apache-2.0}
}Sankur Kundu
ML systems engineer & independent researcher · Bhubaneswar, India · sankur.kundu.tw@gmail.com
Builds the systems that train and serve language models. Also the author of lipi, a tokenizer that needs fewer tokens than GPT-4o's and Gemini's in 23 of 24 Indian languages; Babelguard, an open early-warning system for scams across languages; ratchet, RL post-training from scratch; kessler, an independent space-traffic audit that runs every day. Fixes merged into Apple's MLX, Meta's ExecuTorch, Hugging Face's OpenEnv and Oxford's PyZX, and pull requests in review at Google's JAX, NVIDIA, MIT, Carnegie Mellon and ETH Zurich.