11 KiB

Raw Permalink Blame History

tenstorrent

Tenstorrent Blackhole compute backend. Runs ops on a p150a's Tensix cores through Foundry's own firmware, kernels and driver, rewritten from public documentation and measured hardware behaviour.

Headers: tt/ft_tt_device.h (device), tt/ft_tt_noc.h (data path), tt/ft_tt_launch.h (load/run), tt/ft_tt_queue.h (resident dispatch), tt/ft_tt_op_*.h (one per op family), tt/ft_tt_fttw.h (weights) Source: tt/*.c (host, C99) + tt/kernels/*.c (Tensix, C99, compiled ahead of time by SFPI) Status: 8 op families green on silicon (2026-08-12; the per-op figures below are one core of 140). Since then: multi-core matmul and Gated DeltaNet are built, and Pantheon runs a 9B qwen35 model end-to-end on the card (see tt/README.md)

Folder guide, layout and the build/test commands are in ../tt/README.md. This page is the deeper technical companion — what the layers are, why the shapes are what they are, and where the numbers came from.

Overview

cuda/ and metal/ accelerate an existing CPU op by handing a buffer to a vendor runtime. There is no vendor runtime here. The backend is the whole stack:

Layer	Files	What it owns
device	`ft_tt_device.c`, `ft_tt_tlb.c`, `ft_tt_arch_blackhole.c`	open `/dev/tenstorrent/N`, BAR mapping, TLB windows, arch tables
coordinates	`ft_tt_coord.c`, `ft_tt_arc_telemetry.c`	logical → NOC0 → translated, from ARC telemetry
data path	`ft_tt_noc.c`, `ft_tt_buffer.c`	L1/DRAM reads and writes through an aimed TLB window
load & run	`ft_tt_elf.c`, `ft_tt_kernel_blob.c`, `ft_tt_launch.c`	place kernel text, prime mailboxes, go, wait
dispatch	`ft_tt_queue.c`, `ft_tt_op_dispatch.c`, `ft_tt_group.c`	resident descriptor ring (fast) and one-shot launch (slow)
ops	`ft_tt_op_*.c`	per-family host drivers: CBs, runtime args, parameters
weights	`ft_tt_fttw.c`, `ft_tt_tilize.c`, `convert/`	GGUF → tiled on-card artifact, and uploading it
operations	`tools/ft_tt_core.c`, `ft_tt_watchdog.c`	per-core status, probe, recovery ladder

The dependency boundary: the SFPI compiler only (plus python3 for the tt/tools/emit_tensix_*.py generators), to rebuild the device images (they are embedded, so a Foundry build needs nothing). The backend was rewritten from public documentation and measured hardware behaviour; no tt-metal, TT-NN, tt-umd or tt-llk source is compiled, linked or vendored, and there is no C++ anywhere. The in-tree userspace driver needs no kernel module; tt-kmd's ioctl ABI is an optional path. See tt/README.md "Dependencies".

The execution model, and what it costs

Reader / compute / writer

A Tensix has five RISC-V cores. One op is five programs: reader on BRISC (NoC 0), compute compiled three times for TRISC0/1/2, writer on NCRISC (NoC 1). Splitting the data movement across two NoCs is what lets the reader run ahead of the math.

The compute unit is a 32x32 tile stored as four 16x16 faces. Everything downstream follows from that: the weight converter emits faces in face order, block-float formats share one exponent along a face row (so a transpose after packing is not free), and DST — the MATH↔PACK register file — holds 16 tiles.

Circular buffers are a compile-time contract

Data formats, tile geometry and math fidelity are baked into the kernel image via chlkc_descriptors.h, generated from a .cbtable. The host writes a matching CB config into L1 at launch. Nothing checks that these agree at runtime — a mismatch is a wrong answer, not an error. That is why every format decision is recorded in the .cbtable rather than being implied by a call site.

math_fidelity 4 (HiFi4) is a recorded decision, not a default: a fidelity sweep moved accuracy but moved wall time by 0.6%, because the op is data-movement bound at every setting.

Resident kernels, not per-op dispatch

The slow path (ft_tt_op_dispatch.h) is one host round trip per op — describe CBs, plant args, kick, wait. The fast path (ft_tt_queue.h) leaves a kernel resident on the core polling a 32-byte descriptor ring in its own L1; the host writes a descriptor and bumps a counter.

Three rules shape it, and each was paid for:

Completion is produced by the work item. The kernel writes completed_count and a per-run nonce after executing. RUN_MSG_DONE is 0, identical to a freshly booted mailbox, so "wait for DONE" would pass on a core that never ran.
Retire by asking. Every blocking wait polls stop_request in the same loop, and FT_TT_QUEUE_OP_STOP is an in-stream sentinel. assert_reset on a kernel blocked in a CB wait is what wedged cores.
Residency is bounded. idle_spin_limit makes the kernel exit on its own if the host dies.

Measured cost: 1.363 µs per op host-side at batch 32, 0.391 µs device-side op-to-op, against upstream's bh_p150 CI gate of 2.08 µs. The direct consequence, stated in the header: do not fuse ops to save dispatch — there is no dispatch to save.

Performance, and what each number is

Every figure is one core, HiFi4, at this card's 800 MHz.

Measurement	Value	Kind
matmul, `ffn_down`, prefetch depth 48	468.8 GFLOP/s = 28.6% of one Tensix	measured
matmul, `ffn_gate/up`, Mt=32	631.0 GFLOP/s = 38.5%	measured
matmul through the resident triad, real op	595.2 GFLOP/s = 36.30%	measured
one Tensix at HiFi4, 800 MHz	1.64 TFLOP/s	derived (2048 FLOP/cycle x 800 MHz)
AICLK under load	800 MHz, every sample	measured (473 busy samples)
dispatch, host-side, batch 32	1.363 µs	measured
dispatch, device-side op-to-op	0.391 µs	measured
attention, 8k KV	0.840 ms at `block_tiles=16`/descriptor → 40.3 ms/token → 24.8 tok/s	measured / composed
attention L1 footprint	413,696 B, flat in context	measured
top-k, 130,560 vocab	2.590 ms/token vs 55.6 ms to host = 21x	measured
binary switch	0.484 ms, ~93% of it the loader	measured
whole forward pass, one core	~111.9 ms = 8.94 decode tok/s	composed, not run
core health	140/140	measured

The distinction matters here and is enforced: 1.64 TFLOP/s is arithmetic from the vendor's matrix-engine spec, not a benchmark, so every "% of one Tensix" inherits that. The 111.9 ms forward pass is 3.745 ms/layer x 24 layers plus a 22.0 ms lm_head with the dispatch floor subtracted, composed from per-op costs; it predates the end-to-end forward Pantheon now runs.

Where the speed came from, and where it stopped

Phase 4 started at 81.3 GFLOP/s. The 5.9x to 468.8 was memory-identical and accuracy-identical, and it was not bandwidth: deeper reader prefetch was worth 5.5x on its own, and cutting transaction count another 1.28x. At 39.4 GB/s the card was never near a bandwidth wall. The residual ~103 ns/tile-pair floor beneath that is unidentified.

Two corollaries that generalise:

Byte reduction does not reliably convert to speed. Float16_b (2048 B/tile) → Bfp8_b (1088 B) bought 13%; Bfp8_b → Bfp4_b (576 B) bought nothing.
DST_ACCUM_MODE=1 bought 13-35x accuracy for zero wall clock (max-abs on ffn_down 2.187 → 0.063 at 473 GFLOP/s either way). A precision knob you did not set is still a precision knob.

Correctness gating

PCC alone is not a gate here — it caught none of four real kernel mutations on the reduce family, and a softmax mutation came back green because softmax(x-c) == softmax(x). Op checks report max-abs error alongside PCC, top-k is gated on exact index agreement, and every suite has a matching red table produced by mutating the kernel until it fails.

Weights: the FTTW artifact

convert/ streams a GGUF checkpoint into a tiled Blackhole artifact (magic FTTW). The unit of work is a 32-row strip, so peak RSS is set by strip buffers, not model size: a full model converts in 55.5 s at 210 MiB peak RSS on a 2-core x86-64 host.

The artifact is written in the card's own format, face-ordered and 64-byte aligned, so ft_tt_fttw_upload() is a byte copy — the loader does not transpose and cannot usefully. Two 2-D layouts exist as a per-tensor format field, not a hint:

FT_TT_LAYOUT_SOURCE — GGUF order, row-major tile grid. What an embedding table needs, because it is indexed by row.
FT_TT_LAYOUT_MATMUL_NMAJOR — transposed to K x N, tile grid stored column-major. Makes a reader's block of K-tiles one NoC transaction instead of depth strided ones.

The transpose has to happen in the converter because a Bfp8_b/Bfp4_b tile shares one exponent across a face row; transposing after packing means unpacking and repacking. 1-D tensors stay untiled fp32 — deliberate, on the hardware's broadcast rule.

Per-tensor precision comes from a ledger (convert/ft_tt_precision_ledger.h) and is gated by ft_tt_pcc_gate against the fp32 source: BFP8 worst 0.999967 over 219 tensors, BFP4 worst 0.991455.

Limits

Per-op figures are one core of 140. Multi-core matmul exists (split groups; see tt/README.md). In the single-core op there is no reuse: every A tile is re-read Nt times and every B tile Mt times, so scaling needs an in0 broadcast and core groups, not a fan-out of the current op.
Aggregate DRAM read bandwidth is >= 537 GB/s sustained across 14 concurrent cores and NOT saturated — the ceiling is still unknown, and every core in that test read the same ~17 MB working set, which is the friendliest possible pattern. Only achieved reader rates (28.4 → 39.4 GB/s) and the finding that the binding constraint was transaction count.
Card→host DMA does not exist on Blackhole. Host→card is 287 MB/s write-combined on this host's x1 link; reads back are 2.3-4.7 MB/s. Anything that ships tensors to the host for a decision is a design error — that is why top-k runs on the card.
MoE is not built (the SSM family's Gated DeltaNet is), and no fusion was built: measured operand placement showed the round trip a fusion would remove is 346 ns/tile of barrier, not DRAM.
Not wired into tensor/. There is no tm_device dispatch to Blackhole; tt/ is driven directly. Unlike metal/, it has no integration points inside tensor.c.

Operating rules

Never derive a NoC coordinate by probing. A wrong endpoint does not fail a read, it hangs the chip and costs a physical power cycle. Coordinates come from ARC telemetry.
Never park a resident triad, and never assert_reset a kernel blocked in a wait.
Wedged cores are recoverable without rebooting the host: ft_tt_core recover climbs a documented ladder, rungs 3+ resetting the whole card under an advisory lock.
Off-hardware test binaries are built without -DFT_TT_PHASE1B_NOC_ENABLED so they cannot reach the card even if they are wrong. On-card builds require PHASE1B=1 explicitly.

11 KiB Raw Permalink Blame History