implement-mock

name	implement-mock
description	Workflow for adding a faithful mock of a silicon API to tt-emule. Composes `/arch-lookup` (find HW spec) → emule-mapping strategy selection → stub implementation → verification against sentinel + tt-metal regression.
user_invocable	true

/implement-mock — Add a faithful mock to tt-emule

You've identified a silicon API that tt-emule needs to mock (or a mock that's diverging from silicon). This skill walks the end-to-end workflow.

Usage

/implement-mock <silicon_function_or_concept>
/implement-mock noc_async_write_tile
/implement-mock llk_unpack_AB_matmul
/implement-mock TensorAccessor::get_noc_addr

Step 1 — Look up the HW spec

Run /arch-lookup "<silicon_function>" first. The arch-lookup skill gives you the authoritative HW behavior across wormhole / blackhole (and quasar if asked), AND a starting recommendation for the mock.

Read the sage output's "Emule mapping" + "Recommendation" sections carefully — they may already point at the right strategy.

Step 2 — Pick a strategy

A mock sits at one of three rungs of the kernel API stack (see docs/kernel-api-layers.md for the layer-1 / 1.5 / 2 / 3 vocabulary). Pick the strategy that matches the layer the silicon API bottoms out at:

Strategy	Kernel-API layer	When
A — stub in include/jit_hw/	layer-1 (PoR tile-op)	The silicon API is generic enough that one emule-side body works for every kernel that calls it.
B — per-op #ifdef __EMULE_JIT_MODE patch in the consumer's op header	layer-1.5 / layer-2 mix	The API can't be generically stubbed (takes a constexpr L1 firmware address, calls a small set of LLKs) but the consuming op can route through an emule-friendly equivalent.
C — semantic rewrite under __EMULE_JIT_MODE	layer-2 / layer-3 reframing	The op's TRISC body is deeply LLK / sfpi / TTI_*-bound (>5 LLK template errors on first compile). Re-express the op at a higher layer rather than stub every header below it.

Strategy A — stub in include/jit_hw/

Default for layer-1 lifts: new compute primitives, dataflow helpers, generic API surface. Example: noc_async_write, mul_tiles, recip_tile, pack_tile.

• Add the function to the appropriate header under jit_hw/
• If it operates on DST: use __emule_dst[idst][i] row-major float32
• If it operates on a CB tile: use __emule_compute::cb_read_ptr_at / cb_write_ptr_at + __emule_nfaces::rowmajor_to_nfaces[] permute
• If it's pipeline state (UNPACK/MATH/PACK config): no-op stub
• No __EMULE_JIT_MODE guard needed (jit_hw/ is inside -I ahead of tt-metal's hw/inc)

For LLK compute shims (functions like <op>_tile / <op>_tile_init under include/jit_hw/api/compute/), use /compute-llk-bringup directly — it specializes Strategy A with the shim-pattern catalog, sfpu_split_includes.h wiring, op→test-file mapping, polynomial-port recipe, and PCC triage flow for that scope.

Strategy B — per-op #ifdef __EMULE_JIT_MODE patch

Edit the consumer's op.hpp (or equivalent header) directly under __EMULE_JIT_MODE. Most common shape (constexpr L1 firmware address cast routed through __emule_local_l1_to_ptr):

volatile tt_l1_ptr T* p =
#ifdef __EMULE_JIT_MODE
    reinterpret_cast<volatile tt_l1_ptr T*>(__emule_local_l1_to_ptr(L1_ADDR));
#else
    reinterpret_cast<volatile tt_l1_ptr T*>(L1_ADDR);
#endif

Strategy C — semantic rewrite under __EMULE_JIT_MODE

• Gate off the entire silicon #include block AND the operator() body under #ifdef __EMULE_JIT_MODE
• Reimplement the op directly with CB read/write + __emule_dst[] math
• Validated 4× to date: RMSNorm, clamped_silu, Mcast, eltwise_mul

If unsure: default to A (stub in jit_hw/) if generic, B (per-op patch) if a single op uses it, C (semantic rewrite) if the LLK / sfpi chain is too deep.

Step 3 — Wire it in

Common scaffolding tasks:

1. Allowlist the op (if the consumer uses an __EMULE_JIT_MODE aggregator header): add the op to the #else/emule branch — append, don't reorder existing entries.

2. Wipe the JIT cache before testing:

   rm -rf /tmp/tt_emule_jit_* /tmp/tt_emule_src_*
   

   If the persistent cache at /tmp/tt_emule_jit_cache_* doesn't hash the aggregator/op headers, wipe it too; otherwise per-test temp dirs at /tmp/tt_emule_jit_*/ are enough.

3. Run the test that exercises the new mock.

4. Verify the sentinel test still passes (your smallest end-to-end emule test, used to catch broad collateral damage).

Step 4 — Verify

Pass bar:

• Target test: 100% parametrizations passing (or PCC > test's threshold)
• Sentinel test: still passes (your mock changes may have collateral damage on existing tests)
• tt-metal regression: totals must match the last known good baseline (no previously-passing test now failing):

  # default arch is wormhole n150; use the per-arch script (also
  # run_regression_blackhole.sh / run_regression_quasar.sh as needed)
  TT_METAL_DIR=<tt-metal-checkout> bash scripts/run_regression_wormhole.sh 2>&1 | tee regression-wormhole.log
  

If any of the above slips, see "Recovery" below.

Step 5 — Document

• If the mock is faithful to silicon: short comment at the impl site noting the HW behavior + reference (DeepWiki page, Confluence page, or LLK file:line).
• If the mock is intentionally simplified (no UNPACK/MATH/PACK pipeline state, no fabric, no real sfpi:: SIMD): note the simplification.
• If you discover a new pattern (e.g. a fresh instance of the semantic rewrite, or a brand-new strategy), add it to .claude/references/emule-mapping.md.

Batch mode

For sweeps (≥4 similar mocks at once), see /parallel-mock-implementation for the Workflow-tool dispatch pattern: one sub-agent per file, structured DONE/STUCK output, orchestrator handles centralized wiring afterward.

Step 6 — Time-box

Set a budget (e.g., 30 min for a non-trivial mock, 15 min for a sweep). If you're past the time-box and not converging, mark BLOCKED with a precise triage note + the deepest error message, and move on.

Recovery / rollback

• If a primitive change breaks the sentinel:

  git checkout <file>
  

  and move on with the next primitive.
• If an op's allowlist add poisons all kernels (sentinel breaks), remove the allowlist entry and move on.
• If tt-metal regression slips, bisect by reverting hunks; worst case, restore from a pre-change snapshot.

Anti-patterns

These cause more harm than good — avoid:

1. Stubbing without consulting the HW spec. A no-op stub may compile but produce wrong output, and that's harder to debug than a missing-symbol error.
2. Adding an op to the allowlist before verifying its op.hpp parses under TRISC. Sentinel breaks; everyone else also breaks.
3. Blanket rm -rf /tmp/tt_emule_jit_* while other workers run. Wave-1 swarm hit this. Sequential work is fine.
4. Implementing real sfpi:: SIMD math just to get one test to pass. The shim provides types only; for an op that needs real SIMD, use the semantic rewrite pattern instead.
5. Editing the silicon path in a consumer's aggregator header's #ifndef __EMULE_JIT_MODE branch. That branch isn't ours; keep edits under #else (or #ifdef __EMULE_JIT_MODE).

Related references

• docs/kernel-api-layers.md — layer-1 / 1.5 / 2 / 3 vocabulary the Step 2 strategies map onto
• .claude/references/emule-mapping.md — HW concept → emule mock catalog (§§1–6 + out-of-scope table)
• docs/api-injection-points.md — where in the pipeline emule injects
• CLAUDE.md — project conventions (clang-20, slow dispatch, always run regressions)
• BUILD_GUIDE.md — TT_METAL_DIR, regression scripts