SyclKittens

SyclKittens is a header-only tile programming library for writing fast Intel GPU kernels in SYCL. It follows the ThunderKittens 2.0 repository model (announcement): reusable primitives live under include/, kernels carry their own build files, and the API test suite lives under tests/.

Paper: SyclKittens: A Tile Programming Model for Programmers and Coding Agents on Intel GPUs (arXiv:2610.04277).

Requirements

• Intel oneAPI DPC++ compiler with icpx and SYCL 2020 support
• Intel Data Center GPU Max (PVC) for the supplied AOT target
• C++20
• XPU-enabled PyTorch and pybind11 for Python extension kernels (CONFIG=pytorch)
• oneMKL for the decode GEMV fallback
• Level Zero development headers for the direct-copy all-gather and distributed GEMM targets; set LEVEL_ZERO_INCLUDE=-I/path/to/include when they are outside the compiler path

To configure the build environment, run:

source env.src

Layout

SyclKittens/
├── include/                  # Header-only tile library
├── kernels/
│   ├── attention/            # Adaptive GQA forward and fused backward
│   ├── gemm/                 # BF16 GEMM and fused QKV projection
│   ├── norm/                 # RMSNorm and LayerNorm
│   ├── rotary/               # RoPE, fused Q/K, and fused RoPE-K/cache write
│   ├── activation/           # Fused SiLU multiply
│   ├── decode/               # Registered TP/CP decode with automatic GEMV fallback
│   ├── collective/           # Production *_kernel.hpp plus executable harnesses
│   └── context_attention/    # Production .dp.cpp plus *_bench.dp.cpp harnesses
├── tests/                    # Reusable API unit tests
├── benchmarks/               # Portable suites and result schemas
├── demos/                    # Minimal usage guidance
└── docs/                     # API documentation

The distributed targets keep production launch and dispatch code physically separate from allocation, initialization, timing, correctness checks, reporting, and main(). Collective production code lives in *_kernel.hpp; context attention production code lives in the non-_bench .dp.cpp files. Each local benchmark harness includes its production implementation into one translation unit, preserving the validated SYCL device-code compilation boundary.

Build

List the supported targets:

make help

Build the header/API test suite:

make tests

Build one operation family. Each family writes only to its local build/ directory.

make -C kernels/attention forward
make -C kernels/gemm bf16 M=4096 N=4096 K=4096
make -C kernels/decode decode
make -C kernels/collective all_reduce
make -C kernels/context_attention ring_overlap

Build modes

Each kernel in attention, gemm, norm, rotary, activation, and decode builds in one of two modes, selected with CONFIG:

CONFIG	Output	Needs
standalone (default)	executable; main() comes from <kernel>_harness.impl	SYCL only
pytorch	Python extension (sk_<kernel>*.so)	XPU PyTorch + pybind11

A standalone run times the kernel and checks it against a CPU fp32 reference; arguments are [iters] [warmup], and the exit code is non-zero on failure:

make -C kernels/attention forward
./kernels/attention/build/gqa_forward 100 10

To build and run every standalone harness at once (summary at the end, non-zero exit on any failure):

./kernels/run_standalone.sh [iters] [warmup]

Python extension targets use CONFIG=pytorch and PYTHON to select an XPU-enabled environment:

make -C kernels/norm all CONFIG=pytorch \
  PYTHON=/path/to/xpu/python3 RMSNORM_D=4096 LN_D=4096

API

Include the master header and keep kernels in the shared kittens namespace:

#include "kittens.dp.hpp"

using namespace kittens;

The collective construction API is documented in docs/collective_api.md.