implement-mock

/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 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, clampedsilu, Mcast, eltwisemul

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.

1. 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.

1. Run the test that exercises the new mock.

1. 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.

1. **Adding an op to the allowlist before verifying its op.hpp

parses under TRISC.** Sentinel breaks; everyone else also breaks.

1. Blanket `rm -rf /tmp/tt_emule_jit_` while other workers run.**

Wave-1 swarm hit this. Sequential work is fine.

1. **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.

1. **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