Struct SpatialTarget 

pub struct SpatialTarget {

    pub name: &'static str,
    pub sites: MeshShape,
    pub site: SiteSpec,
    pub tile: TileShape,
    pub interconnect: BTreeSet<Interconnect>,
    pub handshakes: BTreeSet<HandshakeKind>,
    pub compute: Vec<ComputeCap>,
    pub dram: MemHierarchy,
    pub precisions: BTreeSet<Precision>,
    pub noc_rule: Option<NocRule>,
    pub write_ordering: WriteOrdering,
    pub cycle: CycleModel,
    pub modeled: &'static [&'static str],
    pub roofs: Roofs,
    pub part: Option<PartResources>,
}

The described mesh.

Fields

name: &'static str

sites: MeshShape

site: SiteSpec

tile: TileShape

interconnect: BTreeSet<Interconnect>

handshakes: BTreeSet<HandshakeKind>

compute: Vec<ComputeCap>

dram: MemHierarchy

precisions: BTreeSet<Precision>

noc_rule: Option<NocRule>

write_ordering: WriteOrdering

cycle: CycleModel

modeled: &'static [&'static str]

Fields whose values are MODELED (no public datasheet number) — pinned by the target lock so the descriptor never overstates what it knows.

roofs: Roofs

The rates the chip sustains, each cited: what a program’s work is held against.

part: Option<PartResources>

The programmable part this fabric is printed onto, when it is one: what the lowering spends as a resource budget. None for an abstract target (a GPU roofline, an RDU, the sky fabric) that no budget bounds.

Implementations

impl SpatialTarget

pub fn tile_limits(&self, acc: AccPolicy, exact: bool) -> TileLimits

The part’s limits in the units best_tile_config filters by, so any target can drive the lattice search without a bespoke preset of its own.

impl SpatialTarget

pub fn wormhole() -> SpatialTarget

Tenstorrent Wormhole (n150/n300): 8×8 worker grid, 1464 KiB L1 per Tensix, an 8-tile DST that halves under fp32 accumulate, two data-movement RISC-Vs, two NOCs, rectangle multicast. The matmul FPU is float-primary — Int8 is deliberately ABSENT from the tile precisions (int8 matmul_tiles measured 0/4096 correct under ttsim on both Wormhole and Blackhole); int8 GEMMs run as bf16 + SFPU requant.

pub fn blackhole() -> SpatialTarget

Tenstorrent Blackhole: 13×10 compute-with-storage grid, 1536 KiB of L1 per Tensix, eight DRAM channels of 3.984 GiB, and NOC writes that need an explicit flush before a dependent semaphore multicast. The accumulator is the same eight-tile half-sync window as Wormhole’s.

★ ..wh says “the same part, these fields different”, which is true of the NOC rule, the precisions and the compute capabilities — and was silently ALSO true of Wormhole’s L1 size, its SRAM total and its DRAM capacity, none of which are Blackhole’s. size_buffers shrinks a program to local_mem_bytes and tile.rs blocks against it, so inheriting 1464 KiB cost this chip 72 KiB of usable L1 on every program it ever compiled. The three lines below are read out of the vendor’s blackhole_140_arch.yaml and held there by spatial_chip_facts_lock.

pub fn sky_fabric(rows: usize, cols: usize) -> SpatialTarget

Our own Sky130/FPGA fabric: ONE site (the whole array), nearest-neighbour wires, valid/ready and FIFO handshakes, an exact-arithmetic tile engine (the proven weight-/output-stationary systolic array), and the embedded CGRA scalar datapath FabricTarget::full(rows, cols, 64).

pub fn xc7a100t_fabric(rows: usize, cols: usize) -> SpatialTarget

Our fabric on the Artix-7 XC7A100T of the Nexys A7-100T: sky_fabric’s machine with its local memory in the part’s block RAM — 135 blocks of 36 Kb (Xilinx DS180, Table 3) — read one cycle after its address, as a block RAM’s registered read port returns it, and an 8×8 PE array: 64 of the part’s 240 DSP48 slices (DS180, Table 3).

pub fn xc7a200t_fabric(rows: usize, cols: usize) -> SpatialTarget

Our fabric on the Artix-7 XC7A200T of the Nexys Video: 365 blocks of 36 Kb, read as xc7a100t_fabric reads its own, and a 16×16 PE array: 256 of the part’s 740 DSP48 slices (DS180, Table 3).

pub fn linked(self) -> SpatialTarget

This target on boards joined by a link: each board one chip of a tensor-parallel program, the link the hand-off a partial or an activation crosses from one to the other, a token with each word. A chip boundary needs that hand-off (HandshakeKind::Token); a hand-off on one board keeps its own.

pub fn sn40l_rdu() -> SpatialTarget

The SambaNova SN40L RDU, derived from RduChipFacts::sn40l: 1040 PCUs (the sites — the published compute_tiles), 520 MiB of on-chip SRAM spread over them, a crossbar switch fabric, token / valid-ready / barrier handshakes, and the embedded FabricTarget::rdu_shaped() scalar datapath. The mesh aspect, the per-PCU accumulator depth and the tile geometry are not public — listed in modeled, and cited on the facts’ unit model.

pub fn sn50_rdu() -> SpatialTarget

The SambaNova SN50 RDU, derived from RduChipFacts::sn50: 432 MB of on-chip SRAM, 64 GB HBM2E, DDR5 up to 512 GB, the same crossbar / handshake / compute envelope as the SN40L, and the embedded FabricTarget::sn50_shaped() scalar datapath. The PCU/PMU counts are not public, so the mesh, the per-site memory and the silicon tile figures join the modeled list beside the aspect, the accumulator and the tile.

pub fn rtx3060() -> SpatialTarget

The local RTX 3060 (sm_86): 28 SMs, each site a persistent thread block with 48 KiB static shared memory (99 KiB on opt-in, from the CUDA chip descriptor), 16×16 tensor-core fragments, one producer engine (cp.async warps), intra-block shared memory + global memory only (no clusters / DSMEM on sm_86), pipeline barriers. The accumulator fragment budget per warp is modeled.

pub fn provides_handshake(&self, h: HandshakeKind) -> bool

pub fn provides_interconnect(&self, i: Interconnect) -> bool

pub fn reloads_wide_spills(&self) -> bool

The target can reload a spilled partial of the accumulator’s own width.

pub fn accumulator_in_registers(&self) -> bool

The accumulator survives the whole reduction in the compute engine’s own registers, so a blocked plan never spills a partial — see ComputeCap::MatmulTile::acc_in_registers.

pub fn resident_sites_per_unit( &self, footprint_bytes: usize, osh: usize, osw: usize, ) -> usize

Independent accumulator LANES a site runs a block’s tiles across — CUDA warps on a GPU, one elsewhere (a Tensix site has a single DST register file). ★ HOW MANY SITES OF THIS FOOTPRINT A PHYSICAL UNIT HOLDS AT ONCE — the price of a deeper stage.

On Tensix a site IS a core, so the answer is one and a deeper contraction stage costs nothing but L1. On a GPU a site is a thread block and an SM holds several; the number is set by whichever of registers and shared memory runs out first, and THAT is what makes depth a trade rather than a win.

tile::best_candidate_for ranks register-accumulator plans by arithmetic intensity, and intensity_key cancels the contraction depth exactly — so depth fell to a tiebreak that maximized it, reasoning that “a deeper stage amortises the pipeline round-trip over more multiplies”. True, and silent about the price. MEASURED at 2048³ bf16 on an RTX 3060, one 128×128 output tile:

ibw 1 (K=16): 24 KiB, 2 blocks/SM → 5,429 GOPS ibw 2 (K=32): 40 KiB, 2 blocks/SM → 5,743 GOPS ← the deepest stage that keeps its occupancy ibw 4 (K=64): 72 KiB, 1 block/SM → 4,102 GOPS

The amortisation is real and the occupancy cliff is bigger. K=32 is also exactly the BK the hand-tuned knob lattice picks for this shape, which is the reassurance that the rule is the machine’s and not ours.

The constants are the CUDA lowering’s own (64K registers and 100 KiB of shared memory per SM); this is where they live so the planner and the emitter cannot disagree about occupancy — a duplicate of this rule would drift the first time either moved.

pub fn accumulator_lanes(&self) -> usize

pub fn accumulator_lane_grid( &self, osh: usize, osw: usize, ) -> Option<(usize, usize)>

★ HOW A BLOCK’S osh × osw OUTPUT TILES SPLIT ACROSS THE SITE’S ACCUMULATOR LANES.

A lane that owns a RECTANGLE of (osh/wm) × (osw/wn) tiles loads that many A-fragments and B-fragments per contraction step and does their product — the reuse that makes a GEMM a GEMM. The split is chosen to make the rectangle as SQUARE as possible: for a fixed area the squarest rectangle has the smallest perimeter, and the perimeter is what is loaded.

None when no split divides the block evenly. THE PLANNER MUST NOT OFFER SUCH A BLOCK on a register-accumulator target: its lanes cannot each own a rectangle, so the epilogue cannot write straight to global and the output has to stage through a ring the plan was not sized for. (It was: a 13×3 block sized against a 4-tile scratch wrote 39 tiles into it, which the device reported as an illegal access.)

pub fn packer_accumulates(&self) -> bool

The target’s packer can accumulate into a local-memory page.

pub fn packer_accumulates_into_the_output(&self) -> bool

The target’s packer can accumulate into the output’s own page, its intermediate sharing the output’s memory.

pub fn matmul_cap(&self) -> Option<&ComputeCap>

The tile-matmul capability, if the site has a matrix engine.

pub fn matmul_supports(&self, p: Precision) -> bool

Whether the matrix engine accepts p as an operand precision.

pub fn at_fidelity(self, fidelity: Fidelity) -> SpatialTarget

The same machine with its matrix engine declared at a different fidelity.

A Tensix core runs bf16 matmul at any of LoFi/HiFi2/HiFi3/HiFi4 — the fidelity is how many passes the engine makes, not a property of the part — but the descriptor pinned one value, so the pipeline could only ever emit HiFi4. That left us unable to emit at ttnn’s OWN bf16 default (HiFi2), which is the configuration the vendor’s tuned op actually runs, so the one knob group where a like-for-like head-to-head matters had nothing of ours in it at all.

The refusal this unblocks is worth keeping in mind: the emitter still refuses to PRINT a fidelity it did not run (cost.rs gates on matmul_fidelity()), and that guard is untouched. This changes what the machine is declared to do, not what a program may claim about itself.

pub fn matmul_fidelity(&self) -> Option<Fidelity>

The fidelity the matrix engine is declared at — the one the pipeline emits.

pub fn matmul_fp32_acc(&self) -> bool

Whether the matrix engine can accumulate in fp32.

pub fn matmul_acc_int_bits(&self) -> u32

Accumulator width of the matrix engine’s integer path.

pub fn storable_ints(&self) -> Vec<Precision>

The integer precisions a buffer is stored at, narrowest first: those the elementwise datapath lists, or — on a site that computes integers in one width and widens nothing (SiteSpec::int_compute) — that width alone.

pub fn matmul_max_array(&self) -> Option<(usize, usize)>

The largest PE array the matrix engine builds, (contraction rows, output columns), if it caps one.

pub fn matmul_block_api(&self) -> bool

Whether the matrix engine takes whole sub-blocks per call.

pub fn matmul_calls(&self, block: BlockShape, granularity: Granularity) -> u64

Calls the matrix engine makes for one compute op over block: by the engine’s BlockApi, or one per tile pair, osh · osw · ibw.

pub fn operand_loads(&self, block: BlockShape, granularity: Granularity) -> u64

Operand tiles the matrix engine loads for one compute op over block: by the engine’s BlockApi, or two for each tile-pair call.

pub fn matmul_dataflows(&self) -> Vec<DataflowPolicy>

The dataflow policies the matrix engine realizes, in descriptor order.

pub fn has_elementwise(&self) -> bool

pub fn has_row_reduce(&self) -> bool

pub fn elementwise_precisions(&self) -> BTreeSet<Precision>

The precisions the elementwise (tile) datapath computes in — what a Generic stage reads and writes.

pub fn index_precision_for(&self, max: u64) -> Result<Precision, Precision>

★ THE TYPE AN INDEX RIDES. An index is an ADDRESS — a token id, a cache position, a block start — so it is computed in an address-wide integer: the narrowest one of at least 32 bits the elementwise datapath has that HOLDS the index space [0, max]. On an RDU that is INT32, the widest integer the PCU has. Err names the precision the range would need — what the target is missing — so an index space past the chip’s integers is refused, never wrapped.

pub fn generic_holds(&self, p: Precision) -> bool

Whether a Generic stage may hold a value of precision p: a precision the elementwise datapath computes in, or the matrix engine’s fp32 accumulator width — the accumulation path every reduction scratch rides.

pub fn scalar_target(&self) -> Option<&FabricTarget>

The embedded scalar (CGRA) datapath descriptor, if the site has one.

pub fn effective_acc_tiles(&self, acc: AccPolicy, exact: bool) -> usize

Accumulator tiles available under an accumulation policy — the ONE place fp32 halving lives. The halving is paid only when the fp32 partial sums must survive packing (exact); a program that accepts bf16 spill rounding keeps the full accumulator file.

pub fn tile_bytes(&self, p: Precision) -> usize

Bytes of one native tile at precision p.

pub fn max_local_mem_bytes(&self) -> usize

The largest local-memory budget a site can be configured with.

pub fn with_sites(self, x: usize, y: usize) -> SpatialTarget

Every preset, for matrix locks. The same chip on a different deployment grid: the compute-with-storage mesh a device exposes depends on harvesting and on the dispatch mode (fast dispatch reserves a row of Tensix cores for the dispatch kernels).

pub fn wormhole_ttsim() -> SpatialTarget

Wormhole as the cycle-accurate simulator presents it: wormhole_b0_80_arch.yaml (10 Tensix rows, none harvested) under slow dispatch exposes an 8×9 compute-with-storage grid — 72 cores, one row reserved. A production n150 (one harvested row, fast dispatch) exposes 8×8. Every vendor example measured on the simulator spread its work over these 72 cores; a program placed onto 8×8 leaves a ninth row idle.

pub fn presets() -> Vec<SpatialTarget>

Trait Implementations

impl Clone for SpatialTarget

fn clone(&self) -> SpatialTarget

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for SpatialTarget

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl PartialEq for SpatialTarget

fn eq(&self, other: &SpatialTarget) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl Serialize for SpatialTarget

fn serialize<__S>(&self, __serializer: __S) -> Result<__S::Ok, __S::Error>

where __S: Serializer,

Serialize this value into the given Serde serializer.

impl StructuralPartialEq for SpatialTarget