Skip to main content

Module tile

Module tile 

Source
Expand description

tile — choose the per-site block plan from the legal lattice by the cost model.

Candidates: the per-tile STREAMING family (one output tile per work item, single-tile FIFOs, tt-metal’s work split over the whole mesh) and every BLOCKED point of the divisor lattice (per_site_m | Mt, per_site_n | Nt with the block grid inside the mesh, osh | per_site_m, osw | per_site_n, osh·osw ≤ the accumulator budget, ibw | Kt), each legal only if its buffers fit local memory. Every candidate is scored for both reader arrangements (fused/multicast where the target can, unfused) with the target’s cycle model; the minimum wins. Deterministic tie-break: fewer bytes, then fewer sites, then the lexicographically smallest plan. The winning arrangement is NOT stored — fuse_readers re-derives it from the same cost function.

Structs§

Candidate
One scored point of the lattice.
TileConfig
A tiling / blocking config for the blocked matmul kernel — the knobs an autotuner searches, in 32×32 tiles. best_tile_config fills this from a cost model; the vendor hand-tunes the equivalent.
TileLimits
What a part can hold, in the units the lattice is filtered by. A named-part preset produces one of these; so does a SpatialTarget, which is what makes the search portable rather than Tenstorrent-shaped.

Constants§

TWIN_PLANNING_CYCLES
The cycles of simulation the planner spends timing one program’s candidates on the fabric (fastest_on_the_twin): enough to time every candidate of a program the size of a decode step’s section, a fraction of a second of the twin.

Functions§

admissible
The candidates objective may choose among.
apply_plan
Apply a plan to a template: patterns, block shape, spill buffer, work items.
best_among
The candidate objective chooses among cands: the admissible ones (admissible), ranked by plan_key.
best_candidate
The winning candidate (cost-minimal, deterministic tie-break).
best_candidate_for
The winning candidate under a stated objective.
best_tile_config
Smart-lattice autotune: search the divisor lattice of the tiling space — per_core_m | mt, per_core_n | nt, out_subblock_h | per_core_m, out_subblock_w | per_core_n, in0_block_w | kt — keeping only points LEGAL on this part (grid fits, the accumulator holds the sub-block, blocks fit local memory), and return the minimum-tiling_cost point. This is the compiler choosing the config the vendor hand-tunes, and it adapts to the part’s limits: a larger accumulator or grid unlocks larger, faster lattice points.
candidates
Every legal candidate of p on target, scored.
candidates_where
Every legal candidate of p on target, scored, where the plans a caller will keep decide whether the grid may outgrow the mesh: items wrap onto sites only when no blocked plan the caller keeps fits the mesh. A caller that restricts the arrangement (a calibration run pinning the packer’s accumulate) gets the best plan of that arrangement, not the streaming fallback of a round the plans it discards had filled.
fastest_on_the_twin
★ ON A TARGET WHOSE TIME IS EXACT, THE PLANNER TIMES ITS CANDIDATES. The fabric’s cycle twin (fabric_cycles_within) is cycle-exact against its RTL, so the planner builds candidates in the work model’s order and times each one, and the fastest is the plan. A model stands in for time only where nothing can measure it; on the fabric every engine change would otherwise leave the model choosing a plan the twin can show is slower.
finishes_in_place
★ WHERE THE PARTIAL ALREADY IS THE RESULT, THE REDUCTION FINISHES IN THE OUTPUT’S PAGE. Exact numerics hold the partial sums at the accumulator’s own width; where the output is stored at that width and nothing computes after the matmul, the last K block’s sum is the result itself. A target whose packer adds into the output’s page (PackerAccumulate::IntoTheOutput) then accumulates every block there — the intermediate shares the output’s memory (Buffer::alias_of) — and never reloads the partial into the register file, which Wormhole and Blackhole cannot do at that width.
l1_block_tiles
Rough local-memory footprint (in tiles) of the resident double-buffered operand blocks for cfg — this is what bounds in0_block_w.
tile
The tile pass.
tiling_cost
Estimated per-site cost of a tiling (lower = faster) — a heuristic RANKING model, not an absolute cycle count. It captures the three levers that move blocked-matmul time: (1) parallelism — the per-site output-tile work per_core_m·per_core_n, minimized by using more sites (dominant term); (2) reuse — the number of accumulator acquires per site (per_core_m/osh)·(per_core_n/osw)·(kt/in0_block_w), minimized by a bigger sub-block and wider inner block; (3) reload — the spill/reload traffic when the contraction is blocked.