dataflow-cb-sync-audit

name	dataflow-cb-sync-audit
description	Audit circular-buffer (CB) producer/consumer flow control between data-movement (reader/writer) and compute kernels — cb_reserve_back/cb_push_back/cb_wait_front/cb_pop_front credit balance, data-write-before-credit ordering (NOC flush), reserve/wait-before-access, capacity vs num_pages, single-producer/consumer, counter cache-coherency, and remote/sharded CB credits. Use after touching any cb_* call, a reader/writer/compute kernel, fifo_rd_ptr/fifo_wr_ptr/pages_received/pages_acked, or RemoteSender/ReceiverCBInterface. Scope reaches beyond tt-llk into tt_metal/hw/inc/api/dataflow and ttnn/models kernels.
user_invocable	true

/dataflow-cb-sync-audit — Circular-buffer producer/consumer sync

Ground-truth precedence: the live ISA doc (tt-isa-docs MCP, fetched each run) and the live dataflow API source outrank every rule, table, and example baked into this skill — treat those as dated illustrations. If a live source contradicts a baked rule here, do NOT silently proceed: surface the conflict to the user and ask whether the baked rule should be overwritten, discarded, or kept. Default to the live source.

MANDATORY — before any verdict, read the shared grounding policy. The per-architecture source ladder (which docs to consult), the ground-or-abstain rule, and the Source preflight (list the sources you'll consult with their reachability + hierarchy, then PAUSE for the user) are defined once in race-audit-all → .claude/skills/race-audit-all/SKILL.md. Your FIRST action is to Read that file and follow its "Ground-truth source ladder", "Ground-or-abstain", and "Source preflight" sections — they are load-bearing: a verdict produced without them is ungrounded and MUST NOT be reported. If that file genuinely cannot be read, say so and abstain rather than proceed ungrounded. (If you were spawned by a race-audit-all sweep — your prompt already lists the confirmed sources — skip the Source preflight and do not pause; the orchestrator ran it once.)

Coverage — floor, not ceiling. The grep patterns and site lists in this skill are a seed, not an exhaustive enumeration. After running them, widen the search with full reasoning. The techniques here are illustrative examples, not the allowed set — use any approach your reasoning suggests, including ones not listed: e.g. semantic search (by behavior/effect, not just token), resolving macros / wrappers / typedefs / indirection the literal pattern can't match, following the call graph to callers and callees, and diffing the WH/BH/QSR variants to catch a site present in one arch and missing in another. If you can find a hazard, primitive, or site the encoded patterns don't cover — by any means — pursue and report it; do not clamp a stronger analysis to this list or to these techniques. State any residual coverage gaps explicitly (no silent caps).

Execution — parallel by default. When enumeration yields more than a few sites/files, fan out concurrent Agent calls by default (one per file/subsystem, a fresh context each), saturating the available concurrency (~10–16 at once); go inline only for a trivial set. The per-file fan-out described under Thoroughness is the default, not an exhaustive-only option. The cross-referencing/synthesis of results stays sequential (it must follow the per-unit findings). The heavyweight Workflow tool still requires explicit multi-agent opt-in — it is the opt-in exhaustive tier, not the default. Don't over-spawn a tiny diff.

Persisting results — single writer, incremental. Agents only return their findings; they never write a shared file (no concurrent-write clobbering). If findings are persisted to a file, the orchestrator/caller is the sole writer and appends each wave's returns as they arrive — incremental, never only-at-the-end — so an interrupt preserves every completed wave's findings.

Recall preflight — run the tool first (augmentor, not a verdict)

Before enumerating, run the deterministic cb-sync checker for a complete known-pattern worklist (reserve/push & wait/pop credit balance per CB):

# from the tt-metal repo root; do NOT cd into the tool dir (run.sh self-locates)
tt_metal/tt-llk/.claude/tools/llk-audit/run.sh <arch> --full-jit   # cb-sync runs in the kernel tier

cb-sync is empty over the tt-llk headers — CBs live in JIT-compiled kernels (ttnn/models), so it only emits findings when fed a KERNEL fact base via --full-jit (the on-request capture; runbook in race-audit-all). Treat findings[] as a floor, not a ceiling — widen per its blind_spots: it checks WITHIN-function call-count balance only, so it MISSES cross-kernel balance, a loop-multiplier imbalance (reserve in a loop vs push outside → equal static counts), and the data-before-credit NOC-flush ordering (that is noc-sync's join). If the kernel tier isn't built, cb-sync reports nothing and you audit the CB surface by reasoning — that is NOT "no findings".

The bug class (precise)

CBs are the credit-based FIFOs connecting producers (a reader on RISCV B, or the packer on T2) to consumers (compute on T0/T1, or a writer on NC). Flow control is two L1 counters per CB: tiles_received (a.k.a. pages_received, bumped by the producer in cb_push_back) and tiles_acked (pages_acked, bumped by the consumer in cb_pop_front). Misuse → data corruption (consumer reads a page before the producer finished writing it, or producer overwrites a page the consumer hasn't read) or deadlock (a wait whose credit never arrives). This is the dataflow layer the mailbox-sync / cfg-word-overlap "is the referenced data ready?" question hands off to.

Ground-truth mechanism (tt_metal/hw/inc/api/dataflow/dataflow_api.h, verified)

• Producer: cb_reserve_back(cb, n) spins until fifo_num_pages - (received - acked) >= n (free space) — reads the consumer's pages_acked via invalidate_l1_cache() + reg_read (it may be updated remotely). Then writes the n pages into L1 at fifo_wr_ptr. Then cb_push_back(cb, n) does pages_received += n and advances fifo_wr_ptr.
• Consumer: cb_wait_front(cb, n) spins until (received - acked) >= n (reads pages_received uncached). Then reads the n pages at fifo_rd_ptr. Then cb_pop_front(cb, n) does pages_acked += n and advances fifo_rd_ptr.
• Counters are 16-bit wrapping; comparisons use uint16_t subtraction (wrap-safe) — preserve that.
• Compute view (cb_api.h): unpack/T0 consumes input CBs (cb_wait_front/cb_pop_front), pack/T2 produces output CBs (cb_reserve_back/cb_push_back); the UNPACK()/MATH()/PACK() macros split one symmetric call. Dest-register handoff inside compute (tile_regs_acquire/commit/wait/release) is the MATH_PACK semaphore at API level → that's semaphore-handshake-audit's domain, not this one.
• Remote/sharded CBs: RemoteSenderCBInterface / RemoteReceiverCBInterface (circular_buffer_interface.h) propagate credits across cores via NOC atomic-increments / semaphore writes.

What to check

1. Ordering — data fully written/flushed BEFORE the credit (the main race). Producer must complete the page write before cb_push_back. For pages filled by NOC reads, a noc_async_read_barrier must precede cb_push_back; for cross-core credit (remote CB), the data NOC write must be ordered before the credit — and RemoteCB credits are NOC atomic-increments, which need the payload write committed (landed): noc_async_write_barrier (ACK), not a bare noc_async_writes_flushed() which only guarantees departure (data_movement_doc/general/posted_writes.md). Do not assume a flush-only atomic is safe even same-VC unicast — confirm against <arch>/NoC/Ordering.md + posted_writes.md (see noc-sync). Otherwise the consumer's cb_wait_front sees the page and reads stale/partial data. Symmetric on the consumer: finish reading before cb_pop_front (else the producer reuses the page). This is the CB analog of the mailbox fence ordering caveat.
2. Credit balance. Every cb_reserve_back(n) is matched by a cb_push_back(n) with the same n; every cb_wait_front(n) by a cb_pop_front(n). Pushing more than reserved (or popping more than waited) drifts the counters → producer overruns unread data / consumer consumes unwritten. Walk every branch and early-return.
3. Reserve-before-write / wait-before-read. A write to the CB region without a preceding cb_reserve_back, or a read without cb_wait_front, touches an unsynchronized page.
4. Capacity. n <= fifo_num_pages for every reserve/wait — asking for more pages than the CB holds can never be satisfied → deadlock.
5. Single producer / single consumer. The += on each counter is non-atomic and assumes one incrementer per side. Two producers or two consumers on one CB → counter race. (Multi-core fan-in/out must use the remote-CB interfaces, not raw shared counters.)
6. Counter cache-coherency. Each side reads the other side's counter, which a remote core may update via NOC — must use invalidate_l1_cache() / uncached reg_read (as the library does). A cached read → spin forever (deadlock) or proceed early (corruption).
7. Verify the credit semaphore addresses, the NOC ordering (data-before-credit), and that sender/receiver page counts agree.

Method

1. Enumerate CB sites:

   # from the repo root (ttnn/ and models/ are siblings of tt_metal/, NOT under it)
   grep -rInE '\bcb_(reserve_back|push_back(_hold_wr_ptr)?|wait_fronts?|pop_front)\b|\bremote_cb_(reserve_back|push_back_and_write_pages|wait_front|pop_front)\b|\.(reserve_back|push_back|wait_front|pop_front)[[:space:]]*\(|pages_(received|acked)|fifo_(rd|wr)_ptr|Remote(Sender|Receiver)CBInterface' \
     tt_metal/hw/inc/api ttnn/cpp models --include=*.h --include=*.cpp | grep -v '/tests/'
   

   Both API forms. Modern ttnn kernels use the OBJECT form cb_in.wait_front(1) / cb_out.reserve_back(1) (the deterministic tool recalls it via registry.cb_classify), NOT only the free cb_reserve_back(...) — the second alternation above catches it. Caveat: .push_back( also matches an unrelated receiver (e.g. std::vector::push_back), so when classifying, keep only receivers typed CircularBuffer/CBInterface/DataflowBuffer/DFBInterface (mirror the tool's _CB_RECV_TYPES; DataflowBuffer is the Metal 2.0 CB object with the same reserve_back/push_back/wait_front/pop_front quartet); the other three method names are CB-specific. Exhaustive run — no sampling. That grep yields the full kernel set (typically ~150–200 files across the CCL/dataflow families — all_gather, reduce_scatter, all_to_all, llama, moe, deepseek, broadcast, sdpa, matmul, prefetcher, unary/binary readers-writers). Enumerate them all into the run's coverage ledger and fan out (≥) one cell per kernel-family; do not sample — "sampled 6 of 186" is a blocking incompleteness per race-audit-all, not a caveat. This class is in scope (frontmatter): "the CB races live above tt-llk, run it separately" is not a valid skip — audit them here, now.
2. Per CB (per kernel pair), pair the producer (reserve_back+push_back) with the consumer (wait_front+pop_front) — they live in different kernels (reader↔compute, compute↔writer), so trace across the kernel set of the op.
3. Run checks 1–7. For ordering, confirm a NOC/read barrier sits between the data fill and cb_push_back (and between the read and cb_pop_front).

Verdict

• Balanced credits, reserve/wait before access, barrier orders data before the credit, n≤capacity, 1-producer/1-consumer, uncached counter reads → SAFE.
• cb_push_back reachable before the page write is flushed → DATA RACE (consumer reads stale/partial) — fix = add the missing noc_async_*_barrier / flush before the credit bump.
• reserve/push or wait/pop count mismatch on a reachable path → corruption or deadlock.
• n > capacity, or a credit that no producer/consumer ever bumps → DEADLOCK.
• Risk only in a specific op's kernel set, library primitives correct → LATENT/author-level — name the kernel.

Architecture note

The CB API is arch-agnostic (one copy in tt_metal/hw/inc/api); roles map as reader→RISCV B, writer→RISCV NC, compute→T0/T1/T2. Mechanism is the same across WH/BH/Quasar at this layer; the ordering primitives differ per NoC/arch — verify the barrier used matches the transfer type. Quasar caveat: under ARCH_QUASAR the local CB backing is a LocalDFBInterface (tc_slots / rd_offset / stride_size), NOT LocalCBInterface / fifo_rd_ptr — the fifo_rd_ptr/fifo_page_size pointer reasoning here is WH/BH; on Quasar the per-tile spacing is stride_size.

Verified ground truth (don't re-derive)

• cb_reserve_back/cb_wait_front/cb_push_back/cb_pop_front semantics above are from dataflow_api.h (credit math, uncached counter read, 16-bit wrap). Canonical compute pattern (e.g. eltwise_bw_gelu_poly.cpp): cb_reserve_back(out) + cb_wait_front(in) → per-tile tile_regs_acquire/commit/wait + pack_tile → cb_pop_front(in) + cb_push_back(out). The tile_regs_* half is MATH_PACK (defer to semaphore-handshake-audit); this audit owns the cb_* half.

Cross-references (where this audit meets the others)

• mailbox-sync / cfg-word-overlap "is the referenced memory ready?" → discharged HERE: the address the mailbox broadcasts is derived from fifo_rd_ptr, valid only if cb_wait_front gated correctly AND the producer ordered its write before cb_push_back.
• Data-before-credit ordering ↔ mmio-race / NOC memory-ordering (the barrier discipline).
• Dest-register tile_regs_* interleaved with cb_* ↔ semaphore-handshake-audit (MATH_PACK).

Output

For each CB / kernel pair: file:line of producer (reserve/push) and consumer (wait/pop), credit balance (n match), the barrier ordering data before the credit (present/missing, both sides), capacity check, producer/consumer cardinality, counter-read coherency, remote-CB credit path if applicable, verdict (SAFE / DATA-RACE / DEADLOCK / LATENT) + one-line fix. End with totals.