Spec-driven LLM operators across backends — built by agents

The spec is the source; kernels are derived from it and judged against it.

Quick Start · Why it's different · How it works · Installation · Benchmarks · Docs

Quick Start

import torch
import torch_musa  # registers the "musa" device
from tileops.gemm import GemmFwdOp

gemm = GemmFwdOp()  # shapes and dtype are inferred at call time

a = torch.randn(1024, 512, device="musa", dtype=torch.float16)
b = torch.randn(1024, 512, device="musa", dtype=torch.float16)

d = gemm(a, b)  # equals a @ b.T

Operators are auto-tuned on first use, MUSA Graph compatible, and declare their torch.compile(fullgraph=True) support per op. Upstream describes the same capture path as a CUDA graph; on this branch it is the MUSA graph.

Built for agents

An implementation can be regenerated from its spec; a spec cannot be recovered from an implementation. The project is organised around the spec rather than around the kernels:

• The spec is self-contained. Generation reads it and nothing else, so every constraint on the implementation is declared rather than assumed.
• Acceptance is decidable. Correctness settles against a declared reference, performance against a modelled bound — neither is a judgement call.
• The operator/kernel split is enforced. The boundary is checked rather than agreed, because an unenforced convention does not survive automated edits.
• Conformance is validated at every stage. Spec, generated code, tests and benchmarks each answer to a validator, so an operator is certified as it is produced rather than reviewed once at the end.

How it works

Each operator is declared in src/tileops/manifest/ before it is implemented. The entry drives code generation, testing, and benchmarking:

GemmFwdOp:
  ref_api: "torch.matmul"
  signature: {inputs: {a: {dtype: "float16 | bfloat16"}, b: {dtype: "same_as(a)"}}, ...}
  workloads: [{m: 1024, n: 1024, k: 1024, dtypes: [float16, bfloat16]}]
  roofline: {func: tileops.perf.formulas.gemm_fwd_roofline}
  source: {kernel: ..., op: ..., test: ..., bench: ..., kernel_map: ...}

Field	Role
ref_api	Reference implementation the tests compare outputs against.
signature	Tensor contract, shape rules, and dtype combinations; enforced at the op layer.
workloads	Shapes and dtypes the tests and benchmarks cover.
roofline	Performance model. Efficiency is achieved throughput over the modelled bound.
source	Paths to the kernel, op, test, and benchmark, and the slot-to-kernel map.

A validator checks every entry against its implementation in CI, so the declaration and the code stay in step.

The implementation is split in two layers. L2, the Python entry point, owns the caller-facing contract: validation, dtype casting, and memory layout. L1, the TileLang kernel, owns the GPU implementation. trust-model.md defines the boundary between them.

Installation

TileOPs installs from source; a PyPI release lands with the first stable version. This branch targets Moore Threads MUSA on MP31 (MTT S5000). Upstream's CUDA Toolkit 13.2 and SM90 prerequisites describe the NVIDIA build, not this one.

Prerequisites

• Python >= 3.10
• PyTorch 2.9.1 with the matching torch_musa
• MUSA Toolkit 5.2
• An MP31 GPU (MTT S5000)
• The MUSA build of TileLang shipped with the S5000 image (0.1.12)

git clone https://github.com/tile-ai/TileOPs
cd TileOPs
pip install -e '.[dev]' -c constraints.txt   # constraints.txt pins what CI validates
pre-commit install

python -m pytest -q tests -m smoke           # verify; requires an MP31 MUSA GPU

The image this branch is run in is sh-harbor.mthreads.com/ai-compiler/basesdk:ph1-5.2.0-torch2.9.1-triton3.6.0-tilelang0.12.0-v1.2.1. development.md still describes the upstream CUDA image; use the MUSA image above for this branch. That page also covers test tiers, benchmarks, and build troubleshooting.

Documentation

CONTRIBUTING.md	Naming, PR shape, what a review checks
development.md	Build, test, benchmark, dev image
architecture.md	Module map and the agent production loop
manifest.md	The spec format every operator starts from
ops-design.md	Adding an operator, step by step
roofline.md	How performance is scored against Speed-of-Light
trust-model.md	What each layer may assume about the others

The rendered site carries what this table cannot: the API reference generated from the operator signatures, and the performance tables from the nightly run, each operator against the tuned libraries it competes with.

Contributing

Operators are added through the loop above — start from ops-design.md, which walks the path from a manifest entry to a merged kernel.

License

TileOPs is released under the MIT License.