Most inference frameworks we have toured on this blog, like vLLM for serving, llama.cpp for local chat, or PyTorch itself as the universal substrate, sit on top of somebody else’s hardware abstraction. Tenstorrent’s tt-metal is the opposite kind of project: it is the hardware abstraction. It is a single open-source repository that contains the low-level device driver, the dispatch firmware running on RISC-V cores inside the chip, and a full tensor-library front end named ttnn. If you have ever wondered what lives underneath a serving stack, this repo is as deep as it gets without taping out silicon yourself.

The repository organizes that ambition into two clearly separated stacks. The lower one, historically called tt-metalium, lives under tt_metal/ and speaks in the vocabulary of the machine: devices, command queues, buffers, programs, and raw RISC-V kernels. The upper one, ttnn, lives under ttnn/ and speaks in the vocabulary of machine learning: tensors, matmuls, layer norms, and sharded memory configs. Python only ever touches the upper stack, and every Python call bottoms out in C++ that builds a program for the device. This tour walks that path end to end, the same way we walked the compiler story for TVM and the kernel-authoring story for Triton and CUTLASS, but this time the “backend” is a physical card with Tensix cores on it.

We read the main branch, mapped both stacks into two diagrams (an overview and a detail view), and verified every line anchor against the working tree. The single best angle for understanding tt-metal is the seam: what exactly happens between ttnn.matmul(a, b) in Python and a multiply running on a RISC-V core that Python will never see.

Overview: a ttnn call crosses nanobind into C++ operations, which assemble a tt-metalium Program; fast dispatch serializes it into command pages that firmware kernels on the card consume. Files are linked at the anchor lines cited below.

Reading the overview from left to right:

• MatmulDeviceOperation re-export, L25: the Python ttnn.matmul module is thin; the real operation object is imported straight out of the compiled ttnn._ttnn extension.
• NB_MODULE(_ttnn, mod), L233: the nanobind extension root. It first defines submodules (tensor, device, operations), then binds classes, then binds functions, in that order.
• Tensor matmul(), L388: the C++ entry point, which normalizes arguments and forwards to the primitive ttnn::prim::matmul.
• MatmulDeviceOperation::select_program_factory, L2538: the op framework chooses among program factories based on shape, dtype, and the user’s program config.
• tt_metal::CreateKernel, L634: inside a factory, ttnn stops being a tensor library and becomes a tt-metalium client, attaching reader, mcast, writer, and compute kernels to a Program.
• class Program, L24: tt-metalium’s unit of work: kernels, circular buffers, and semaphores destined for one device launch.
• assemble_device_commands, L2417: the dispatcher-to-be serializes the program into CQ commands, runtime-argument writes, and launch messages.
• class HWCommandQueue, L30: the host-side queue that owns those command pages in system memory.
• kernel_main, L1589: the on-card dispatcher firmware; it consumes command pages and fires launch messages at worker cores over the NOC.

Two Stacks, One Accelerator

The tt_metal/ tree splits into a public API and an implementation. The API surface is unusually clean for a hardware stack: class IDevice, L64 declares what a device can do (worker cores, sub-devices, programmable core types), while CreateKernel, L124 is the free function you call to attach a kernel to a program. The implementation behind it is where the years live: Device::initialize, L553 wires up the process-wide MetalContext, L150, the allocator, and the dispatch firmware.

One detail worth pausing on is what a “kernel” is here. Tensix cores contain several small RISC-V processors per tile, and tt-metal lets you program them individually. enum DataMovementProcessor, L23 names them: RISCV_0 is BRISC, RISCV_1 is NCRISC, and a compute kernel runs on the TRISC set. So one ttnn matmul typically becomes several cooperating binaries per core, talking through circular buffers. This is a very different mental model from one big CUDA kernel, and it explains why the repo carries its own kernel sources next to every C++ operation.

The ttnn/ tree mirrors the same split. Public headers like class Tensor, L30 define the value types, ttnn/cpp/ttnn/operations/ hosts hundreds of operations, and ttnn/ttnn/ is the Python package. The generic machinery in launch, L606 is the interesting part: every device operation funnels through one templated launch path that validates inputs, reuses cached programs, and eventually calls EnqueueMeshWorkload, L227 on the device’s mesh command queue. A single card is just a one-chip mesh, which is why mesh abstractions like class MeshDevice, L72 are not optional extras; they are the normal case.

The Python-to-C++ Seam

The binding layer uses nanobind rather than pybind11, and it is disciplined about ordering. NB_MODULE(_ttnn, mod), L233 comments explain why: every class and enum must be bound before any function that uses them as arguments. That is why the module first registers types such as the BufferType enum, L79 (DRAM, L1, L1_SMALL, TRACE), then binds operation functions.

Operations take one of two routes into Python. Newer, “C++-first” ops like matmul are bound wholesale on the C++ side, as in bind_function “matmul”, L693, docstrings and program-config types included. The Python layer then does almost nothing except re-export and attach a golden function: a pure-torch implementation used for correctness comparisons, registered through REGISTERED_OPERATIONS, L1751. That registry also powers ttnn.query_registered_operations(), a neat introspection surface few frameworks offer.

Device management crosses the same seam. open_device, L44 in Python is literally an alias for the bound ttnn._ttnn.device.open_device, documented in the device bindings, L73. And when you convert data with ttnn.from_torch(..., device=device), the wrapper at def from_torch, L282 creates the tensor, applies layout and dtype, and may immediately upload it through a command queue. Compared with the hand-written binding files we saw in MLX, this is a much more industrial setup: the seam is wide, typed, and mostly machine-repetition rather than glue hacks.

How an Op Becomes Silicon Work

Follow matmul deeper and the layered discipline pays off. The C++ wrapper at matmul.cpp forwards to a primitive; the primitive is a generic DeviceOperation; the operation picks a factory; and the factory does the real work of hardware mapping. The mcast-1D factory, for example, decides which cores multicast A’s tiles, which cores hold B’s shards, and which compute kernels consume both, calling tt_metal::CreateKernel, L634 once per kernel role. Each of those corresponds to an actual .cpp kernel source under device/kernels/, like the compute kernel at kernel_main, L20 and its data-movement counterpart at kernel_main, L14.

Once the program exists, dispatch takes over. FDMeshCommandQueue::enqueue_mesh_workload, L407 is the fast-dispatch front door, and assemble_device_commands, L2417 turns the program into a compact command sequence: launch messages, runtime-argument writes, semaphore updates. Those commands go into a host-side ring buffer managed by SystemMemoryManager, L123, which the card reads over PCIe. Two firmware kernels on reserved cores do the rest: a prefetcher at kernel_main, L3284 streams command pages downstream, and the dispatcher at kernel_main, L1589 executes them, writing kernel binaries and arguments to worker cores over the NOC. Even addresses come from a hardware abstraction: class Hal, L391 is populated per architecture by files like Hal::initialize_wh, L243 for Wormhole.

The elegance is that dispatch is itself a program on the same chip. There is no privileged host daemon copying work into a magic doorbell; the “driver” is firmware you can read, in this same repository.

Memory: Banks, Buffers, and Tiles

Memory is where tt-metal’s opinions get sharp. Every allocation is a class Buffer, L154, created through BufferImpl::create, L542, which asks the allocator for addresses. The allocator splits DRAM and L1 into banks managed by BankManager, L74, so an interleaved tensor is really a round-robin of pages across banks to spread bandwidth. Sharded layouts skip the round-robin and pin each core’s data to its own L1 region, which is what makes big matmuls and attention fast here.

The other half of the story is tilization. The matrix engine does not consume rows and columns; it consumes 32x32 tiles. That means every tensor has a logical shape and a padded physical shape, and layout conversion is a first-class operation. On the host, to_tile_major_layout, L206 does the conversion for CPU-resident data. On the device, a dedicated kernel at kernel_main, L26 performs the same trick as part of an op graph. Kernels exchange these tiles through circular buffers, tt-metal’s producer-consumer primitive, which is why program factories spend so much code choosing circular-buffer sizes. If you have read about warp-level tiling in CUTLASS, this is the same idea with different vocabulary, except here the tile size is fixed by the hardware and the compiler never hides it from you.

From Install to Your First Model

The quickest path is a wheel: the ttnn README documents pip install ttnn for machines with Tenstorrent hardware and drivers already set up. Building from source is the honest path for reading along: clone the repository, run ./build_metal.sh (the root CMake build), and the Python extension is produced from the same tree you just read. There is also a dockerfile directory if you want a hermetic environment.

The APIs used below are all taken from the repo’s own tutorial script, so this snippet is exactly what the library expects to work:

import torch
import ttnn

device = ttnn.open_device(device_id=0)

a = ttnn.from_torch(
    torch.rand(32, 32), dtype=ttnn.bfloat16,
    layout=ttnn.TILE_LAYOUT, device=device,
)
b = ttnn.from_torch(
    torch.rand(32, 32), dtype=ttnn.bfloat16,
    layout=ttnn.TILE_LAYOUT, device=device,
)

out = ttnn.matmul(a, b, memory_config=ttnn.DRAM_MEMORY_CONFIG)
print(ttnn.to_torch(out))

ttnn.close_device(device)

Three details connect back to the source tour. First, layout=ttnn.TILE_LAYOUT is you opting into tilization at upload time, saving an on-device conversion. Second, memory_config=ttnn.DRAM_MEMORY_CONFIG is the interleaved layout managed by the BankManager allocator. Third, ttnn.matmul here runs the whole chain from the overview diagram: nanobind to C++ op to program factory to dispatch firmware. The same pattern scales up: the repo’s models/ tree contains ported networks, and if you want to see where such a runtime sits in a bigger system, our tours of MLC-LLM and ONNX Runtime cover the layers above.

Honest Limits

No hardware, no execution. Everything dispatch-related in this post was read, not run: tt-metal targets Tenstorrent cards (Wormhole, Blackhole, and newer), and the CI constantly exercises real silicon. Without a card you can still build and inspect, but the first open_device will not get far. That is inherent to the project, not a flaw, but it makes the curve steeper than for stacks where rented GPUs are everywhere.

The tree moves fast enough that line numbers are a perishable commodity. The dispatch internals have clearly been refactored recently; command-queue code now lives under several directories, and the Metal 2 host-API work is reshaping program construction. Any anchor cited here reflects the main branch at the time of writing and will drift.

Kernel coverage in ttnn is broad but uneven. Many ops accept program configs that dramatically change which factory runs, and the fallback-to-torch golden path exists precisely because some composite behavior is easier to verify than to reimplement. Reading ttnn.matmul alone involves thousands of lines across six factories; this tour follows one representative path rather than all of them.

Finally, the firmware kernels are long, dense C++ with compile-time constants from the HAL, and the prefetcher file alone runs past three thousand lines. The abstractions are honest, but the bottom of the stack is still the bottom: sometimes you read constexpr address arithmetic written for a chip, not for a reader. That is the price of a driver and a tensor library sharing one repository, and it is what makes tt-metal one of the few places where you can see the whole road from import ttnn to a RISC-V core doing arithmetic.

Enjoyed this post? Never miss out on future posts by following us