Triage Tt Metal Asserts

Triage tt-metal assert changes against tt-mlir coverage, per optimization level

A tt-metal uplift changes the TT_FATAL checks ops run in validate(). Each one constrains a tensor or attribute that tt-mlir chose at compile time, and tt-mlir chooses differently at each optimization_level. Sort every assert into exactly one bucket, and for LOOSE and TIGHT say which level it applies to:

A deleted or relaxed assert whose only tt-mlir consequence was a constraint is TIGHT with the tag stale. A TIGHT finding never gets applied by this skill: relaxing a constraint changes codegen for every model using the op and wants its own change and test run.

Definition of Done

Every in-scope assert lands in one bucket with a citation per optimization level, for the mechanism that covers it or for the mechanism that should have. No assert left as "probably fine". Truncated conditions resolved, not guessed. Indirect and level-conditional coverage labelled as such, never as plain COVERED. Every in-scope op also gets a TIGHT pass (Stage 6) even if none of its asserts changed in a way that hurts.

What each optimization level can do

Resolve every claim below against origin/main; the level toggles live in resolveOptimizationLevelOptions() in include/ttmlir/Dialect/TTNN/Pipelines/TTNNPipelines.h and the pass order in createTTNNPipelineAnalysisPasses in lib/Dialect/TTNN/Pipelines/TTNNPipelines.cpp.

This table is the default effective mapping for a stock optimization_level. resolveOptimizationLevelOptions() derives the toggles only when they are unset, so explicit options (enable-optimizer, memory-layout-analysis-enabled, a forced compute-config or dtype) move behavior between columns: level 0 with enable-optimizer=true runs the optimizer, level 1 with sharding enabled gets L1 candidates. Always read the actual invocation before trusting a row.

Consequences that decide buckets:

• Level 0 has no layout/constraint net, but is not pass-free. A violated layout assert is a device fault: only operand workarounds, decomposition workarounds, and TTIR verifiers stand between the IR and the kernel, and buffer-type/memory-layout asserts are mostly moot because everything is DRAM interleaved unless a workaround sets otherwise. But the always-on pre-optimizer passes still run here (weight/KV/activation dtype conversion, TTNNSetComputeKernelConfig), so dtype and compute-kernel-config asserts are shaped by them at level 0 too.
• Levels 1 and 2 catch what the constraint query can see. The query runs metal's validate() and program creation for the exact layouts in the IR. A new TT_FATAL is caught at compile time if, and only if, (a) the op has a working OpModel, (b) the query passes the same arguments the runtime passes, and (c) metal checks the condition in validate() or program creation rather than deep in a kernel. Being caught is not the same as being fixed: the fallback set must contain a config the kernel accepts, and fallback success still costs L1 residency and possibly a dtype cast.
• The optimizer never changes dtype; only the fallback pass does, and only into DRAM interleaved. A dtype assert at level >= 1 is satisfied by a whitelisted operand workaround, by a pre-optimizer dtype pass (TTNNWeightDtypeConversion, TTNNKVCacheDtypeConversion, activation dtype lowering), or by OperationValidationAndFallback casting a failing op's inputs to one of its four dtypes with DRAM-interleaved layouts. So a dtype requirement can be met at levels 1 and 2, but never on an L1 tensor: an op that needs both L1 and a different dtype has no compile-time path except a whitelisted workaround. The op's original config is validated first, so a bfp8 operand set by a pre-optimizer dtype pass is tried as bfp8; only when that fails does the fallback cast, and its set contains no bfp8, so a failing bfp8 tensor gets widened to one of the four.
• Silent wrong results are never caught by queries. If metal returns garbage instead of asserting (its comment or the linked issue will say "incorrect results"), a rule-book filter or a workaround owns it, not the query (e.g. the where operand workaround forces the predicate dtype for exactly this reason). The uplift diff cannot tell you when it is safe to drop such a constraint, so those TIGHT candidates need a silicon check.
• Compute-kernel-config asserts (math fidelity, fp32 dest acc) are reachable only where a knob is actually applied. By default that is level 0 (hifi4, fp32_dest_acc_en=true on every op without its own value, TTNNSetComputeKernelConfig), or any level when the frontend forces a knob (tt-crank and tt-xla both forward math_fidelity / fp32_dest_acc_en into computeCfgMathFidelity / computeCfgFp32DestAccEn). At levels >= 1 with nothing forced, TTNN's own defaults are what the kernel sees, so an assert there is on metal's defaults, not on anything tt-mlir sent.
• Whitelisted operand workarounds insert ttnn.to_tensor_spec before the optimizer. An L1 result of that op is unplaceable by both L1 spill trackers (tt-mlir #9299). For a layout-only requirement at level >= 1, prefer a rule-book filter over a whitelist entry; reserve whitelisting for dtype requirements the optimizer cannot express.

Stage 0: baseline

Hard rule: resolve everything against origin/main, never the checked-out branch.

git fetch origin main -q
git rev-list --count HEAD..origin/main   # nonzero => working tree is stale

A feature branch tens of commits behind will show ops as "not called" that main calls, producing a confidently wrong "nothing to do here". Use git grep <pat> origin/main -- <paths> and git show origin/main:<path> for every lookup below. Audit a branch instead only if the user asks, and say so.

Stage 1: parse the diff into claims

One row per assert: (metal file, op, condition text, added | modified | deleted).

• For modified, recover both old and new condition. The delta is the whole point.
• If a digest truncates a condition (…, &&, a cut line), fetch the full text before classifying. Guessing the tail is how a gap gets certified as covered.
• Keep the surrounding comment. Metal's comments state the reason (CB page sizing, scaler computed from padded shape). That decides whether a violation is a loud crash or silent numerical corruption, which drives priority.

Stage 2: scope filter by symbol, never by path

Do not filter by directory. Ops under tt-train/sources/ttml/ look like a separate consumer of tt-metal, but tt-mlir calls several directly (ttml::metal::adamw, ttml::metal::sdpa_fw, ttml::metal::sdpa_bw). Path scoping silently drops real exposure.

git grep -l "<op_name>" origin/main -- runtime lib include

Three call paths, all counting:

Grep trap: ttml is a substring of ttmlir, so grep ttml matches nearly the repo.

git grep -nE "\bttml::|ttml/" origin/main -- runtime lib | grep -v ttmlir

List zero-hit ops in the report as out-of-scope rather than omitting them, so the reader can audit the filter. Scope is usually the largest single reduction; apply it first.

Stage 3: drop the IGNORE tier

Everything in a validate() is nominally about caller data, so "internal to tt-metal" is almost never the right dismissal. The real noise tier is asserts we cannot violate:

• Device identity: x.device() == query.device(). One program targets one device/mesh and every tensor comes from that program's pool.
• Allocation / liveness: buffer() != nullptr, storage_type() == DEVICE. The runtime tensor pool guarantees this (DEBUG_ASSERT(ttnnTensor.is_allocated()) on insert). Exception: TTNNLayout deliberately keeps conv2d weight arguments and MeshShard inputs in system memory, so a storage_type() == DEVICE assert on those paths is live, not IGNORE.
• Derived-quantity restatements: e.g. target_pages == input_nc_pages sitting beside a shape assert that already pins the relation; it guards metal's own page indexing.

Report these as one grouped line. State the reason they are noise, because it is a property of tt-mlir, not of the assert: multi-device programs would make device-identity checks live again. Never silently delete them from the digest.

Stage 4: classify the rest by constraint kind

The kind names the mechanism that can own it. Find the mechanism first, then in Stage 5 check whether it is in force at each level.

Ops with no entry in getRuleBook() use defaultRules: accept every input layout, null output hint plus sharded fallbacks, reshards explored. That is the loosest possible setting and the first thing to check for a LOOSE finding at level 2.

Filter helpers (layout_filter_utils::, include/ttmlir/Dialect/TTNN/Analysis/OpRules/LayoutFilterUtils.h): requireTiled, requireRowMajor, rejectAllSharded, rejectL1Interleaved, rejectWidthSharded, requireDRAMInterleaved, allowOnlyShardingType, isFullBboxSharded; output-hint helpers nullHintOnly, nonShardedOutputHints, dramInterleavedOnlyOutputHints.

Stage 5: per-level matrix

For every in-scope assert fill three cells, one per level, each with exactly one of:

• forced(<mechanism file:line>): a workaround, rule book, verifier or pass makes the violating config unreachable at that level.
• caught(<fallback outcome>): the constraint query would reject the violating config at compile time. State whether the fallback set contains an accepted config (caught, fallback ok) or not (caught, compile fails).
• unprotected: the violating config is reachable and nothing checks it.

Then derive the bucket: all cells forced or caught, fallback ok gives COVERED, with the per-level caveats in the table. Any unprotected or caught, compile fails gives LOOSE at that level. Fallback ok at level 1 and 2 with a dtype cast or a lost L1 residency is COVERED for correctness but is also a TIGHT/perf note in Stage 6.

Four ways a forced or caught cell is wrong:

5a. Operand workarounds are gated by an opt-level whitelist

At optimization_level >= 1, operand workarounds run only for ops in enabledOpsForWorkaroundWithOptimizer (lib/Dialect/TTNN/Transforms/Workarounds/TTNNWorkaroundsPatterns.cpp, list at the file end). At level 0, all TTNN dialect ops get theirs.

git grep -n "create<Op>OperandsWorkarounds" origin/main -- lib          # 1. exists?
git show origin/main:lib/Dialect/TTNN/Transforms/Workarounds/TTNNWorkaroundsPatterns.cpp \
  | grep -n "ttnn::<Op>::getOperationName"                             # 2. whitelisted?

Exists but not whitelisted: forced at level 0, unprotected or caught at levels 1 and 2 depending on the OpModel. The whitelist comments are the model for justifying an entry. Absence is sometimes deliberate (ArgMax is out because its rule book supplies RowMajor input siblings instead); check for a comment before proposing an addition, and prefer a rule-book filter for layout-only requirements (see the to_tensor_spec note above).

5b. The constraint query may not run for this op

git show origin/main:lib/Dialect/TTNN/Interfaces/TTNNOpModelInterface.cpp | awk '
  /[A-Za-z0-9]+Op::getOp(Constraints|Runtime)\(/ {
    match($0, /[A-Za-z0-9]+Op::getOp(Constraints|Runtime)/); sig = substr($0, RSTART, RLENGTH);
    fn = sig; sub(/::getOp.*/, "", fn); inC = (sig ~ /Constraints/) }
  inC && /ReasonForLackOfSupport::/ { match($0, /ReasonForLackOfSupport::[A-Za-z]+/); r[fn] = r[fn] " " substr($0, RSTART + 24, RLENGTH - 24) }
  inC && /constraintsDispatch|OpModel</ { real[fn] = 1 }
  END { for (f in r) print f, r[f], (real[f] ? "(conditional)" : "(always)") }' | sort

Only refusals inside getOpConstraints matter for validation; ignore getOpRuntime refusals (creation and weight-preparation ops), since the optimizer does not call getOpRuntime at all today. Do not read the file with a plain grep -B, the window lands on the previous function's name. On main the always-refusing set is DeallocateOp, MoeGptOp, PrepareConv3dWeightsOp, ToTensorSpecOp; TypecastOp refuses only for a system-memory input. Ops without the OpModel interface at all are skipped the same way.

A refusing op is skipped by OperationValidationAndFallback, so its metal asserts are not compile-time checked. But "not checked" is not the same as unprotected, and the two levels differ:

• Level 1 has no spill pass. A refusing op's cell is unprotected unless a workaround or verifier forces the condition.
• Level 2 runs GreedyL1SpillManagement first. On a NotImplemented op it evicts every live L1 tensor to DRAM interleaved (L1SpillManagement.cpp, the isNotImplemented path). So a "must be DRAM/interleaved" assert on such an op is forced by the spill at level 2 even though no query ran. Classify the cell after checking whether the spill produces the layout the assert wants.

For ops that do have a query, compare the argument list OpModel<Op>::getOpConstraints passes in lib/OpModel/TTNN/TTNNOpModel.cpp with what runtime/lib/ttnn/operations/** passes at dispatch. A program config, compute kernel config or optional tensor present at runtime but defaulted in the query means the query validates a different call. Mark caught (arg mismatch: <field>) and treat as LOOSE.