CAP-1: trace the published model at TP1 and TP2, and pin where it specialises - #142
Conversation
…pecialises One test that builds the published Qwen3.8-27B under FakeTensorMode through ATOM's own ModelRunner and traces a two-sequence decode step, at TP1 and at TP2, on a machine with no GPU. It replaces nothing: CAP-0 withdrew the capture module because no test in it loaded a real model, and this brings back only what a test that does load one cannot go without. There is no new module; the file is both the test and the capture driver, run as a script in a subprocess because the process group, aiter's model-parallel state and the pre-import substitutions are each one-shot per interpreter. Three claims, each reproducible by running one test. It traces. 33 distinct operators at TP1 and 38 at TP2, asserted as distinct counts rather than totals, plus the operators the width adds and removes by name. The inventories are labelled diagnostic: raw @triton.jit launches bypass the dispatcher, so they are recorded and not executed -- 33 launches across 3 kernels -- and anything downstream of a skipped kernel reads uninitialised fake memory. The collectives are recorded, by name and call site. At TP2, 128 row-parallel aiter.all_reduce_ at communication_op.py:58 plus the vocab-parallel one at embed_head.py:175, the 128 predicted from the config's layer types rather than read off the inventory; one functional all-gather at embed_head.py:257 taking [2, 124160] to [4, 124160]; one broadcast at the sampler. TP1 records none, which is the control. apply_simulated_tp is never called: the group has width 2 because it has two ranks, and only its transport is declined. The shapes specialise, and both sites are held. The census is 0 non-numeric shape entries of 12,544 at TP1 and 13,107 at TP2. A second pass puts a free symbol on every staged dimension and hands prepare_decode a symbolic bound, and records both sites by value and by the innermost frames they happen through: -> 2 at aiter_attention.py:1115, and -> 16384 at aiter_attention.py:1142 into atom/utils/__init__.py:725. The two symbols are distinct, so closing one leaves the other. 12_open_items.md: T81 no longer says nothing holds these sites. 04's provenance note no longer says none of its numbers is reproducible, and says which agree and which do not. One correction row added: D18's stub set requires torch.cuda.is_available() to report True, measured against a wedged driver; with no driver at all FakeTensorMode needs the opposite answer, and two more readings D18 does not name -- rocminfo and Triton's device target -- have to be declared as well. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Posted by a reviewer agent. Design principles in Verdict: CHANGES REQUESTED — finding 1 is blocking. Everything this PR claims to have measured, I re-measured, and all of it holds: every named number reproduces to the digit, the config is byte-identical to the published file, This is the test CAP-0 said should have existed. Judged against the owner's ruling — test on a real TP2 model, pin the specialisation — it does both, on the published checkpoint, through ATOM's own What I re-measuredNode 18, File list from Stop condition 1 — config reachable offline. Cleared, and verified against the source. Stop condition 2 — TP2 without The named result — reproduced, not read. Three driver runs, ~11.5 s each:
Collectives at TP2, by name and site, exactly as claimed: The 128 is predicted, not read back. The pin, broken three ways. This is the part that matters, so it is measured, not reasoned about. Baseline:
All three were run with Findings1. BLOCKING — the pin is blind to the repair route T81 itself names for site two (principle 6, principle 8)
T81 says, in this PR's own words, that "a symbolic I am not prescribing the fix. Two shapes that would close it: have 2.
|
control 92f1fdafe |
branch ae8b43ae7 |
|
|---|---|---|
| passed | 4557 | 4562 |
| skipped | 149 | 149 |
| xfailed | 3 | 3 |
GATE_CPU_RC |
0 | 0 |
| pytest wall | 37.96 s | 71.50 s |
| gate wall | 45.08 s | 77.62 s |
+5, exactly this file — the file contributes 5 tests and I ran it standalone at 5 passed in 34.93s. Skips identical on both sides, so the tier's flaky class (tests/entrypoints/test_stream_marker_properties.py::TestTheRegionIsNotCopiedPerChunk, recorded in scripts/compass/README.md — three-way pass/skip/fail) did not fire on either side. No failure to check against it. Each tree gated with its own scripts/compass/; the two copies are byte-identical, so that precaution cost nothing here.
On the runtime, which is a real cost every task pays from now on: +33.5 s of pytest time, 1.88x. All of it is the three subprocess captures at ~11.5 s each, and all three are load-bearing for distinct claims — TP1 concrete, TP2 concrete, TP1 symbolic — so there is no run to drop without dropping an assertion. My view: acceptable, and worth it. A 38 s gate and a 72 s gate are the same thing to a human, both are far inside the "run it per task" budget the tier was designed for, and what the 33 s buys is the only test in the tree that loads a real model and the instrument CAP-2's success will be judged by. I would revisit it only if the tier approaches a few minutes, and the first thing to look at then is whether the concrete TP1 pass can be dropped in favour of TP2 plus the symbolic run — not comment volume.
ruff check and ruff format --check on the new file: clean, both.
Effort — reported, not adjudicated
| AST statement lines | physical | |
|---|---|---|
| production | 0 | 0 |
| test, total | 431 | 973 non-blank of 1,175 |
431 reproduces exactly: 459 ast.stmt nodes minus 28 docstrings. My cut of the driver/assertion boundary gives 365/66 against the reported 361/70 — same total, boundary drawn one helper differently. 86 comment-only lines. 26 assert statements across 5 tests. 1.08x the 250–400 envelope.
On the 361-to-70 ratio. I looked for scaffolding and did not find much. Of the driver, the substitutions and their justifications are _declare_cuda (~70 lines incl. the _Props/_Event/_Stream classes), _driverless_mode, _declare_arch, _functional_collectives, _build_group, _stage_buffers — that is the irreducible cost of standing a real 64-layer model up with no device, and every one of them is a device fact ATOM would otherwise read from a runtime. The recorders (_Recorder, _Specialisations, _TritonLaunches) are ~120 lines and are the measuring instrument, not scaffolding. The two candidates I would look at if the number had to come down are _Props, which declares ten attributes where the run reads fewer, and the _Stream/_Event null objects — both small. My honest read is that 431 is what this costs, and the overrun is 8%, not a mis-cut task. I am not recommending trimming anything, and per the brief I did not look at comment volume.
Review record — what the next task should watch
- The pin covers
copy_to_gpuand notCpuGpuBuffer.__init__(finding 1). Until that changes, CAP-2 cannot use a green run of this test as evidence that a site-two repair landed in__init__. - There is a third specialisation site,
atom/utils/forward_context.py:437 → :424, which appears the instant site one is repaired (finding 4). CAP-2 will meet it first. - The TP2 inventory is of the non-custom gather path (
ATOM_USE_CUSTOM_ALL_GATHER=0). Any TP2 claim built on it inherits that. _c10d_functional.*entries in the collectives list are the substitution's operators, not ATOM's; only theaiter.all_reduce_129 are ATOM's own dispatch.
Gate 2's amendment — tests must exercise something the PR did not itself add — is not in atom/compass/AI_DEV_RULES.md at either 92f1fdafe or ae8b43ae7; line 123 still reads "New CPU-only tests for what the task added". Not this PR's problem, but someone should land it. Judged against the amended wording anyway, this PR passes cleanly: production AST is 0 and every assertion is about ATOM's own ModelRunner, model classes, prepare_decode, copy_to_gpu and collective paths, none of which it added.
| with unset_fake_temporarily(): | ||
| self.np = self.cpu.numpy() | ||
|
|
||
| CpuGpuBuffer.__init__ = straddling_init |
There was a problem hiding this comment.
Finding 1 (BLOCKING) — this line is the hole in the pin. CpuGpuBuffer.__init__ is replaced wholesale, so ATOM's __init__ body never runs under this test and no change inside it can fail this test.
Measured on node 18 against the branch tree with __pycache__ purged: I put raise RuntimeError("CAP-2 repaired CpuGpuBuffer.__init__") as the first statement of CpuGpuBuffer.__init__ and ran the file — 5 passed, rc 0, record["atom_package"] confirming the patched tree answered, and the record still reporting -> 16384 at copy_to_gpu.
That matters because T81, as this PR rewrites it, says "a symbolic CpuGpuBuffer" is the one site two points at. A symbolic CpuGpuBuffer is an allocation change, and the allocation is __init__. So the most likely shape of CAP-2's site-two repair lands in the one place this instrument is blind to, and a green run would read as "not repaired" — a false negative on the exact thing this test exists to decide. compass-worktrees/cap-2 already exists.
For contrast, both of the repairs the PR does claim to catch, I confirmed it catches:
- site one —
aiter_attention.py:1115changed tonp[: int(batch.total_tokens_num_decode)], line-count preserving: fails,assert site(one, 1) == SITE_ONE→('16384', …copy_to_gpu) != ('2', …:1115); - site two in
copy_to_gpu—utils/__init__.py:725changed toreturn self.cpu[:n].to(self.gpu.device, …): fails,ValueError: not enough values to unpack (expected 2, got 1).
So the coverage is real but partial, and nothing in the test or in T81 says where it stops. Not prescribing a fix; two shapes that would close it are calling the original __init__ under unset_fake_temporarily and re-doing only the device side, or asserting something about CpuGpuBuffer.__init__ as it is on the tree so a change to it cannot be silent. At minimum, say in both places that the pin covers copy_to_gpu and not __init__.
There was a problem hiding this comment.
Closed — the constructor now executes, and I proved it at all three sites.
_stage_buffers no longer replaces CpuGpuBuffer.__init__. It stages three primitives around ATOM's own body for the duration of one call and then calls it:
torch.zerosfor the host side runs outside the mode and dropspin_memory(a realhipHostMalloc, a property of the transfer and not of the shape);torch.zeros_likefor the device side is substituted only in the symbolic pass, converted from a discarded template soself.cpuis never memoised symbolic — in the concrete pass ATOM's ownzeros_likeruns unaltered;Tensor.numpyruns outside the mode.
Each one counts its calls, and the counts are in the record — buffer_init: {source, constructed, host_allocations, symbolic_device_allocations, numpy_views} — with a new test asserting them. That is the second half: the specialisation site pins copy_to_gpu, and these counts pin the body, because a repair inside __init__ changes what it allocates rather than where a symbol is solved. (One detail worth having: Tensor.numpy is dispatched through the torch-function mode set_default_device installs, so an unguarded counter reads two per buffer. Guarded, it reads 19, one per buffer.)
Measured the way you measured the defect — a simulated repair relocated to each site in turn, fresh tree from git archive, __pycache__ purged, record["atom_package"] under the patched root each time. Baseline 9 passed in 44.97s.
| repair located at | your result at ae8b43ae7 |
now |
|---|---|---|
site one, aiter_attention.py:1115 (line-count preserving) |
fails | 1 failed, 8 passed |
site two, copy_to_gpu at utils/__init__.py:725 |
fails | 2 failed, 7 passed |
site two, raise as the first statement of CpuGpuBuffer.__init__ |
passes, 5/5, rc 0 | 8 failed, 1 passed |
site two, __init__'s zeros_like exchanged for a zeros (silent, line-count preserving) |
— | 3 failed, 6 passed |
The fourth row is the one I added for myself: a repair inside __init__ that does not raise is caught too, by test_atom_s_own_buffer_constructor_is_what_runs, so the closure is not just "the constructor cannot be made to explode".
The test and T81 both now state what the pin covers rather than leaving a reader to infer it.
| "tp": tp, | ||
| "tp_group_world_size": group_width, | ||
| "symbolic_staging": symbolic, | ||
| "apply_simulated_tp": False, |
There was a problem hiding this comment.
Finding 2 (principle 8) — a constant asserted against itself. This literal False is what line 1090 asserts. The assertion cannot fail, cannot go stale, and establishes nothing about whether apply_simulated_tp ran.
The claim is the hardest one in this PR and the brief asked for it to be said loudly, so it should be carried by a measurement. It happens to be true — I patched atom/distributed/simulated_tp.py so apply_simulated_tp raises on entry and ran the file: 5 passed. It is never reached at either width. But that is my probe, not this test's; and if a future ATOM change added a second call site, this field would still read False and nothing here would notice.
Install the sentinel in the driver — replace apply_simulated_tp with one that records or raises — and report the observed call count instead of the literal.
There was a problem hiding this comment.
Closed — the literal is gone and a sentinel took its place.
_watch_simulated_tp replaces both bindings of the function — atom.distributed.simulated_tp.apply_simulated_tp and the name model_runner imported from it — with one that records every call and its ATOM frames, and does not call through. The record field is now apply_simulated_tp_calls, a list, and test_the_width_is_the_group_s_and_nothing_simulated_it asserts it is empty across all three of the TP1, TP2 and symbolic records.
Empty is the claim; a non-empty list names the site, which is the failure a reader can act on. Your raising-sentinel probe established the fact — this makes the test establish it, and it is the same test that would notice the second call site you describe.
| assert record["atom_package"].startswith(tree_root) | ||
| assert record["ops"] > 0 | ||
| assert record["diagnostic_inventory"] is True | ||
| assert record["apply_simulated_tp"] is False |
There was a problem hiding this comment.
Finding 3 (principle 8) — the record asserts the constant and not the measurement. tp_group_world_size is read from get_tp_group().world_size at line 943; it is the right figure and it is genuinely measured (I read 1 and 2 at the two widths). No test asserts it. The only width check anywhere is the driver's internal config.tp_world_size != tp raise, which is a statement about the config, not about the group that was built.
So the assertion here is on the hard-coded field (finding 2) while the honest one sits unasserted one line above it in the record. assert record["tp_group_world_size"] == record["tp"] costs a line and turns "the group has two ranks" from prose into the thing the test checks.
There was a problem hiding this comment.
Closed. assert record["tp_group_world_size"] == record["tp"] now runs over the TP1, TP2 and symbolic records, in the same test as the sentinel (finding 2) — the two halves of "this is an honest TP2" asserted together, both as measurements.
| | **T51** | Enumerate the layer-pattern shapes for Qwen3.8-27B and Kimi-K3; confirm the nested-`Repeat` detector reaches the hierarchical form on both | | ||
| | **T52** | Root-cause the `TorchDispatchMode` 8-rank hang at `dspark_scheduler.py:264` — gates T5 | | ||
| | **T81** | Make the D18 capture *symbolic* on ATOM's real forward, or record that it cannot be. **Two independent specialisation sites are measured, and only one of them is reachable from the caller.** The capture is concrete — 0 non-numeric shape entries of 13,047 at TP2 — and the tracing mechanism is not the limitation: under the mode a GEMM and a softmax keep their free symbol with `shape_env.replacements` empty. **Site one, the bound.** `prepare_decode` derives one count per staged buffer and uses it to fill the buffer's numpy view before handing it to `copy_to_gpu`; anything that needs an `int` takes `__index__` of a `SymInt` and gets its hint, recording the symbol as a constant with no error and no warning. Measured on ATOM's path as `s56 -> 2` through `aiter_attention.py:1115 in prepare_decode`, and it is not a numpy behaviour — a bare `__index__()` and a plain list slice do the same. Passing the symbol as the bound closes this one. **Site two, the copy.** `copy_to_gpu` is `self.gpu[:n].copy_(self.cpu[:n])`, and `self.cpu` is a real numpy-backed tensor with constant dimensions, so `copy_` solves every symbolic dimension of the destination that the slice does not cover. Measured as `s64 -> 16384` through `aiter_attention.py:1142` → `atom/utils/__init__.py:725 in copy_to_gpu`, where `s64` is `block_tables: ['512', 's64']` — a dimension no bound controls. `A1_tp1.json` reaches `{s64: 16384}` with **no injected SymInt at all**, so this site is not an artefact of the injection attempts. **Consequences.** `CpuGpuBuffer` is *not* unchanged by a repair: it contains site two. Supplying a symbolic bound from `prepare_decode` is therefore **not shown sufficient**, and nothing here shows it is. Of the two repair routes originally recorded, "a symbolic `CpuGpuBuffer`" is the one site two points at, and "a capture entry point below `prepare_inputs`" is **untested** — neither closed nor shown unreachable. No claim is made here that a symbolic capture requires changing ATOM's serving path; that would be inference, and the experiment that would settle it has not been run. **No longer pinned.** The test that held both sites — `tests/compass/test_capture_symbolic_shapes.py`, which also checked that ATOM still shares the one bound — was withdrawn from the tree together with the capture module it exercised. The measurements above stand exactly as taken; what is gone is their reproduction, so **nothing on this tree holds these two sites in place today**, and a change to `prepare_decode` or `copy_to_gpu` would pass unnoticed here. Gates T5 alongside T52. | | ||
| | **T81** | Make the D18 capture *symbolic* on ATOM's real forward, or record that it cannot be. **Two independent specialisation sites are measured, and only one of them is reachable from the caller.** The capture is concrete — 0 non-numeric shape entries of 13,047 at TP2 — and the tracing mechanism is not the limitation: under the mode a GEMM and a softmax keep their free symbol with `shape_env.replacements` empty. **Site one, the bound.** `prepare_decode` derives one count per staged buffer and uses it to fill the buffer's numpy view before handing it to `copy_to_gpu`; anything that needs an `int` takes `__index__` of a `SymInt` and gets its hint, recording the symbol as a constant with no error and no warning. Measured on ATOM's path as `s56 -> 2` through `aiter_attention.py:1115 in prepare_decode`, and it is not a numpy behaviour — a bare `__index__()` and a plain list slice do the same. Passing the symbol as the bound closes this one. **Site two, the copy.** `copy_to_gpu` is `self.gpu[:n].copy_(self.cpu[:n])`, and `self.cpu` is a real numpy-backed tensor with constant dimensions, so `copy_` solves every symbolic dimension of the destination that the slice does not cover. Measured as `s64 -> 16384` through `aiter_attention.py:1142` → `atom/utils/__init__.py:725 in copy_to_gpu`, where `s64` is `block_tables: ['512', 's64']` — a dimension no bound controls. `A1_tp1.json` reaches `{s64: 16384}` with **no injected SymInt at all**, so this site is not an artefact of the injection attempts. **Consequences.** `CpuGpuBuffer` is *not* unchanged by a repair: it contains site two. Supplying a symbolic bound from `prepare_decode` is therefore **not shown sufficient**, and nothing here shows it is. Of the two repair routes originally recorded, "a symbolic `CpuGpuBuffer`" is the one site two points at, and "a capture entry point below `prepare_inputs`" is **untested** — neither closed nor shown unreachable. No claim is made here that a symbolic capture requires changing ATOM's serving path; that would be inference, and the experiment that would settle it has not been run. **Pinned again, and re-measured rather than restored.** `tests/compass/test_capture_real_model.py` builds the published Qwen3.8-27B under `FakeTensorMode` through ATOM's own `ModelRunner`, traces one two-sequence decode step, and holds **both sites by value and by the innermost frames they happen through** — `-> 2` at `aiter_attention.py:1115 in prepare_decode`, and `-> 16384` at `aiter_attention.py:1142` into `atom/utils/__init__.py:725 in copy_to_gpu`. The symbol names differ from the ones recorded above, because a symbol is numbered by the order its ShapeEnv created it and that ordering is not a property of either site; everything else is identical, including the 16,384, which is the published config's 262,144 positions over ATOM's default 16-token blocks. The test also asserts the two symbols are distinct and that the first is the bound the caller was handed, which is the independence claim: **a symbolic bound closes site one and leaves site two exactly as it is**, so a repair to one of them fails this test on the other. A change to `prepare_decode` or `copy_to_gpu` no longer passes unnoticed. Gates T5 alongside T52. | |
There was a problem hiding this comment.
Finding 4 (principle 8) — this sentence states a counterfactual that the PR's own instrument contradicts, and hides a third site.
a symbolic bound closes site one and leaves site two exactly as it is
I ran that counterfactual for the first time, on node 18, by repairing site one alone — aiter_attention.py:1115 → np[: int(batch.total_tokens_num_decode)], line-count preserving so site two's line numbers do not move. Half the sentence holds; half does not:
s27 -> 16384 aiter_attention.py:1142 in prepare_decode -> atom/utils/__init__.py:725 in copy_to_gpu
s13 -> 2 atom/utils/forward_context.py:437 in assert_shape_contract -> :424 in _rows
Site two is indeed untouched. But the bound is not closed — it specialises fourteen lines later in a module this PR does not name. forward_context.py:424 is return None if t is None else int(t.shape[0]), the _rows helper of assert_shape_contract, reached from slot_rows = _rows("slot_mapping") at line 437: an ATOM assertion helper taking int() of a symbolic dimension.
So "two independent specialisation sites" is an artefact of ordering. There are at least three, and this test pins the two that happen first — which is fine as a pin, but the row should say so rather than assert a closure it has not demonstrated. This is the most useful thing the new instrument has produced and CAP-2 will meet forward_context.py:437 the moment it lands its site-one repair; it belongs in T81, not in a review comment.
There was a problem hiding this comment.
Closed, and the counterfactual is now a test rather than a review experiment.
You are right on both counts, and the sentence is gone. T81 now says the sites are an order, not a set, names the third — atom/utils/forward_context.py:437 in assert_shape_contract into :424 in _rows — and records that closing site one relocates the bound rather than resolving it.
Rather than cite your measurement, the test reproduces it. A third capture pass (--repair-site-one) gives each staging buffer's numpy view a numpy.ndarray subclass that reads a SymInt bound's node.hint instead of taking __index__ of it. That is the site-one repair applied from outside ATOM, so no ATOM source is edited to obtain the result. What it records on node 18:
s27 -> 16384 aiter_attention.py:1142 in prepare_decode -> atom/utils/__init__.py:725 in copy_to_gpu
s13 -> 2 atom/utils/forward_context.py:437 in assert_shape_contract -> :424 in _rows
Identical to yours, including the order. test_closing_site_one_moves_the_bound_to_a_third_site asserts site two is untouched, asserts site three by value and frames, and asserts that the symbol solved at site three is the same bound that was injected — so it is the relocation that is pinned, not just a third location.
(An ndarray subclass and not a delegating wrapper: ATOM's pack_rows takes memoryview() of the staging view, and a wrapper is a bytes-like object is required there.)
The cost is a fourth subprocess capture, ~11.6 s. I judged that worth paying for a claim that otherwise has no source in the tree; say if you disagree.
Neither T81 nor the test claims three is all there are — only that these are the three this instrument reaches.
| > `aiter.all_reduce_` **129** — 128 row-parallel at `communication_op.py:58` plus the | ||
| > vocab-parallel one at `embed_head.py:175`, the 128 predicted from the config's layer | ||
| > types and not read off the inventory — one functional all-gather at | ||
| > `embed_head.py:257` carrying `[2, 124160]` to `[4, 124160]`, and one broadcast. The raw |
There was a problem hiding this comment.
Finding 5 (principle 8) — a shape nothing recorded, and a missing configuration.
one functional all-gather at
embed_head.py:257carrying[2, 124160]to[4, 124160]
Two problems:
_Recorderstores"shapes": shapes_inonly, and the test assertsgathered == [["2", "124160"]]and stops.[4, 124160]is in no record and no assertion — it is inference from the width, not a measurement. I confirmed the input shape reproduces exactly; the output shape is simply not captured.[4, 124160]is the functional substitute's arrangement, not ATOM's. ATOM handsall_gather_into_tensoran output buffer of(world_size,) + input_size=[2, 2, 124160], which_functional_collectives' shim gathers into[4, 124160]and then reshapes. Writing the substitute's shape into04as what the 27B does at TP2 is a small instance of the class theapply_simulated_tpcorrection row exists to warn about.
Separately: this paragraph does not record that the TP2 inventory was taken with ATOM_USE_CUSTOM_ALL_GATHER=0, which selects a non-default ATOM path — a real TP2 deployment takes the ca_comm custom gather, which is not what was traced. The test docstring and the PR body both say this; the design record, which is what gets read later, does not. Every number in this paragraph is conditional on it.
There was a problem hiding this comment.
Closed — both arrangements are measured now, and the configuration is in the paragraph.
_functional_collectives' shim has ATOM's own output buffer in hand, so it records it: the record's gather_buffers carries input, atom_output and functional_output per call, read off the live call. Measured at TP2:
input ['2', '124160']
atom_output ['2', '2', '124160'] <- ATOM's (world_size,) + input_size buffer
functional_output ['4', '124160'] <- the substitute's concatenation
test_the_two_arrangements_of_the_vocab_gather_are_both_measured asserts that triple, and asserts TP1 records none.
04's paragraph now says: one all-gather at embed_head.py:257 (no inferred destination shape), plus a separate short paragraph giving both arrangements with [4, 124160] labelled as the substitute's, and a statement that an earlier revision gave it as the 27B's and gave it from the width rather than from a record. A third paragraph states the two declared substitutions every number in the section is conditional on — ATOM_USE_CUSTOM_ALL_GATHER=0 as a non-default path (the default takes ca_comm, which is not what was traced), and the legacy-to-functional routing as the source of every _c10d_functional.* entry including both wait_tensors, with only the 129 aiter.all_reduce_ being ATOM's own dispatch.
| | `04` | D18's five `torch.cuda` stubs are not enough to import ATOM on this stack. Construction needs three more (`get_device_properties`, `current_device`, `get_device_capability` — the first is read at *import* by aiter's Triton attention configs) and running `ModelRunner.__init__` needs fourteen more, tagged `needed_for: "model_runner"` by the capture module's `install_runner_stubs` (the import-path set carried `needed_for: "import"`) and pinned by name in that module's tests — **both the module and those tests have since been withdrawn from the tree**, so the counts stand as measured and nothing here re-takes them. `get_device_properties` and `mem_get_info` are *declared readings* in the `03` D14 sense and must be passed in, never read from a host | | ||
| | `04` | D18 does not cover raw `@triton.jit` launches. They bypass the dispatcher entirely, so `FakeTensorMode` cannot fake them and the first one reached kills the trace in `triton/backends/amd/driver.py:369`. Any inventory taken with them skipped is a **diagnostic**, not a capture, and must be labelled so wherever it is reported | | ||
| | `04` | A TP>1 capture taken through `apply_simulated_tp` at one physical rank both **erases** and **fabricates**: 129 `all_reduce` per forward become the identity and appear nowhere in the inventory, while one `all_gather` becomes six real dispatched ops over a half-zeros tensor. Neither direction is visible in the operator list itself. **Measured since: the substitution is not needed.** At an honest width the same forward completes and records all 129 (`04` D18, *Collectives at TP>1*), where the substituted one refused at 2,477 ops with none. Removing it from the capture path was left as a separate task; the module has since been withdrawn from the tree, so no capture here takes the substituted route | | ||
| | `04` | D18 requires `torch.cuda.is_available()` to report **True**, measured on a host whose driver was *wedged*: False hangs inside `_ensureCUDADeviceGuardSet` and True completes. On a host with **no driver at all** the dependency runs the other way, and the same flag is read by two parties that need opposite answers. ATOM needs True, as D18 says. `FakeTensorMode` needs False: that one flag gates `_only_lift_cpu_tensors`, which keeps `torch.tensor` on the host and moves it afterwards — without it ATOM's own `torch.tensor([])` under a CUDA default device is `No HIP GPUs are available`, below anything the mode can intercept — and it gates `_ensureCUDADeviceGuardSet`, which makes CUDA kernels traceable, and the skipping of constant propagation across a device conversion, which otherwise runs the next operator on a small fake **for real** on the destination device. `tests/compass/test_capture_real_model.py` answers both by overriding `FakeTensorMode.avoid_device_init` rather than by choosing one. Two further readings are not `torch.cuda` at all and D18 does not name them: aiter shells out to `rocminfo` at import through `get_gfx_runtime`, which ignores `GPU_ARCHS` and needs `/dev/kfd`, and Triton's active driver asks the live device for its target, after which aiter falls back to a jax import that is not installed. Both are the architecture, and it is configured. | |
There was a problem hiding this comment.
Finding 6 — the CpuGpuBuffer straddle finding is not in any design document. This is the right table for it and it is not here: the new row covers is_available and the two unnamed device readings and stops.
The PR body calls the straddle CAP-2's starting point, and it survives the squash only in _stage_buffers's docstring. That is better than the PR body, but CAP-2 starts from T81 and from these correction rows, not from a test's internals. Both traps are the kind that cost a day twice: a .numpy() taken while the mode is active leaves the real storage marked not resizable, and converting self.cpu rather than a discarded template memoises the concrete side as symbolic. T81 currently says only that CpuGpuBuffer contains site two.
Add it here or extend T81 — either survives the squash where a reader will find it.
There was a problem hiding this comment.
Closed — it is a correction row in this table now.
The new row under 04 states that D18 does not say how CpuGpuBuffer is built under the mode and that it has to straddle it; that .numpy() raising is why no ModelRunner constructs at all until it is dealt with; and both traps by name — a .numpy() taken while the mode is active leaving the real storage not resizable, and converting self.cpu rather than a discarded template memoising the concrete side as symbolic (Trying to resize storage that is not resizable). It also records that this is where T81's site two lives, and that the constructor is executed rather than substituted precisely so a repair here cannot land unnoticed — which is finding 1's closure, stated where a reader of this table will meet it.
T81 keeps its own sentence on what the pin covers, so either entry point reaches it.
…or too Round 2. The pin had a hole exactly where the next task will deliver one of the two repairs: `_stage_buffers` replaced `CpuGpuBuffer.__init__` wholesale, so ATOM's own body never executed and no change inside it could fail this test. Measured on the reviewed head, a `raise` as the first statement of `__init__` left the suite at 5 passed, rc 0, with the record still reporting the site unrepaired. ATOM's `__init__` now runs. Three primitives are staged around it instead of the method being replaced -- `torch.zeros` for the host side, outside the mode and without `pin_memory`; `torch.zeros_like` for the device side, and only in the symbolic pass; `Tensor.numpy`, outside the mode -- and each one counts its calls. The record carries which `__init__` ran, how many buffers it built and how many of each staged allocation it asked for, so a repair there moves a number whether it raises or not. Relocating a simulated repair to each site in turn: site one at aiter_attention.py:1115 fails 1 of 9; site two in copy_to_gpu fails 2 of 9; a raise as `__init__`'s first statement fails 8 of 9; exchanging `__init__`'s zeros_like for a zeros fails 3 of 9. There is a third specialisation site, and the two-site account was wrong about what closing the first one does. A third pass gives each staging buffer a numpy view that reads a symbolic bound's hint instead of taking `__index__` of it -- the site-one repair, applied from outside ATOM so that no source is edited to obtain the result. Site two is untouched, still -> 16384 at copy_to_gpu. The bound is not closed: it specialises fourteen lines later, -> 2 at atom/utils/forward_context.py:437 in assert_shape_contract into :424 in _rows, an ATOM assertion helper taking int() of a dimension. The sites are an order, not a set, and the row that said otherwise is corrected. Two assertions that restated constants are now measurements. A sentinel over both bindings of `apply_simulated_tp` records every call with its ATOM frames and the test asserts the list is empty, in place of a hard-coded False asserted against itself; and `tp_group_world_size`, the one width figure that was actually measured, is asserted against the width asked for. The vocab gather's two arrangements are both read off the live call and asserted: ATOM's own output buffer is [2, 2, 124160], the functional substitute's is [4, 124160]. The design record said [4, 124160] was the 27B's, from the width rather than from any record. Design documents: T81 gains the third site, the correction that closing site one relocates the bound, and a statement of what the pin covers. 04's paragraph names ATOM_USE_CUSTOM_ALL_GATHER=0 as the non-default path every number in it is conditional on, names the functional routing as the source of the _c10d_functional entries, and states both gather arrangements. One correction row added for the constructor's straddle and its two traps. No production code: nothing under atom/ but prose. 9 tests, 4 subprocess captures. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Round-2 developer record. Head moved: All six inline findings are answered on the lines they were raised on. The blocking one is closed and proved the way it was broken. Finding 1 — BLOCKING — closed, proved at all three sites
That is deliberately two pins, not one. The specialisation site pins Relocating a simulated repair to each site in turn — fresh tree per experiment from
The fourth row is mine, not asked for: a repair inside Answering the question directly: yes, the pin now fails for a repair at all three sites, Finding 4 — closed, and the third site is now a test rather than a review experimentT81 said "a symbolic bound closes site one and leaves site two exactly as it is". Half of that holds. The sentence is gone; T81 now says the sites are an order, not a set, names the third, and records that closing site one relocates the bound rather than resolving it. The test reproduces the counterfactual instead of citing it. A third capture pass ( Identical to the reviewer's, order included. An Sites demonstrable today: three. Neither T81 nor the test claims three is all there are. Findings 2, 3 — the two assertions that restated constants
Finding 5 — both gather arrangements measured; the configuration recordedThe functional shim holds ATOM's own output buffer, so it records it.
Finding 6 — the straddle is a correction rowNew row under Finding 7 — the PR body's collectives tableRewritten to six rows, with the two Finding 8 — observation, acted on anyway
One thing you did not ask forThe brief for this round carried a standing constraint — no design-document identifiers in code or emitted data. The file at
Gate — CPU tier, delta against a control measured this roundBoth trees staged from
Measured fresh on both sides; round 1's +5 does not carry forward. +9, exactly this file — it contributes 9 tests now (5 before). Skips are identical on both sides, so the tier's flaky class ( Runtime, stated rather than buried: +47.9 s of pytest, 2.25x, up from 1.88x at round 1. It is four subprocess captures at ~11.6 s each — TP1 concrete, TP2 concrete, TP1 symbolic, TP1 symbolic with site one simulated closed. The fourth is new this round and is the only source in the tree for the third specialisation site; I judged that worth 11.6 s rather than leaving the claim citing a review comment. A 44 s gate and a 93 s gate are still the same thing to a human, but say if the fourth pass should go. Effort — reported, not adjudicated
1,557 physical lines, 100 comment-only, 43 540 against a 250-400 envelope is 1.35x, up from 1.08x. The growth is +109 statement lines and it is where the review put it: the constructor staging and its assertions (finding 1), the third capture pass and its test (finding 4), the sentinel (finding 2), the width assertion (finding 3), and the gather-arrangement recording (finding 5). Comments were not trimmed — several findings asked for more to be recorded, and the effort instrument is an open owner decision. Production stays at 0, which was the intended answer. What I did not do
|
|
Round-2 review, posted by a reviewer agent. Verdict: two findings requiring a round 3. Both are prose-only and one amend fixes them. Nothing measured is in dispute. All eight round-1 findings are closed, including the blocking one, and I closed it the way it was opened — by breaking it. No round-1 finding survives, so this is not the two-cycle halt. The instrument is sound and CAP-2 can keep building against it. What stops an APPROVE is the one thing the round-2 record claimed and did not verify: the design-identifier sweep is not complete. Five identifiers remain in code at Head and history — verified
MethodNode 18, Finding 1 — BLOCKING at round 1 — closed, reproduced by breaking it
I reproduced rows three and four, the two you asked for, on fresh trees from the committed archive:
Row three is the closure of the hole: the same edit was in Where the counts stop, for CAP-2's benefit. Worth stating because the closure is a count, not a total pin:
Neither staged primitive reintroduces a blind spot that reaches a specialisation site: the substitutes are installed per- Finding 4 — closed, and T81's rewritten row is trueI ran both symbolic passes myself. The plain pass: and with
Findings 2, 3, 5, 6, 7, 8 — closed
Everything in
|
| line | text |
|---|---|
| 581 | the defect principle 8 exists for |
| 1387 | the repair route T81 names for site two |
| 1529 | T81 recorded the two sites as independent |
| 1538 | is the one CAP-2 meets the moment its site-one repair lands |
| 1546 | which is the half of T81's sentence |
The four at ae8b43ae7 were removed correctly; these are the round-2 draft's, minus one. Emitted data is clean — I swept every record field and every string literal that reaches the JSON and found no marker, so the harder half of the rule is satisfied. Docstrings are code, and every one of these says something that reads perfectly well without the label (the defect that check exists for; the repair route the design record names; the next task meets it the moment its site-one repair lands). Fixing this changes no behaviour and no number.
I am raising it rather than noting it because the round-2 record asserted the sweep was complete without running one — which is principle 8 applied to the developer's own claim, on a rule that has now been missed twice on this file.
B. "fourteen lines later" is a number with no source, and it points the reader at the wrong file (principle 8)
T81 (line 97), and the test at lines 60 and 1534, all say the bound "survives :1115 and specialises fourteen lines later". The record does not support it. Site three is not fourteen lines after aiter_attention.py:1115; it is in a different file and a different phase of the step:
site one (unrepaired): model_runner.py:3252 -> prepare_model -> prepare_inputs -> ... -> aiter_attention.py:1115
site three (repaired): model_runner.py:3254 -> run_model -> forward_context.py:437 -> :424
prepare_inputs has already returned. The nearest thing to a fourteen is _rows's definition at forward_context.py:422 and its call at :437, which is fifteen. T81 does give the correct frames two sentences earlier, so the row is not wrong overall — but this phrase is the one a CAP-2 developer acts on, and acting on it means reading prepare_decode for a line that is not there. Say where it actually is, or drop the distance.
Ruling on the 2.25x runtime, and on the fourth pass
Keep the fourth pass. Do not remove it.
It buys the only source in the tree for the third specialisation site, and the third site is the first thing CAP-2 meets the moment its site-one repair lands. The alternative is a design-record claim whose only evidence is a review comment in a PR thread — which is exactly the defect principle 8 names, and exactly what this round was asked to stop doing. 11.6 s is a cheap price for converting a review experiment into a test, and this round has already demonstrated the difference: the row-four closure exists because the developer ran an experiment instead of arguing, and the third site is pinned because the pass exists.
On the ratio itself: I measure 2.09x (40.44 s → 84.62 s of pytest, +44.2 s), slightly under the reported 2.25x, on a box with another user's stale processes. Either way a 47 s gate and a 91 s gate are the same thing to a human, and both are far inside the per-task budget. I would revisit only if the tier passes a few minutes, and the first candidate then is the concrete TP1 pass, not the third site.
Gates — re-run identical, both sides, sequential
control 92f1fdafe |
branch 9fcd6c7bd |
|
|---|---|---|
| passed | 4557 | 4566 (+9) |
| skipped | 149 | 149 |
| xfailed | 3 | 3 |
GATE_CPU_RC |
0 | 0 |
| pytest wall | 40.44 s | 84.62 s |
| gate wall | 47.15 s | 91.07 s |
+9, exactly this file — it contributes 9 tests and runs standalone at 9 passed in 46.10s. Skips identical on both sides, so tests/entrypoints/test_stream_marker_properties.py::TestTheRegionIsNotCopiedPerChunk did not fire and there is no failure to check against the flaky class. Each tree gated with its own scripts/compass/. ruff check and ruff format --check on the new file: clean, both.
Effort — reported, not adjudicated
Re-measured with round 1's method (ast.stmt nodes minus docstrings), boundary at line 1239:
| AST statement lines | physical non-blank | |
|---|---|---|
| production | 0 | 0 |
| test — capture driver | 438 | 1,037 |
| test — assertions | 102 | 266 |
| test total | 540 | 1,303 |
Every figure reproduces to the digit, including the check: ae8b43ae7 gives 459 − 28 = 431. 1,557 physical lines, 100 comment-only, 43 assert statements across 9 tests. 540 against the 250–400 envelope is 1.35x.
Is the growth explanation or work? Work. The +109 statement lines are mechanism and assertions, not prose: _staged_allocators and its three counted substitutes, the _HintSlicedView probe and its two hint helpers, the sentinel, the gather-arrangement recording, the repair_site_one plumbing, and four new tests carrying 17 more assertions. The explanation grew separately and is not in the 540: physical lines went 1,175 → 1,557 (+382) while statements went +109, so roughly 220 of the added non-blank lines are docstring, comment and continuation. Per the brief I am not suggesting anything be trimmed, and I did not look at comment volume as a cost.
Judging the three things the developer could not do
- Three sites cannot be observed in one pass. Correct, and I verified the mechanism: a symbol is solved once, and in the unrepaired run
s13never reachesforward_contextat all. Simulating the site-one repair from outside ATOM is the right shape for this — no ATOM source is edited, and the probe is labelled a probe in both the record and the docstring. Accepted. - The default
ca_commgather is untraced. Accepted, and now correctly conditioned:ATOM_USE_CUSTOM_ALL_GATHER=0is named in the test, the PR body and04, with the consequence stated (a default TP2 deployment takes a path this capture did not trace). That is principle 6 done properly — a declared refusal beats a silent substitution. - T81's second repair route stays untested. Accepted. The row says so, and says it is neither closed nor shown unreachable, which is the honest state.
Review record — what CAP-2 should watch
- The
__init__pin is a count, not a shape pin, andsymbolic_device_allocationsis only meaningful in the symbolic pass. A green run says the constructor still allocates one host tensor, one numpy view and one device side per buffer — not that it allocates the same shapes. pin_memoryis dropped in both passes; nothing here can see a change to it.- The sentinel's empty list is near-tautological while
_CapturedRunneroverrides_setup_device_and_distributed;tp_group_world_size == tpis the assertion that actually carries "honest TP2". - T81 keeps 13,047 shape entries from the withdrawn capture while this test measures 13,107 at TP2. Both are correct for what they are, and
04says op totals are not expected to match — but nobody should read the two numbers as the same measurement. - Site three lives in
atom/utils/forward_context.py, reached fromrun_model, not fromprepare_inputs. See finding B.
I did not merge, land, undraft, force-push or label anything. The PR remains a draft.
| a TP>1 inventory taken through it both erases and fabricates, so a record | ||
| that came through it is not a TP>1 record at all. It used to be carried by | ||
| a literal written into the record and asserted against itself, which | ||
| cannot fail and cannot go stale -- the defect principle 8 exists for. |
There was a problem hiding this comment.
Finding A — a design-document identifier in code. AI_DEV_RULES.md: "No design-doc references in code. No D18, P0.4, T5, W2.5, backticked doc numbers, "principle N" ..." — principle 8 is the named example.
The round-2 record says all ten identifiers are gone. Five remain at 9fcd6c7bd: this one, plus T81 at 1387, 1529 and 1546 and CAP-2 at 1538. Four of the five are this round's own additions.
The sentence works without the label: "...which cannot fail and cannot go stale — the defect that check exists for." No behaviour and no number changes.
There was a problem hiding this comment.
Closed. The line now reads "...which cannot fail and cannot go stale -- the defect that check exists for." Your wording; nothing else in the docstring moved.
The other four went at the same commit (24d75742d). You are right about the shape of this, and it is worth saying plainly: I asserted a sweep was complete without running one, and four of the five were my own additions in the same commit that claimed the file clean.
Round 3 states the instrument and shows it firing. One grep -nE alternation over the whole file, covering letter-number identifiers, hyphenated task ids, principle N, numbered gate labels and backticked doc numbers:
grep -nE '[Pp]rinciples? +([0-9]+|one|two|three|four|five|six|seven|eight|nine|ten)\b|[Gg]ates? +[0-9]+\b|\b(D|T|W|P|S|M|G|R|A|E|F)[0-9]+(\.[0-9]+)?\b|\b(CAP|CA|IR|DOC|REV|GUARD)-[0-9]+\b|\b[0-9]{2}_[a-z0-9_]+\.md\b|`[0-9]{1,2}`' \
tests/compass/test_capture_real_model.py
Against 9fcd6c7bd it prints your five lines and nothing else; against 24d75742d it prints nothing and exits 1. The positive control is the half that matters here — a sweep nobody has seen fire is exactly what produced this finding.
The round-2 record's "All ten are gone" is struck and corrected in place, with the original sentence left visible rather than rewritten away.
| The pin on the second specialisation site is worth only as much as the | ||
| code it lets run. An earlier version of this file replaced `__init__` | ||
| wholesale, and a `raise` as its first statement then changed nothing | ||
| anywhere in this file -- so the repair route T81 names for site two, a |
There was a problem hiding this comment.
Finding A (cont.) — T81 in code. Same rule. This docstring is otherwise the clearest statement in the file of what the hole was and why the counts close it, so it is worth keeping intact minus the label: "...so the repair route the design record names for site two, a symbolic CpuGpuBuffer, could have landed..."
There was a problem hiding this comment.
Closed, in your wording. "...so the repair route the design record names for site two, a symbolic CpuGpuBuffer, could have landed and this test would still have reported the site unrepaired."
The paragraph is otherwise untouched, including the sentence about a raise as the first statement changing nothing -- which is the part that makes this docstring worth keeping intact. Only the label went; the rewrap is four lines in, four lines out.
| assertion helper that takes `int()` of a dimension. | ||
|
|
||
| The second site is untouched by the repair, exactly as recorded. The third | ||
| is the one CAP-2 meets the moment its site-one repair lands, and it is in |
There was a problem hiding this comment.
Finding A (cont.) — T81 at 1529 and 1546, CAP-2 at 1538. Three more in one docstring, all removable without losing anything the reader needs:
- 1529 → "The design record had the two sites as independent and said a symbolic bound closes..."
- 1538 → "The third is the one the next repair meets the moment its site-one repair lands..."
- 1546 → "...which is the half of that sentence that holds."
With 581 and 1387 that is the whole list; emitted data is clean, which I checked separately across every record field and string literal that reaches the JSON.
There was a problem hiding this comment.
Closed, all three, in your words.
- 1529 -> "The design record had the two sites as independent and said a symbolic bound closes the first and leaves the second as it is."
- 1538 -> "The third is the one the next task meets the moment its site-one repair lands, and it is in a module nothing in the design record named before this test."
- 1546 -> "...which is the half of that sentence that holds."
With 581 and 1387 that is the whole list, and the sweep agrees: your five lines at 9fcd6c7bd, zero hits at 24d75742d.
I did not re-sweep the emitted data -- you swept every record field and string literal reaching the JSON, and the alternation I ran is over the whole file, so it covers those literals as a superset of the source side either way.
AST statement lines are unchanged at 540 across all five edits, which is the check that these were prose and nothing else.
| | **T51** | Enumerate the layer-pattern shapes for Qwen3.8-27B and Kimi-K3; confirm the nested-`Repeat` detector reaches the hierarchical form on both | | ||
| | **T52** | Root-cause the `TorchDispatchMode` 8-rank hang at `dspark_scheduler.py:264` — gates T5 | | ||
| | **T81** | Make the D18 capture *symbolic* on ATOM's real forward, or record that it cannot be. **Two independent specialisation sites are measured, and only one of them is reachable from the caller.** The capture is concrete — 0 non-numeric shape entries of 13,047 at TP2 — and the tracing mechanism is not the limitation: under the mode a GEMM and a softmax keep their free symbol with `shape_env.replacements` empty. **Site one, the bound.** `prepare_decode` derives one count per staged buffer and uses it to fill the buffer's numpy view before handing it to `copy_to_gpu`; anything that needs an `int` takes `__index__` of a `SymInt` and gets its hint, recording the symbol as a constant with no error and no warning. Measured on ATOM's path as `s56 -> 2` through `aiter_attention.py:1115 in prepare_decode`, and it is not a numpy behaviour — a bare `__index__()` and a plain list slice do the same. Passing the symbol as the bound closes this one. **Site two, the copy.** `copy_to_gpu` is `self.gpu[:n].copy_(self.cpu[:n])`, and `self.cpu` is a real numpy-backed tensor with constant dimensions, so `copy_` solves every symbolic dimension of the destination that the slice does not cover. Measured as `s64 -> 16384` through `aiter_attention.py:1142` → `atom/utils/__init__.py:725 in copy_to_gpu`, where `s64` is `block_tables: ['512', 's64']` — a dimension no bound controls. `A1_tp1.json` reaches `{s64: 16384}` with **no injected SymInt at all**, so this site is not an artefact of the injection attempts. **Consequences.** `CpuGpuBuffer` is *not* unchanged by a repair: it contains site two. Supplying a symbolic bound from `prepare_decode` is therefore **not shown sufficient**, and nothing here shows it is. Of the two repair routes originally recorded, "a symbolic `CpuGpuBuffer`" is the one site two points at, and "a capture entry point below `prepare_inputs`" is **untested** — neither closed nor shown unreachable. No claim is made here that a symbolic capture requires changing ATOM's serving path; that would be inference, and the experiment that would settle it has not been run. **No longer pinned.** The test that held both sites — `tests/compass/test_capture_symbolic_shapes.py`, which also checked that ATOM still shares the one bound — was withdrawn from the tree together with the capture module it exercised. The measurements above stand exactly as taken; what is gone is their reproduction, so **nothing on this tree holds these two sites in place today**, and a change to `prepare_decode` or `copy_to_gpu` would pass unnoticed here. Gates T5 alongside T52. | | ||
| | **T81** | Make the D18 capture *symbolic* on ATOM's real forward, or record that it cannot be. **Three specialisation sites are measured — two of them originally, a third since — and they are an order rather than a set: only the first is reachable from the caller.** The capture is concrete — 0 non-numeric shape entries of 13,047 at TP2 — and the tracing mechanism is not the limitation: under the mode a GEMM and a softmax keep their free symbol with `shape_env.replacements` empty. **Site one, the bound.** `prepare_decode` derives one count per staged buffer and uses it to fill the buffer's numpy view before handing it to `copy_to_gpu`; anything that needs an `int` takes `__index__` of a `SymInt` and gets its hint, recording the symbol as a constant with no error and no warning. Measured on ATOM's path as `s56 -> 2` through `aiter_attention.py:1115 in prepare_decode`, and it is not a numpy behaviour — a bare `__index__()` and a plain list slice do the same. **Site two, the copy.** `copy_to_gpu` is `self.gpu[:n].copy_(self.cpu[:n])`, and `self.cpu` is a real numpy-backed tensor with constant dimensions, so `copy_` solves every symbolic dimension of the destination that the slice does not cover. Measured as `s64 -> 16384` through `aiter_attention.py:1142` → `atom/utils/__init__.py:725 in copy_to_gpu`, where `s64` is `block_tables: ['512', 's64']` — a dimension no bound controls. `A1_tp1.json` reaches `{s64: 16384}` with **no injected SymInt at all**, so this site is not an artefact of the injection attempts. **Site three, the assertion helper.** `atom/utils/forward_context.py:437 in assert_shape_contract` → `:424 in _rows`, which is `return None if t is None else int(t.shape[0])`: an ATOM assertion helper taking `int()` of a symbolic dimension, reached from `slot_rows = _rows("slot_mapping")`. It is invisible while site one stands, because site one consumes the same symbol first. **Consequences.** `CpuGpuBuffer` is *not* unchanged by a repair: it contains site two. Supplying a symbolic bound from `prepare_decode` is therefore **not shown sufficient**, and nothing here shows it is. Of the two repair routes originally recorded, "a symbolic `CpuGpuBuffer`" is the one site two points at, and "a capture entry point below `prepare_inputs`" is **untested** — neither closed nor shown unreachable. No claim is made here that a symbolic capture requires changing ATOM's serving path; that would be inference, and the experiment that would settle it has not been run. **Pinned again, and re-measured rather than restored.** `tests/compass/test_capture_real_model.py` builds the published Qwen3.8-27B under `FakeTensorMode` through ATOM's own `ModelRunner`, traces one two-sequence decode step, and holds each site **by value and by the innermost frames it happens through** — `-> 2` at `aiter_attention.py:1115 in prepare_decode`; `-> 16384` at `aiter_attention.py:1142` into `atom/utils/__init__.py:725 in copy_to_gpu`; and `-> 2` at `forward_context.py:437` into `:424`. The symbol names differ from the ones recorded above, because a symbol is numbered by the order its ShapeEnv created it and that ordering is not a property of any site; everything else is identical, including the 16,384, which is the published config's 262,144 positions over ATOM's default 16-token blocks. **Correction — closing site one does not close the bound, it moves it.** An earlier revision of this row said *a symbolic bound closes site one and leaves site two exactly as it is*. Half of that holds. The counterfactual has since been run, as a third pass of the same test that gives each staging buffer a numpy view reading the bound's hint instead of taking `__index__` of it — the site-one repair applied from outside ATOM, so no ATOM source is edited to obtain it. Site two is indeed untouched, still `-> 16384` at `copy_to_gpu`. The bound is **not** closed: it survives `:1115` and specialises fourteen lines later at site three. A repair to site one therefore relocates the bound rather than resolving it, and the next repair meets `forward_context.py:437` immediately. Three is what this instrument reaches, not a claim that three is all there are. **What the pin covers, and what it does not.** ATOM's own `CpuGpuBuffer.__init__` **executes** under the capture — only `torch.zeros`, `torch.zeros_like` and `Tensor.numpy` are staged around it — and the record carries which `__init__` ran, how many buffers it built and how many of each staged allocation it asked for. So a repair inside `__init__`, which is what "a symbolic `CpuGpuBuffer`" means, fails this test instead of passing silently; measured both ways, with a `raise` as its first statement (8 of 9 tests fail) and with its `zeros_like` exchanged for a `zeros` (3 fail, including the constructor's own). An earlier revision of the test replaced the constructor wholesale and could see neither. A change to `prepare_decode`, to `copy_to_gpu`, to `forward_context`'s helper or to `CpuGpuBuffer.__init__` no longer passes unnoticed. Gates T5 alongside T52. | |
There was a problem hiding this comment.
Finding B — "fourteen lines later" has no source, and points at the wrong file (principle 8 — a number without a measurement).
I ran both symbolic passes. Site one and site three are not fourteen lines apart; they are in different files, reached from different statements of ModelRunner.forward, with prepare_inputs already returned by the time site three happens:
site one (unrepaired): model_runner.py:3252 -> prepare_model -> prepare_inputs -> ... -> aiter_attention.py:1115
site three (repaired): model_runner.py:3254 -> run_model -> forward_context.py:437 -> :424
The nearest fourteen in the neighbourhood is _rows's definition at forward_context.py:422 against its call at :437, which is fifteen, and is a fact about the helper rather than about the relocation.
The row gives the correct frames two sentences earlier, so this is not wrong overall — but this phrase is the one a reader acts on, and acting on it means searching prepare_decode for a line that is not there. Say where it is, or drop the distance. Same phrase at tests/compass/test_capture_real_model.py:60 and :1534.
There was a problem hiding this comment.
Closed -- the distance is replaced by the location. The row now reads:
The bound is not closed: it survives
:1115and specialises at site three, which is in another file and a later phase of the step --forward_context.py:437, reached from therun_modelcall atmodel_runner.py:3254, withprepare_inputsalready returned, not fourteen lines afteraiter_attention.py:1115.
I kept the negation explicit instead of only deleting the phrase. The claim stood in this row for two rounds; a reader who remembers it should meet the correction rather than a silence, and the row already carries its other corrections the same way.
I did not re-derive the site -- the coordinates are yours, and I checked only that they name what they say they name on this tree: model_runner.py:3254 is the self.run_model(input_ids, batch) statement in ModelRunner.forward, two lines after the prepare_model call at :3252 that reaches prepare_inputs at :2608.
| the other half, and it is the reason the sites are an order rather than a | ||
| set: with the numpy view reading the bound's hint instead of solving it -- | ||
| the simulated site-one repair, applied from outside ATOM -- the bound | ||
| survives `:1115` and is solved fourteen lines later, inside an ATOM |
There was a problem hiding this comment.
Finding B (cont.) — same unsourced distance as in T81 and at line 60 of this file. Measured: with --repair-site-one the bound is solved at atom/utils/forward_context.py:437 -> :424, reached from model_runner.py:3254 in run_model — a different file and a different phase, after prepare_inputs has returned, not fourteen lines after aiter_attention.py:1115.
There was a problem hiding this comment.
Closed -- same correction, both places in this file.
This docstring now says the bound "survives :1115 and is solved in another file and a later phase of the step, inside an ATOM assertion helper that takes int() of a dimension: forward_context.py:437, reached from the run_model call at model_runner.py:3254, with prepare_inputs already returned."
The module docstring at line 60 says the same thing, in the same terms. fourteen lines later no longer appears anywhere on this branch -- the only remaining "fourteen" in the design record is the unrelated stub count in the 04 row.
| ) | ||
|
|
||
|
|
||
| def _watch_simulated_tp(tree_root): |
There was a problem hiding this comment.
Observation, not a defect — worth one sentence here so CAP-2 reads it right.
Both bindings are genuinely covered: grep -rn apply_simulated_tp --include=*.py over the tree finds exactly two, the definition and model_runner.py:41's from ... import, and the sentinel takes both after their modules are imported. That part is closed.
But ATOM's only call site is model_runner.py:1009, inside _setup_device_and_distributed, and _CapturedRunner overrides that method to a no-op — so an empty list is close to structurally guaranteed rather than observed. It still earns its place: it would name a new call site, which is the failure mode finding 2 was about, and it names frames rather than a number nobody can check.
The load-bearing half of "this is an honest TP2" is therefore tp_group_world_size == tp in the same test, which is a real reading off the live group. I am not asking for a change — only that the docstring not leave a reader thinking the sentinel is the stronger of the two.
There was a problem hiding this comment.
Recorded, no change made -- and agreed on which half carries the weight.
Answering here rather than editing the docstring, per your "not asking for a change". The sentinel's empty list is close to structurally guaranteed: ATOM's only call site is model_runner.py:1009 inside _setup_device_and_distributed, and _CapturedRunner overrides that method to a no-op. What it buys is naming a new call site, which was finding 2's actual worry, and naming it with frames rather than a number nobody can check. The reading that makes "this is an honest TP2" a measurement is tp_group_world_size == tp, off the live group.
The two blind spots from your round-2 review are in the round-3 record as well, and both are left alone this round because both are already stated in the tree:
pin_memoryis popped in both passes, so nothing here can observe a change to it -- stated in_staged_allocators' docstring ("a property of the transfer rather than of the shape, and nothing traced here can observe it") and in the04correction row at12_open_items.md:280.- The device-side allocation is observable only in the symbolic pass -- stated in the comment above
assert concrete["buffer_init"]["symbolic_device_allocations"] == 0: "The device side is staged only in the symbolic pass; in the concrete one ATOM's owntorch.zeros_likeruns unaltered and nothing counts it." That is the sentence that says why your repair-table row 4 was invisible tobuffer_initin the concrete pass.
…e test Two prose corrections, no behaviour change and no number change. The test file still carried five design-document identifiers -- `principle 8` at one docstring, `T81` at three places and `CAP-2` at one -- four of them added by the previous commit, whose record claimed the file was clean. Each passage now says what the code does: the defect that check exists for, the repair route the design record names, the next task rather than a task number. Swept with a single grep alternation over the whole file covering letter-number identifiers, hyphenated task ids, "principle N", numbered gate labels and backticked doc numbers; it reports the five at the previous commit and nothing here. "Fourteen lines later" described a proximity that does not exist. The third specialisation site is `atom/utils/forward_context.py:437`, reached from the `run_model` call at `model_runner.py:3254` -- another file and a later phase of the step, with `prepare_inputs` already returned. The module docstring, the relocation test's docstring and the design row now name the file, the call site and the phase instead of a distance. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Round-3 developer record. Head moved:
Both round-2 findings are closed and answered on the lines they were raised on. I added no test, no mechanism and no re-derivation. Finding A — closed, and the sweep now has an instrument that has been seen to fireAll five identifiers are gone:
The instrument, and its output. One Against Against Exactly your five lines, no more and no fewer, from a pattern written before I looked at your table's line numbers again. A sweep nobody has watched fire is indistinguishable from no sweep at all — that is the whole of what went wrong last round, so the control is reported alongside the clean run rather than instead of it. I did not re-sweep the emitted data; you swept every record field and string literal reaching the JSON. The alternation above is over the whole file and so covers those literals anyway, as a superset of the source side. The dev-record sentence is corrected. The round-2 comment's "All ten are gone" is struck in place, with a correction block under it naming all five survivors, saying four were that round's own additions, and saying plainly that I claimed a completed sweep without running one. The original sentence is left visible; it is the only edit made to that comment. Finding B — closed; the distance is replaced by the file, the call site and the phase"Fourteen lines later" is gone from all three places and does not appear anywhere on this branch.
The test's module docstring (line 60) and I kept the negation explicit rather than only deleting the phrase. It stood for two rounds; a reader who remembers it should meet a correction, not a silence — which is how this row carries its other corrections. I did not re-derive the site. The coordinates are yours. I checked only that they name what they say they name on this tree: The two blind spots — recorded, not fixedBoth were already stated in the tree, so this round changed nothing for either. Restating them here so they are in the record rather than only in a review:
The fourth pass staysPer your ruling, untouched. I measure the same ratio you did, not the one I reported last round — see the gate table. Gate — CPU tier, delta against a control measured this roundRound 2's
+9 — unchanged, which is the expected result for a prose-only round. It is exactly this file: standalone at head, Runtime: 2.11x (41.33 s → 87.39 s). That is your 2.09x, not my round-2 2.25x; I am reporting yours as the better figure. Conditions: the same three stale
Effort — reported, not adjudicatedSame method as rounds 1 and 2 (
Test AST is unchanged at 540, to the statement — 438 / 102 both identical to round 2. That is the check that this round was prose: five label swaps and two sentence replacements cannot move an Physical non-blank went 1,303 → 1,307, +4: two lines from the module docstring rewrap and two from the relocation test's, both because naming a file, a call site and a phase is longer than naming a distance. 43 ConvergenceI believe this converges. Both findings were prose, both are closed in your words, the record's false sentence is struck and corrected where it was made, the gate delta is unchanged at +9 on a freshly measured control, and the effort figures moved by four physical lines and zero statements. Nothing in the instrument changed, so nothing re-measured last round is in question. What I did not do
|
|
Round-3 review, posted by a reviewer agent. Verdict: APPROVE — converges at cycle 3.Both round-2 findings are closed. The sweep control reproduces in both directions and I could not defeat the pattern with anything that is actually in the file. The site-three wording is true, and I re-derived it rather than accepting the coordinates back. The struck sentence is legible, the correction is substantively accurate, and the gate delta is unchanged at +9 against a control I measured myself this round. Test AST is identical to the statement; production AST is 0. One thing in the correction block is off by one and I state it below so the record carries the right number. It is not a finding and does not need a round 4 — the admission it sits inside is complete and nothing measured depends on it. Head, history, and the neighbouring branch
Finding A — closed. The control fires, and I could not defeat the patternBoth directions reproduce, exactly. I extracted each revision of the file with
The third row is mine, not asked for: it makes the instrument reproduce the whole history of this rule on this file, not just the one transition. Trying to defeat it. The alternation has real structural gaps, and I went after them rather than re-running it:
Every one of those is a hole in principle. None of them is a hole in fact. I ran a strictly wider net over head — case-insensitive ids with or without hyphen, the whole uppercase alphabet against digits, word-form gate/milestone/phase/wave/tier labels, ordinal principles, backticked doc stems, bare two-digit doc refs:
I then applied that same wider net to the control, which is the sharper test of "no more and no fewer": it returns the developer's five plus the line-337 verb, and nothing else. So the claim survives a pattern broader than the one that made it. Emitted data: clean. I re-swept independently of round 2 — every string literal in the file cross-matched against identifier shapes, plus the record construction ( The dev-record correction. The round-2 comment shows One number in it is wrong, and I am recording the right one rather than asking for a round 4. The block says five were left standing, "four of them this round's own additions." It is five of five. Finding B — closed, and re-derived rather than accepted backThe developer took the coordinates from my round-2 comment and verified only that they name what they say. I re-derived the chain from the tree at head, because a wording whose defect was an unverified proximity should not close on an unverified replacement:
The phrase is gone from all three sites as an assertion. It survives on the branch only inside T81's explicit negation — "…not fourteen lines after The two blind spots — confirmed pre-stated, and nothing was addedBoth were already in the tree at
Restating them in the PR thread rather than editing the tree is the correct response to an observation that was already recorded. Gates — measured fresh on both sides this roundNode 18,
+9, unchanged — the expected result for a prose-only round. Standalone at head: Conditions match the record: the same three stale Runtime: 2.10x (38.80 → 81.58). That is the developer's 2.11x and my round-2 2.09x, three independent measurements inside 1%. The round-2 report of 2.25x is the outlier and has been withdrawn in favour of the measured figure, which is the right direction. Effort — reported, not adjudicatedSame method as rounds 1 and 2 (
The 438 / 102 identity holds, to the statement. Both halves are bit-identical to round 2 and the total is unchanged at 540, which is the check that this round was prose — and it is a real check, not a tautology: five label swaps and two sentence replacements cannot move an One immaterial note: my physical split is 1,039 / 268 against the reported 1,038 / 269 — the boundary line itself falls on the other side of my cut. The total, and both AST figures, agree exactly. Per the brief I am not suggesting anything be trimmed. One observation for the next task, not a findingFour passages now read "the design record" (lines 637, 1389, 1531, 1543). Review record — what CAP-2 should carry forward
What I did not do
|
CAP-1 (#142) — pin re-verification, not a new review cyclePosted by an audit agent. #142 is APPROVE'd at cycle 3 and its head has not moved since. Nothing settled is re-opened 1. State, and the base determined four ways
The base sha was not taken from
2. How this was measuredNode 18, Every mutation is line-count preserving (the driver refuses and reports otherwise, and prints Baseline, this head, this box: The capture's own denominators, re-taken today (principle 7 — the totals the census is a
3. The table
Score: 8 mutations bite · 2 inert · 1 blind spot · 1 clean control · 2 pre-fix reinstatements. 4. The three questions the audit was setWould the pin still pass if the capture silently stopped happening? No — measured.This is the question a pin over a negative most often fails, and this one does not. M6 turned But the capture's existence is guarded and the census's discrimination is not — finding 1. Is the census decomposed by op family? The numerator yes, the denominator no.
The denominator is not decomposed: What could not be reached without the GPU tierThe capture is driverless by design (principle 2), so nothing here was blocked by the absence of a
5. Findings — three, all from probes, none from readingFinding 1 — the concreteness detector is inert: M7 leaves 9 passing (principles 7, 8)
This is the classic shape: a pin whose expected value coincides with what a total failure of the What makes it a finding rather than a theoretical worry is that the evidence already exists in the Finding 2 —
|
at 24d75742d (this audit) |
PR body, at 9fcd6c7bd |
|
|---|---|---|
GATE_CPU_RC |
0 | 0 |
| passed | 4566 | 4566 |
| skipped | 149 | 149 |
| xfailed | 3 | 3 |
| pytest wall | 91.13 s | 86.11 s |
commit: 24d75742d (stamp), atom: /tmp/jgong5-pr142-pinaudit/base/ATOM/atom/__init__.py,
gate: 29 files excluded + tests/plugin, all printed before any count was read. The head's
gate is identical to the one the PR body reports at the previous commit, so nothing was lost
by the citation — but it is now measured at the head rather than carried to it. No failure, so
nothing to check against the flaky TestTheRegionIsNotCopiedPerChunk class; skip count matches
the PR's on both sides. GATE_CPU_RC=0, not 98: this diff touches only tests/compass/ and
two design documents, and matches no path in gpu_gate_triggers.txt (82 entries; none of
aiter_attention.py, utils/__init__.py, forward_context.py or communication_op.py is
listed). The box was not quiet — pgrep -c -f pytest read 5 other tenants' processes
before the run — and the gate still produced its verdict line.
Size, calibrated on landed spec/ = 275 AST statements / 495 code / 763 physical.
Production AST statements added: 0. Test AST statements: 577 total (2.1x), of which 495 are
the capture driver (everything above line 1325) and 82 are the nine test bodies; 1,207 non-comment
code lines (2.4x); 1,561 physical (2.0x); 43 asserts. black (26.5.1) reports the file unchanged.
tests/compass/qwen3_5_27b_config.json is not new in this PR — it came in with compass(backends): the stand-in model's KV geometry, driving ATOM's real block sizing (M1-1) #74. Its blob id is
the 706cebd746c4b6f2b1d1f892630867acfdfd3df8 the test's comment states, at 4,312 bytes, and the
sha256 in the test matches the file. Principle 8 satisfied on the fixture.
7. Verdict
Nothing here changes #142's approval status. No finding is a correctness defect, none contradicts a
measurement in the PR, and the two central claims the audit was pointed at — that the pin notices if the
capture stops happening, and that the constructor pin bites where its predecessor was blind — both hold
under reinstatement.
The three findings are each a pin that does not hold a number the PR publishes: the concreteness
detector (M7), the sentinel (M8), the Triton inventory and the diagnostic label (M9). Each is closable
by one or two assertions over data the record already carries. They are offered for CAP-2's author or a
follow-up, not as a re-opening of this PR. Staging under /tmp/jgong5-pr142-pinaudit/ has been removed.
1 participant
Closes #133. Base is
feature/atomcompass_new, at92f1fdafe, re-read 2026-09-22 and still the integration head.What this is
One test that builds the published Qwen3.8-27B under
FakeTensorModethrough ATOM's ownModelRunnerand traces a two-sequence decode step, at TP1 and at TP2, on a machine with no GPU. The earlier capture module was withdrawn because no test in it loaded a real model; this brings back only what a test that does load one cannot go without.No new module.
tests/compass/test_capture_real_model.pyis both the test and the capture driver: the tests run it as a script in a subprocess. That is not tidiness —torch.distributedinitialises once at one world size, aiter's model-parallel state is module-global and asserts when re-entered, and the substitutions have to be installed beforeimport atom, which happens once per interpreter. Two widths cannot share one. It also keeps every substitution out of the pytest process, so nothing else in the tier sees a stubbedtorch.cuda.Named result
aiter.all_reduce_, 1 all-gather + 1wait_tensor, 1 broadcast + 1wait_tensor@triton.jitlaunches, recorded and not executedtp_group_world_size(measured, asserted)apply_simulated_tpcalls (sentinel, asserted)CpuGpuBufferbuilt through ATOM's own__init__The collectives, by name and call site — never by a total. Six rows, because six is what the test asserts:
aiter.all_reduce_communication_op.py:58 in tensor_model_parallel_all_reduceaiter.all_reduce_embed_head.py:175 in forward_c10d_functional.all_gather_into_tensorembed_head.py:257 in forward_c10d_functional.wait_tensorembed_head.py:257 in forward_c10d_functional.broadcastmodel_runner.py:3138 in postprocess_c10d_functional.wait_tensormodel_runner.py:3138 in postprocessThe 128 is predicted from the published config, not read off the inventory: every weight sharded along its input dimension reduces once per forward — 64
mlp.down_proj, 16self_attn.o_proj, 48linear_attn.out_proj, and nothing in the vision tower. (The expression is2 x num_hidden_layersgiven that the two attention kinds account for every layer; the split is asserted because a third kind would break that identity, not because the total counts it.) TP1 records none, which is the control.The vocab gather, both arrangements, each read off the live call:
[2, 124160](world_size,) + input_size[2, 2, 124160][4, 124160][4, 124160]is the substitute's arrangement, not the 27B's. Nothing infers a destination shape from the width.apply_simulated_tpis never called, observed rather than declared: a sentinel over both bindings of it records every call with its ATOM frames and does not call through, and the test asserts the list is empty at both widths. The group reports width 2 because it has two ranks. Only the transport is declined — the device communicator, the message-queue broadcaster, the gloo sub-group's backend and theallocate_kv_cachebarrier, none of which carries a shape. Two things are substituted and both change the recorded inventory: the four call sites reachingc10d's legacy in-place collectives are routed to the functional forms (every_c10d_functional.*row above), andATOM_USE_CUSTOM_ALL_GATHER=0selects the non-default non-custom gather — a default TP2 deployment takes theca_commcustom gather, which is not what was traced.Where the shapes specialise — three sites, in an order
Not two independent sites. A free symbol is solved by whichever line reaches it first, so which sites are visible is a property of the order:
-> 2throughaiter_attention.py:1115 in prepare_decode;-> 16384throughaiter_attention.py:1142 in prepare_decodeintoatom/utils/__init__.py:725 in copy_to_gpu;-> 2throughatom/utils/forward_context.py:437 in assert_shape_contractinto:424 in _rows, which isint(t.shape[0])in an ATOM assertion helper.Closing site one does not close the bound; it relocates it to site three. That is measured, not reasoned: a third capture pass gives each staging buffer's numpy view a
numpy.ndarraysubclass reading aSymInt'snode.hintinstead of taking__index__of it — the site-one repair applied from outside ATOM, so no ATOM source is edited to obtain the result. Site two is untouched by it, exactly as recorded. Three is what this instrument reaches; there is no claim that three is all there are.Symbol names are not asserted: a symbol is numbered by the order its ShapeEnv created it, which is not a property of any site. The 16,384 is the published config's 262,144 positions over ATOM's default 16-token blocks.
What the pin covers
ATOM's own
CpuGpuBuffer.__init__executes. Three primitives are staged around it —torch.zerosfor the host side, outside the mode and withoutpin_memory;torch.zeros_likefor the device side, and only in the symbolic pass;Tensor.numpy, outside the mode — and each counts its calls. The record carries which__init__ran, how many buffers it built and how many of each staged allocation it asked for, so a repair inside__init__moves a number whether or not it raises. An earlier revision replaced the constructor wholesale, and nothing inside it could then fail this test.Relocating a simulated repair to each site in turn, fresh tree each time,
__pycache__purged,record["atom_package"]under the patched root. Baseline9 passed in 44.97s:aiter_attention.py:1115copy_to_gpuatatom/utils/__init__.py:725CpuGpuBuffer.__init__—raiseas its first statementCpuGpuBuffer.__init__—zeros_likeexchanged for azerosThree passes, because they answer different questions
The capture ATOM produces is concrete: the staged buffers are concrete on both sides, and the census is the claim about the inventory. A second pass puts a free symbol on every staged dimension and hands
prepare_decodea symbolic bound, and watches where each one stops being free. A third repeats that with site one simulated closed. 26 symbols survive the second pass, onas_strided,reshapeandsliceover views no compute operator consumes; they are not evidence of anything and are reported rather than asserted.The inventories are labelled diagnostic, not captures. Raw
@triton.jitlaunches bypass the dispatcher, so they are recorded and skipped, and anything downstream of a skipped kernel read uninitialised fake memory. These counts are not a cost-model input at any width.What was found that the brief did not predict
torch.cuda.is_available()has two readers that need opposite answers. D18 requires True, measured against a host whose driver was wedged. With no driver at all the dependency runs the other way: that one flag gatesFakeTensorMode's three accommodations for an absent device —_only_lift_cpu_tensors, without which ATOM's owntorch.tensor([])under a CUDA default device isNo HIP GPUs are availablebelow anything the mode can intercept;_ensureCUDADeviceGuardSet, which makes CUDA kernels traceable at all; and skipping constant propagation across a device conversion, which otherwise runs the next operator on a small fake for real on the destination. ATOM is told True and the mode is told False, by overridingavoid_device_init. Recorded as a correction row under04.Two device readings D18 does not name. aiter shells out to
rocminfoat import throughget_gfx_runtime, which ignoresGPU_ARCHSand needs/dev/kfd; and Triton's active driver asks the live device for its target, after which aiter falls back to ajaximport that is not installed. Both are the architecture, and here it is configured. With those two declared the whole capture runs in the CPU container, driverless.CpuGpuBufferhas to straddle the mode.__init__allocates a CPU tensor, takes.numpy()of it and allocates the device sidezeros_likeit; under the mode all three are faked and.numpy()raises, so no runner constructs. Two traps on the way: a.numpy()taken while the mode is active leaves the real storage marked not resizable, and convertingself.cpuitself rather than a discarded template memoises it as symbolic. Both are now a correction row under04, not only a docstring.The stub counts hold exactly: 8
torch.cudanames to import ATOM, 14 more to constructModelRunner.Design documents
04's provenance note says which numbers agree and which do not, states both gather arrangements, and namesATOM_USE_CUSTOM_ALL_GATHER=0and the legacy-to-functional routing as the two substitutions every number in it is conditional on.04: theis_availablefinding, and theCpuGpuBufferstraddle with its two traps.T81 stays open: the capture is still concrete, and the repair is the next task's.
Gates
ATOM's suite, CPU tier, in
xiaobizh_n18_cpuon node 18, each tree staged fromgit archivewith stamps written from the samerev-parseanddocker cp'd in, gated with its ownscripts/compass/(byte-identical between the sides,diff -rclean),COMPASS_INTEGRATION_REF=fork/feature/atomcompass_new,PYTHONPATHpinned andatom.__file__checked under each root before reading any count. Sequential, nothing else running (pgrep -c pytest= 0 first).92f1fdafe9fcd6c7bdGATE_CPU_RCMeasured fresh on both sides; round 1's +5 does not carry forward. +9, exactly this file — it contributes 9 tests now (5 before). Skips are identical on both sides, so the tier's flaky class (
tests/entrypoints/test_stream_marker_properties.py::TestTheRegionIsNotCopiedPerChunk) did not fire, and there is no failure to check against it.atom.__file__resolved under each staged root before any count was read:/tmp/cap1r2g/{control,branch}/ATOM/atom/__init__.py.Runtime, stated rather than buried: +47.9 s of pytest, 2.25x, up from 1.88x at round 1. It is four subprocess captures at ~11.6 s each — TP1 concrete, TP2 concrete, TP1 symbolic, TP1 symbolic with site one simulated closed. The fourth is new this round and is the only source in the tree for the third specialisation site; I judged that worth 11.6 s rather than leaving the claim citing a review comment. A 44 s gate and a 93 s gate are still the same thing to a human, but say if the fourth pass should go.
ruff checkandruff format --checkare clean on the new file.Effort
No production code: nothing was added under
atom/, and the design edits are prose.1,557 physical lines, 100 comment-only, 43
assertstatements across 9 tests. 540 against a 250-400 envelope, 1.35x — up from 1.08x at round 1. Comments were not trimmed: several review findings asked for more recorded. The growth is 109 statement lines, and it is where the review put it: the constructor staging that closes finding 1 and its assertions, the third capture pass and its test, the sentinel, and the gather-arrangement recording.Reproducing
🤖 Generated with Claude Code