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.
Instrucciones de origen · Vista previa de solo lectura
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.
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 rather than proceed ungrounded. (If you were spawned by a sweep — your prompt already lists the confirmed sources — skip the Source preflight and do not pause; the orchestrator ran it once.)
abstain
race-audit-all
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.
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
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.
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.
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.
Capacity.n <= fifo_num_pages for every reserve/wait — asking for more pages than the CB holds can never be satisfied → deadlock.
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.)
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).
Remote/sharded CBs. Verify the credit semaphore addresses, the NOC ordering (data-before-credit), and that sender/receiver page counts agree.
Method
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.
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.
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.
