| Commit message (Collapse) | Author | Age |
|---|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
readback-before-scale class)
Core CUDA starts each compute-channel push with a C56F SEM_EXECUTE
acquire method (five words [ADDR_LO][ADDR_HI][PAYLOAD_LO][PAYLOAD_HI]
[EXECUTE]; op 0/2/3/4/5 = acquire/equal, strict GEQ, circular GEQ,
AND, NOR). The decoder handled only op 1 (release) and 6 (reduction),
so every wait was ignored and a push's work could run ahead of the
dependency it waits on. Observed on streams2: the D2H readback push
carries ACQ_CIRC_GEQ on the channel progress counter (0x20400ffd0 >=
4, the 4th push = the scale launch); the readback executed before the
scale, and the client verified the pre-scale buffer (bad=32767, c2 =
a2+b2 without the x2).
Now the decoder records the acquire and svc_run_ops honors it before
running any of the push's work: satisfied (read + circular/equal/GEQ
compare) -> proceed; unsatisfied -> set op_tr_pending, so the entry is
held and its release withheld (bounded by SVC_OPDEF_HOLD_MS, loud) and
the re-decode runs the work once the counter advances. Unreadable
semaphore: fail-open (pre-r336 behaviour). New stat wait=held/pass.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
The m180 unification redirected a divergent write to the ledger's page
(where an earlier write landed) even when the current translation was
pin-backed (guest pagemap truth, op_tr_src == 1). When the guest
remaps/reuses a VA (free + reallocate, or migration), the ledger page is
freed and can be reallocated by the guest kernel: the redirected write
then smears client copy data over kernel memory. Observed on the r335
vector_add run: an RM map/AVL node's pointer field held client c-data
(floats 551.25/552.0), the free-path traversal faulted
(mapFindLEQ_IMPL/_mapInsertBase), and the guest panicked
(Attempted to kill init). 4448 write-side divergences in that run.
Now a pin-backed translation wins: the write goes to the fresh page and
the ledger store after the write re-homes the entry, so readers converge
on the same page. Non-pin translations (2..5) keep the m180 redirect -
they can carry stale PTEs (m169) and the write history stays the
authority there. New counter sema-svc ledger pinwins.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Under MSGQ_FLAGS_SWAP_RX (both sides set it; flags=1 in both headers) the
guest driver's msgqRxMarkConsumed writes its STATUS readPtr to
pOurRxHdr->readPtr, and the driver's pOurRxHdr is the COMMAND-queue buffer's
RX header (msgq.c msgqRxLink swap branch: pReadOutgoing = &pOurRxHdr->readPtr
with pOurRxHdr from msgqTxCreate's pBackingStore = pCommandQueue). The
status queue's own +rxHdrOff is the mirror we publish - our command-queue
read position.
The status free-space check read the mirror, so it tracked our own writes:
with the guest mid-drain the ring was declared full (free=0, peer=0,
write=62/63), replies were written into it anyway, unread records were
overwritten, and the guest RM read a newer record - GspMsgQueueReceiveStatus:
Bad sequence number. Expected 447 got 510. Possible memory corruption. -
after which every RPC failed (status 0x25) and the first client's cudaMalloc
returned device busy or unavailable. That is the boot lottery: on a lost
boot the first client (mail) FAILs and the boot is poisoned for every later
client.
Evidence (r330 vs r332): r330 mail FAILed at the first cudaMalloc and only
one channel unregistered in the whole boot (the guest RM never tore its
channels down after the status desync); with the fix r332 mail + matmul pass
and every client's channels unregister (79 unregisters, one session flush
per client).
Also adds bounded diagnostics: sema-d2hmap (per-page destination resolution
of >=1MB host-dest copies) and a raised DIVERGE witness budget.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(cold-run gap)
The fwd-defer arm took priority over svc_opdef_release, so a pass whose op
was still pending recorded its fence release in svc_fwd_rel (applied at the
hold break / fallback) instead of the opdef slot - the fence fired while
the copy was still owed, the driver refilled the ping-pong staging, and the
late capture/read got refilled bytes (cold-run vector_add gap: one to four
chunks per cold boot). svc_fwd_defer_arm() is now skipped when
svc.op_tr_pending, so those releases go through svc_opdef_release and apply
only when the retried pass has no pending op (the op has run by then).
Measured cold-run vector_add bad across this iteration: 193536 -> 100352
(r323) -> 479232 (r324 wrong queue-pending experiment, reverted) -> 66560
(r325, ~1 chunk). All other acceptance items green on r325 (mail OK, smi OK,
matmul OK, transformer OK).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
denied passes
Two fixes for the cold-boot classes (CONT.60):
1) Pending releases starved forever: svc_apply_pending resolved the pool
semaphore with h=0, but the verified pool resolver is channel-keyed - it
never matched the per-channel pool entries while the very same lookup at
decode time succeeded (matmul r322: 'pool-rel cand ... rrc=0 cur=0' then
'pending release starved va=0x121000010 payload=1 tries=4/64...' forever,
the client hung and the boot died). Pending entries now record the channel
they were decoded for (struct svc_pending_release.h, set from svc.cur_ch
which svc_decode_stream stamps at its top) and svc_apply_pending resolves
with it. Also applies to the 'completion hint rc=-2' class.
2) The fence escaped on denied passes: in the P6b wave-2 overflow branch
the release was queued only when n_ops>0; a denied pass (n_ops==0) but
op-pending (a copy still owed) applied the release inline, letting the
driver refill the ping-pong staging before the copy ran - the cold-run
capture gap (vector_add run 1 bad=157696/193536). The queue condition is
now n_ops>0 || svc.op_tr_pending.
Measured (r323 fresh boot, full sequence): mail OK, nvidia-smi OK, matmul
OK (no death - the r322 killer), transformer OK (no stall), vector_add
cold-run bad=100352 (two chunks, down from 157696-193536; warm runs remain
bad=0).
|
| |
|
|
| |
(wd_rewind_off_until struct member repair)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
first decode
vector_add was deterministic-wrong (bad=786335 checksum=359549561 identical
across boots) because the 1MB H2D chunk copies always read the driver's
ping-pong staging AFTER it had been refilled (the fence release applied long
before the late op execution; 60+ rc=-2 retry storms per chunk).
Two wiring fixes (r314-r315):
1) The decode-time snapshot was NEVER CONSUMED: svc_copy looked it up
(svc_copy_snap_lookup at the top) and freed it on success, but no code read
from it - every copy fell through to the live per-page source resolution
and read refilled staging. The source read loop now takes the bytes from
the snapshot when present (memcpy from snap + o). The copy also no longer
needs a live source translation, so it executes on its first allowed pass
instead of storming retries.
2) Capture on the DENIED pass too: the old capture gate lived in the
op-build block (allowed passes only - by then the fence fired and staging
was refilled). The denied device-dst copy branch now captures immediately
(pipeline-order moment), same translation-probe guard.
Measured r315 (fresh boot): vector_add first run bad=157696 (cold-run
capture coverage gap), second run bad=0 checksum=392344090.5 STATUS OK
with ALL EIGHT chunks clean and no retry storms (~75 s, guest survives).
Previously 786335 bad / identical checksum every boot.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
for the sw-copy client
Two fixes, both demanded by live evidence:
1) Stale-wedge republish loop (r311-r312): the pump's per-channel write-back
did not persist gr_pub_armed, so the residual watchdog re-armed every pass;
its recovery kept zeroing the completion word and re-publishing (pub-adv/
zero/stale-wedge 3s cycles, exit: pub-adv pa=... val=104 cur=0 last=0). The
word oscillated under a live client that then never completed (r311 vector_add
wedged; guest sshd died). Persist gr_pub_armed/counter_gen in the write-back
and keep them (plus slot_pend_tries) across same-generation re-registration.
Also guard the residue zero: never zero a completion word authored by the
live session generation (r312 instrumentation: pub-adv-wrc + stale-recover-
zero + skip witness). Verified r313: stale-wedge rows 0 (was 66+), client
completes, guest survives.
2) Staging snapshot capture for the sw-copy client (r311): the decode-time
H2D capture gate required src_aperture == CE_APERTURE_SYS, but the driver
sends no SET_SRC_PHYS_MODE for the 1MB chunk copies, so src_ap defaulted to
VID and the gate never matched - zero captures on every run r302-r310 while
executions read the (driver-refilled) staging late. Gate judge is now the
translation: capture when the first source page resolves through SYS. First
verified captures (r311: 2-3 per run, real client bytes in first4) - partial
coverage only; the remaining pipeline-order gap is the open lead (see
handoff CONT.57).
Diagnostics kept: snapgate witness, pub-adv-wrc, stale-recover-zero[-skip].
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Root cause of the vector_add chunk drop (CONT.55 chase): the sw-copy
client's GR-classified GRCE channel (eng=0, h=0x5c000059) decodes its
1MB H2D chunk copies as CE transfers, but ops on that channel run only
once the pin/trust stream (injected) reaches each push page. The first
pass over a chunk push lands before that; with allow_ops=0 the copy was
never built and the entry was consumed silently. All later passes
classify as replays (slot_ops_complete=1), so the entry was never
retried: 6 of 8 chunks skipped, kernel computed on stale VRAM
(vector_add r302-r307: 2-5 of 8 chunk ops ever existed, the rest never
ran; the +2MB 'phase' signature is stale residue).
Fix: a denied launch that WOULD have built a type=1 device-dst copy
(no remap, both offsets, out >= SVC_VAS_FLOOR) with a trusted push now
marks the pass op-pending. op_tr_pending un-completes the slot, so the
next pass is treated fresh and the copy runs as soon as trust (injected
with allow=1) or the operand translation arrives - the existing
m163/opdef hold+retry machinery. Replay passes are deliberately
eligible (that is what makes the retry possible); a per-slot
slot_pend_tries cap (4 per put epoch) prevents a slot whose trust never
arrives from wedging the channel in 10s hold cycles.
Measured (r308, sema_dbg): all 8 consecutive chunk dsts build ops
(0x...cc..d3) and retry to success - 28 CEAUDIT rc=0 soft copies, 38
rc=0 vs 20 rc=-2 executions; previously max 5 dsts and only the last
2-3 ran. Regression checks on r309 with the fix: probe RAWA2+PROBE2B
bad=0, victory_mail OK, victory_matmul OK (one transient first-run
fail, clean re-run), transformer OK. vector_add now advances through
the full H2D set but still does not complete (retry storms stretch the
copy phase; the run later wedges in the GR stale-wedge recovery loop -
named as the next lead).
Diagnostics kept (dbg-gated, bounded): sema-pend set/hold witness,
sema-consume consumed-slot witness, sema-entry de-gated for eng=0
channels, dmatrace extended with replay/eng/pend_would, budgets raised
under sema_dbg so boot init cannot spend the trace before the run.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
ops-order release guard
1. CONT.53 step (1): device-dst copies from host sources now snapshot the
source at op-build (bounded pool 16 x 4 MB, 30 s TTL) and execute
from the snapshot, so the staging read happens at the push's
processing position.
2. The capture gate is aperture-based (dst VID, src SYS) after the r300
captures showed a naive VA-floor gate also caught D2H readbacks and
froze pre-kernel zeros.
3. Ops-order release guard: a pass that built ops queues its releases
through the pending pass instead of applying inline, so a fence
never clears before the pass's copy ops read their source.
A/B (r297/r299/r300): snapshots fire (6-7 per run) but the +2 MB phase
persists (win@0/1M = exp[+2M]); the r300 capture contents show the H2D
op sources already carry late-chunk data at decode - the scrambling is
at the driver's staging hand-off into the push stream, before our
service sees the op (next instrument: correlate the client's PTRS with
per-capture buffer identity).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
The m169 device-dst guard defers a staged H2D up to 3 s; the driver's
ping-pong staging buffers are overwritten meanwhile, so a deferred read
returned a later chunk's content (the +2 MB vector_add phase = the 2x1
MB staging stride). Held copies now snapshot their source at hold time
(bounded: 8 slots x 4 MB, 30 s TTL) and the retried execution copies
from the snapshot, matching the hardware's in-order staging read.
r296 A/B: the snapshot never fires for vector_add - its H2D dsts
resolve with pins promptly and execute without the hold; the +2 MB
persists identically. So the staging overrun happens WITHOUT the
guard: the driver's 2-buffer staging discipline relies on the copy's
completion fence (a semaphore applied when the copy executes) to pace
the restaging; if our completion for chunk i arrives before/independent
of the staging read, the driver restages and our later decode reads
chunk i+2. Next lead: the staging fence's increment timing vs the
copy's execution (the release-before-work class extended to the H2D
staging fences), not the m169 hold.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
both-side op advance
Survival (the death cascade is gone for the full user sequence):
1. pool residue re-seed applies ABSOLUTE (cur=0) - the UVM assert
(uvm_channel.c:204) that took the device down no longer fires
(committed earlier in 67148bd8).
2. completion-hint window 10s->1.5s (67148bd8).
3. status-ring free-space discipline in gsp_rm_send_flags: the
guest's consumption cursor (msgqRxHeader.readPtr at stat+rxHdrOff,
the msgqTxGetFreeSpace source) gates EVENT posts; events drop
freely when short, replies wait (bounded) and proceed. A guard
rejects a cursor the guest cannot be behind (peer > our writePtr -
the observed mirror tracks writePtr+1 and gating on it silenced
every wake-up and wedged the clients).
Verified live: mail OK, nvidia-smi OK, victory_matmul OK,
vector_add completes (content still wrong) and victory_transformer
STATUS OK AFTER it - no ISR panic, no guest death, guest ssh alive.
Content: CE linear auto-advance now applied to BOTH offset registers
after each executed op (the driver re-sets one side per op and the
engine advances the other; H2D ops showed an explicit advancing dst
with a fixed src base: the device a held h_a[0..1MB) replicated).
vector_add content still FAIL (~66% bad) - the H2D operand stream
under retry/replay passes is the next lead (opadd rows show five dst
x two src combos with heavy retry multiplicity).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
auto-advance tried and reverted
1. svc_gr_cand is now a 4-slot ring addressed by an accessor macro
(svc_gr_cand -> ring[pos]); captures that would overwrite a
DIFFERENT launch's unconsumed candidate move to a free slot, and
svc_gr_request selects this channel's oldest unconsumed candidate
before consuming; the defer-arm gate checks any unconsumed
candidate. Struct tag renamed svc_gr_cand_rec to keep the macro
safe. (The streams2 cross survives this - it is not (only) the
single slot; see the new evidence in the handoff.)
2. Diagnostic budgets raised for the triage (opadd/op_trace 2048,
memsets rc-log 2048). These exposed the real memset shape: the
8 MB cudaMemset is SIXTEEN grid=1024/block=512 kernel launches
whose win0 param bank is fully correct (dst patched to the reloc,
w0x168=0x5a5a5a5a, launches OK) - yet the D2H readback returns
zeros, so the loss is between the engine's stores and the
readback pages (the rec hpa/landing class, GRVAUDIT unlanded
mass), not the params.
3. A CE-style per-op dst auto-advance was implemented, evaluated
(no memset improvement) and reverted - re-sets are decoded and
the risk to the green clients was not justified.
Verified on r262 (boot 1 of the acceptance round): ctaid_probe RAWA
0 1 2 3 4 5 6 7 + PROBE2 bad=0; transformer STATUS OK; streams2
bad=65536 (red); memcpy_roundtrip wedged the guest (flaky: green on
r254, wedges on r260/r262).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
reloc half
Three fixes on top of 54cf842c, targeting the big-buffer clients:
1. m163 parity for fills: a small failing FILL op with a virtual device
destination is now held/retried like the small copies (-ENOENT with
the pins still landing was silently counted nomap and the release
applied - the 8 MB cudaMemset landed only partially). Verified:
memcpy_roundtrip now reaches STATUS OK (RT1 bad=0 csum exact,
MSET bad=0, h_m all 0x5A).
2. svc_copy landing ledger parity with svc_fill: a large virtual-dst
copy records its landing pages so the GR rec build (m171 ledger
preference) follows the page the copy actually wrote. streams2's
c1 moved from a stale-page mix ([28.0, 28.625, 29.25]) to the
correct-page values ([0.0, 0.125, 0.25]).
3. Per-launch relocated-data base: concurrent forwarded launches
(streams2's vadd pair) shared one base and their bounce windows
overlapped; the base now alternates across the grown region's two
halves keyed to the capture push (retries keep their half), cap
halved accordingly. Also raise the per-page data attribution budget
(16 -> 192 pages) for span forensics.
Remaining on streams2/victory_multikernel: the single-slot svc_gr_cand
capture is overwritten by the second concurrent launch before the work
item consumes it - ctaid-probe evidence: c1 == a2+b2 exactly (the
stream-2 result in the stream-1 output) with the in-tree
candidate-overwrite witness present. Next: make the capture a small
FIFO drained in order by the gsp-phys-work item.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
green
Two fixes on top of c66a0172:
1. Slot rule for LOAD_INLINE_DATA runs: in the UVM upload pushes
each run fills a whole 32 KB slot - its data runs to the next
run header (or to the next 0x8000 boundary for the last run),
regardless of the count field (observed slot starts every
0x8000: cnt=4095 and cnt=2497/2753 variants). Compact pushes
keep the count*4 reading. The old reading truncated each slot
run to half its data and walked the tail as methods.
2. Argument-ring order: the bridge's arg-block search must see the
run's START as the inline-write destination; the new cursor
advance must not leak into rec->dst_va (an end-of-run dst_va
made the QMD param patch miss and the engine never consumed the
launch - regression found and bisected against the pre-session
ko via victory_mail).
Verified one boot (r251/r252, standard params): ctaid_probe RAWA
0 1 2 3 4 5 6 7 + last8 16376..16383, PROBE1/PROBE2/PROBE2B bad=0;
transformer STATUS OK; victory_mail STATUS OK.
Residual (distinct class, NOT the H2D delivery): streams2
bad=65534, memcpy_roundtrip bad=8388608, victory_multikernel
bad=2883066. Their big spans collapse to a single rec/bank entry in
the launch tables ('bank 0 hit @0x168: VA -> reloc' for one page of
the span) and the engine write-back replicates one page across the
span (OUT-ALIAS-WIT partners) - the original span-aliasing class on
the large-buffer paths.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
instruments
The probe/mallocecho H2D (cudaMemcpy H2D of a 64 KB device buffer)
arrives as LOAD_INLINE_DATA method runs (op=3, subch=1, addr 0x1b4)
carrying the payload inline in the push, split across several runs
(and, depending on pool state, across several pushes of one sequence).
The inline-write handler wrote every run at st.offset_out without
advancing, so a multi-run upload kept only the last run's words and
the readback page held deterministic junk (RAWA 16382 16383 8193...).
Fix: a LOAD_INLINE_DATA run advances the destination cursor by the
bytes it wrote (hardware FIFO semantics); consecutive runs of one
upload now tile from the first destination.
Verified (fresh boot, standard params): mallocecho d-afterH2D head
0 1 2 3 4 5 6 7 (was junk); ctaid_probe RAWA first8 0 1 2 3 4 5 6 7
(was 16382 16383 8193 8194...). Residual: the upload spans multiple
runs/pushes and only the first 4095-word run lands today (bad=12289,
tail zeros) - the remaining runs are still lost to the decoder's
run-boundary handling (see sema-next-hdr captures).
Instruments added (all gsp_sema_dbg-gated, bounded, session-reset):
sema-pushall (de-windowed push decode trace with page_pa),
sema-opadd, sema-copyent, sema-dmatrace (denied transfers),
sema-bigrun (count>=1024 runs on len>=8192 pushes with cursor state),
sema-lid-hdr, sema-skiprun, sema-mystery, sema-next-hdr (header after
each big LID run). LDA/xfer-raw/stream dumps de-gated from
ch->h_object!=0 so unidentified channels are visible.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
class is ~950 recs)
The probe's launches carry ~950-970 recs (its 2 MB VA arena is one span),
so the <=64-rec gates on A-TRACE pt=P3 (engine-read source content) and
READ-DIVERGE-WIT (fill vs copy-path vs ledger per rec) never fired for
the probe - exactly the launches they were built for. Raise both gates to
1024; budgets unchanged (128 lines / first 8 recs), so the transformer's
1418-rec launches stay out.
r239 findings this rides on: the probe's H2D (h_A->d_A) leaves ZERO host
trace (no copy, no launch, no audit in a continuous 6.6k-line dmesg
stream) and A's vidmem reads as deterministic junk from the very first
post-upload read; the driver's 512-thread memset kernel zero-writes with
fe_esr=0x80000000 and the 3x fresh-attempt retry path (fresh qmdwin bank +
win0 bank + rebuilt PB) does NOT recover it. A minimal new repro (blk512)
shows the band bug without any probe machinery: grid(8,1)x512 writes 4/4
buffers perfectly while grid(16,4)x(16,16) 'out[i]=i' is deterministically
bad=12288 (rows >= 16), with the engine-side scan reporting 31/31 changed
pages - the by>=1 stores exist somewhere the client does not read.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r237-2: T2's GR channel carried T1's residual put=115; my 3 s empty-decode
hold delayed the host's GP_GET during the residual walk, and the guest
unwound at binding+3.46 s (zeroed PROMOTE_CTX, then full FREE) - exactly
when the host finally consumed the residual. The guest's ctx-init
patience is ~3 s, so the hold must never delay the walk that long.
- New SVC_GR_NEWBIND_HOLD_MS 1200u and per-channel
wd_newbind_hold_until: the hold windows (empty-decode and unmapped
entry) now apply only while it is fresh (<1.2 s; the rebuild lands in
~0.8 s) and only when put == get (residual context, no producer work
pending). Any pending producer work is walked without delay, so
GP_GET always advances promptly.
- wd_newbind_until (3 s) still governs retry eligibility; the opdef
drop, chanswitch log, pending-map/rebuild path are unchanged.
make kernel green.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r235: only 4 of the probe's ~9 64KB copies reached the ring (the
sema_bigcopy_log budget had absorbed the rest - the A H2D is missing
entirely), and the events lack direction/buffer identity. Now:
- every total==65536 copy bypasses the bigcopy budget (its own 128-line
A-TRACE budget still bounds volume) so the probe's H2D and all D2H
readbacks always appear;
- A-TRACE carries dir=H2D/D2H (device=VID low-pa, host=SYS high-pa) plus
pt=P1/P2 (upload read/landed) or P4/P4R (readback source/result).
r235 decode so far: the A upload lands ([0 1 2 3] in device pa 0x7800000
at t=.94) and the SAME device page is later rewritten to
[16382 16383 8193 8194] - h_A's own values at shifted offsets - with the
pattern repeated across consecutive device pages (0x7800000/0x7810000/
0x7820000), i.e. a repeated-page/offset-wrapped write inside the emulated
copy or engine path, not client-side. The trace fix makes the next capture
name the exact path (missing H2D included).
(cherry picked from commit 9b35145136c949405c7fb4d22b3405c174ecb055)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r235-2 mechanism: at R2's ctx-init one pass fired 8 'stale-replay
rewind=1' lines (waited_ms=3491, incl. fresh channels with sub=6
put=1) and set each channel's pin-batch rewind cooldown to +24 s. The
rewind is the retry path for a fresh context (its init pushes are
decoded before the pins land; the rewind re-decodes them once the
translations exist), so disabling it killed R2 exactly like it would
every next client (r233 P2+, r235 R2-R4). R1 (working) fired zero
tokens. The 3 s window misread a normal init oscillation (rewind
replaying a just-consumed entry) as a rewind livelock.
- Removed the rewind cooldown entirely: wd_rewind_off_until field,
write-back, watchdog clear/set and the svc_handle_map_add skip are
gone. The pin-batch rewind is never manipulated by any recovery.
- stale-replay now requires SVC_GR_STALE_REPLAY_MS (10 s) of stalled
position with consumption happening (a genuine livelock; r225 wedge
was 42 s+), keeps its window accumulating until then, and clamps
last_get = put at most once per put epoch (wd_epoch_done).
- Token updated: '... reason=stale-replay reset=0 clamp=1 ...' (no
rewind manipulation; clamp is bookkeeping, a rewind may undo it).
Epoch legs (residual/dead-range), stale-entry, newbind window,
pending-map/rebuild path, kill-recovery, r220 guard and the m181 rule
are unchanged. make kernel green.
(cherry picked from commit 459384e02bba4eca04fa3eca2365a731a4a7a33a)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r233: ctaid_probe attempt-1 completed (newbind fix works) but the NEXT
client (P2, fresh process after P1's CLEAN teardown+flush) stalled with
zero tokens. witness-r233.log shows why: after the last FREE
(180966.07) and rebuild #40 (180967.07) the guest sends ZERO traffic
for 9+ s (only 'completion hint rc=-2'); no channel binding ever exists,
so no entry-level recovery (pending-map/chanswitch/stale-entry) can
apply. Pattern across boots: a CLEAN client exit is followed by a
pre-registration stall; a KILLED client is followed by a working client.
Diagnostics added (bounded 16/session each):
- 'GR-STALL-newbind h=... put=... get=... action=window
reason=stale-entry-newbind' when a binding's freshness window is
armed (proves the reconcile runs per binding, not just once).
- 'GR-STALL-rebuild consumed=1 by=svc_thread ...' when the service loop
consumes svc.rebuild_now (proves the forced rebuild is taken; rules
out the stale-flag hypothesis).
wd_chanswitch/wd_newbind_until are per-channel/per-binding (armed on
every !same_gen registration) - no global-state stomping between
clients. make kernel green.
(cherry picked from commit 2e6921f4fe7d2a0dcd478430981367a6e354465c)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r233: the probe ran (attempt-1 completes now) with the same deterministic
foreign readback, while REC-DUMP/READ-DIVERGE-WIT/READ-ORDER-WIT/ledger
DIVERGE were all absent from the capture. The r231/r233 rings only cover
the transformer_mt window (6-7 GRID-WITs, 15 sema-bigcopy lines, all of
the transformer); the probe's launches are outside the saved ring, so the
witness absence is not yet attributable to the code. Two additions make
the next capture self-contained:
- RDW nrecs=<n> launch_va=<push>: one arming line per forwarded launch
(budget 256) - a capture can never confuse 'witness did not run' with
'no recs'; probe-class launches must show nrecs>0.
- A-TRACE, four labeled points for every 64 KB copy and every small
launch (budgets 128/128):
pt=P1 H2D source read (src_int4 as ints; expect 0 1 2 3)
pt=P2 H2D destination right after the copy (dst_int4_after)
pt=P3 engine-read point: for <=64-rec launches, the first recs'
source-page content the copy-in stages (w0..15 per rec)
pt=P4 D2H source read (the engine-authored page as ints)
pt=P4R D2H host-side readback after the copy
The probe A buffer is ints 0..16383; the first point where the pattern
breaks is the bug site.
Stats check for the elision class: r231's probe-window stats show
copy_skipped=0 copy_soft=0 copy_fail=0 and no device-dst guard lines, so
no forwarded/skip elision is in play in that window; the trace decides
where the read/write split actually occurs.
Build green.
(cherry picked from commit 98bfc75410b214f7de3dfc6ef80bff76118ff14d)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
vs ledger)
r231/r232: the probe-class launches produced zero REC-DUMP lines while the
transformer's 8+2 appeared - the per-line pr_info_ratelimited was
suppressed wholesale during the probe's post-stall logging burst. REC-DUMP
now uses plain pr_info under a boot-wide 512-line budget, so a <=64-rec
launch always dumps its full rec table.
READ-DIVERGE-WIT (sema service, rec-build loop): for small launches the
three resolutions of each VA are compared side by side - the fill's
dg->pa (the page whose content the copy-in stages; exactly what the engine
reads), svc_op_translate (the H2D/D2H copy path's resolver) and the
ledger's write landing. A page-level split is logged with op_src/op_rc and
an op-miss token. Bounded: first 8 recs, boot budget 128.
Static permutation check (negative): reloc_va is assigned sequentially in
fill order (svc_gr_data_fill), copy-in/out iterate recs in order, the
engine reads reloc-base + linear offset, and the per-page copy loop folds
in-page offsets correctly on every resolver path; no sort/rotate/LIFO
order found in the page-table or copy paths. The remaining split suspect is
the m171/m180 ledger DIVERGE preference - exactly what READ-DIVERGE-WIT
compares.
Build green.
(cherry picked from commit 59e50207d60067b5d1e3a0546f20e9b22c92e112)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Net port of the stallfree work onto master without the reverted hint-stop:
- Channel-reuse reconciliation, provably safe: svc_chan_registration marks
a genuinely new binding (!same_gen) with wd_chanswitch; the first pump
pass drops opdef slots bound to the reused handle ONLY when their
first-pending stamp predates the reconcile instant (time_before), so a
slot the fresh binding creates (after registration, at/after the
reconcile) can never be touched; logs reason=stale-entry-chanswitch when
the ring already carries content. Captures stay enabled (the fresh
client's own entries may already be at slot 0).
- First-attempt readiness (r229/r231: attempt 1 stalls, attempt 2 works
once the 10 s rebuild cadence caught up): a genuinely new binding gets a
3 s freshness window (wd_newbind_until, SVC_GR_NEWBIND_MS) in which the
pump (a) does NOT silently consume a GR entry that decoded nothing (no
release/QMD/ops, non-replay) - it holds it and requests an immediate
rebuild (svc.rebuild_now, consumed at the top of the service loop before
the next pump), and (b) does the same for a GR entry whose page is
unmapped. The entry re-evaluates with resolved pages inside the window,
so the launch executes on the FIRST attempt; after the window entries
consume as before (<=3 s extra on first contact).
- Token: 'vgsp-gr: GR-STALL-pending-map h=... va=... entry=... action=rebuild
reason=stale-entry-newbind' (bounded 16/session).
- No hint-stop (master's revert already removed it; kept absent - the
re-post window is the delivery bridge for late-registered events).
- f0789007 kill-recovery, the r220 hold guard and the m181 arm rule are
untouched (already present on this base).
make kernel green (0 errors, pre-existing warnings only).
(cherry picked from commit 4d420373e50be777e967b0c671a8ff037ef92a45)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r231 probe: RAWA (fresh 64 KB upload -> straight D2H readback, no kernel)
returned foreign deterministic content (0x437F0000-family ramps) where the
uploaded ints 0..16383 must be, and PROBE2 == PROBE2B on that same content.
The upload/readback round trip inside the emulation is therefore not
content-preserving - before any engine execution. Two candidate mechanisms,
both now witnessed:
- READ-ORDER: a device-src copy reading a VA whose device-dst write is
still inside the m169 pin-guard window (the writer's bytes land only when
the guard resolves) - the read executes first. READ-ORDER-WIT logs the
coincidence; with sema_read_order_guard=1 (new param, default 0) the read
is deferred through the existing bounded-hold re-decode so it runs after
the write. Opt-in pending an A/B proving no green regression.
- RESOLUTION DIVERGENCE: the dst (H2D) resolution consults the ledger
before svc_op_translate while the src (D2H) resolution does the opposite,
so a stale ledger landing can make the writer and reader touch different
pages. The sema-bigcopy proof (va->pa/aperture/provenance for BOTH sides)
now fires from 32 KB (was 1 MB), so the probe-class copies are on record.
Build green; module param registered and visible.
(cherry picked from commit bfe5d295dd3922c0dbcd0bd1a29912a02acb58ec)
|
| |
|
|
|
|
| |
(GR-STALL-stale-entry-chanswitch)"
This reverts commit 21eb5bb0582aaa2409601a1a1466f7bc72841647.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(GR-STALL-stale-entry-chanswitch)
r229a evidence: victory_mail OK -> ctaid_probe attempt 1 printed DEVICE and
never completed (killed at 50 s); attempt 2 (fresh context) worked. The
retained window shows mail's teardown + session flush (gen 3) then, for
the whole first-attempt window, no channel registration and only the host
'completion hint rc=-2' flood (~0.5 s cadence, window reopened by the
completion-apply sites). No GR lines, no stale-* lines: the first launch
never reached the sema walk for the new binding.
- svc_thread hint: on rc == -2 (no event context to post to) stop
reopening svc.post_hint_until; the next real completion reopens it. The
flood was pure noise and obscured the real state.
- Channel-reuse reconciliation: svc_chan_registration marks a genuinely
new binding (!same_gen) with wd_chanswitch; the channel's first pump
pass drops every opdef slot bound to that handle (a dead context's
withheld sets must not ride the reused channel - the r228 opdef-hold
class) and, when the ring already carries content (put > 1), logs
'GR-STALL-stale-wedge ... reason=stale-entry-chanswitch'. Captures stay
enabled: the fresh client's own entries may already be at slot 0, and
suppressing them would lose exactly the first launch this protects; a
stale launch re-walk stays bounded by gsp_phys_gr_wait_max_ms and the
m181 put-drop recovery. r220 hold guard and the m181 arm rule
untouched.
Note: the phase2 witness for r229 was not readable at commit time
(/tmp/opencode/witness-r229-phase2.log absent); r229a is boot-side and
post-flush, so the attempt-1 window shows the guest sent no RM/channel
traffic for ~9.5 s - the reconciliation removes the host-side reused-
handle state classes; r230 will show whether the first-attempt launch
now reaches the walk. make kernel green.
(cherry picked from commit 21eb5bb0582aaa2409601a1a1466f7bc72841647)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r228 witness: entry 82 of channel 0x5c000059 stayed pending with
put=84 get=82 last=82 frozen for ~10 s (its 1 MiB op operand never
resolved, sema-optrace rc=-2) while 'sema-opdef: hold entry=82' and
'opdef-clear caller=4 entry=82' repeated: svc_opdef_drop_stale cleared
the opdef slot on every rewind replay that processed a different entry
of the channel, so the 10 s deadline restarted and the hold re-armed
forever (opdef_held climbing, no EXPIRED). No GR-STALL line fired
because this is below the GR-publish layer: the launch never reached a
capture/submit.
- svc_opdef_drop_stale: only a slot whose entry was CONSUMED this put
epoch (slot_done_epoch bit set by svc_replay_note) is delivered-late
and dropped. A rewind re-decoding an unconsumed entry no longer resets
its hold deadline, so the bounded 10 s flush/expiry is reachable
again.
- svc_gr_stale_wedge_check: new stale-entry leg (reason=stale-entry,
entry=1) for a genuinely stalled pending entry: GR channel,
put > get, position AND submissions static for the 3 s window, and no
bounded owner active (no defer_since grace, no no-capture window, no
live m150 hold, no opdef slot for the stuck entry). Action is the
canonical unstuck: last_get = 0 + wd_nocap_put, decode releases only
(forward-only, no re-capture/re-forward), once per put epoch. No
residual publish/arm (m181-only rule); r220 hold guard untouched.
Green safety: a green pending entry either advances submissions (consume
within ms) or sets one of the bounded owners the leg gates on; the
stale-entry leg therefore needs all of {put > get, 3 s position stall,
0 consumption, no owner} to fire. make kernel green.
(cherry picked from commit f0789007c49c7c6e01da5643d56ceff0b2b83819)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r227 field result: the watchdog fired 10x (reason=stale-epoch, arm+publish)
but the wedge survived, and victory_mail after the stalled rerun also hung.
Root causes fixed here:
1. Do NOT arm/publish on a residual epoch. Arming let the next pass's
svc_counter_publish write the residual put (35/101) into the fresh
client's completion word - the r207/mm2 far-jump wedge
(UVM_GPU_SEMAPHORE_MAX_JUMP), the exact poison gr_pub_armed exists to
prevent. The epoch reset now only zeroes a word the HOST published
this session (counter_last != 0, m181 semantics); the arm stays at
m181's put-drop and the pristine put <= 1 discovery.
2. The stale-replay leg could never fire: its predicate required
wd_since (set only when a candidate armed the window) - a circular
dependency. The window now arms on either candidate and the action is
decided at fire time with proof of the livelock (submissions moved
during the stalled window).
3. Dead-range reset (reason=stale-dead-range): when the host caught up
but the whole consumed range delivered nothing (counter_last == 0,
submissions >= 8) and the epoch was never armed, rewind last_get to
slot 0 so the fresh session's own entries (which start at slot 0) can
be walked, and force replay for those dead slots (wd_nocap_put) so the
re-walk decodes releases only (idempotent, forward-only) and never
re-captures/re-forwards a dead launch. Once per put epoch
(wd_epoch_done), cleared when new producer work arrives.
4. Log extended: reason=stale-epoch|stale-dead-range|stale-replay plus
reset=/rewind=/cpa=/rpa= so the next run shows exactly which recovery
ran and whether a publish target even exists.
Green safety: both epoch legs require !gr_pub_armed (a channel that never
proved a fresh epoch); a green mid-stream-discovered unarmed channel can
at worst pay one forced-replay re-read of its consumed range per put epoch
(no capture, no word writes). The r220 hold-storm guard and the rewind
cooldown are untouched. make kernel green.
(cherry picked from commit d2ba3cadecf1ca0e72358024f80dbe7a8c468842)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r225 evidence: after a client was killed mid-run, a fresh client's GR
channel (0x5c000059) was registered with a residual USERD put=101 from
the dead session. The m181 fresh-put-drop recovery never fired (the
fresh session never wrote a put - it was stuck at ctx-init, polling
MC_SERVICE_INTERRUPTS once a second), the completion-counter publish
stayed withheld (gr_pub_armed false), and the pin-batch rewind
(SVC_PIN_REWIND=8) pinned last_get at 93 on every injection batch: the
entry loop re-walked slots 93..100 forever, decoding nothing (rel=0, no
capture) and consuming with zero delivery. No capture -> no request ->
no bounded submit -> zero GR-STALL lines; the launch never reached the
bounded path because the session epoch itself was wedged.
- svc_gr_stale_wedge_check (pump pass, before the entry loop, GR only):
tracks (last_get, put) per channel and fires after a 3 s no-progress
window with two signatures:
1. stale-epoch: residual put (put > 1, publish withheld, host believes
the ring fully consumed): arm gr_pub_armed and void the withheld
publish state (m181 semantics, counter_last-gated zeroing; no
last_get rewind - the proven m181 rewind stays with its put-drop
signature), so the fresh client's ctx-init wait gets its host edge.
2. stale-replay: last_get pinned behind put while put == get and the
loop consumed entries during the window (the rewind re-opening an
already-consumed window): drop the window (last_get = put) and put
the pin-batch rewind on a bounded cooldown (wd_rewind_off_until,
cleared by new work) so the futile retry cannot pin the pump.
- svc_handle_map_add: channels under the rewind cooldown are not rewound.
- New per-channel watchdog state written back in the snapshot list
(GRFIX NOTE rule); new diag token gr_wedge (drain line); new
greppable line 'vgsp-gr: GR-STALL-stale-wedge h=... va=... last=...
put=... get=... sub=... waited_ms=... action=recover reason=...'.
Green safety: fresh channels have put <= 1 (signature 1 inert) and
consume their first entries within milliseconds (submissions > 0);
signature 2 needs a stalled (last_get, put) pair for 3 s with no pending
producer work AND entries consumed during the window, which only the
rewind re-opening of a consumed window produces. make kernel green; only
pre-existing warnings.
(cherry picked from commit 6aa5709f57cc1d91a01205042d973d5b203c9bec)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
bigcopy resolution budget
Four guest CUDA clients fail all-bad while four pass; the aggregate
counters cannot separate 'engine result never produced' from 'produced
but written to a different guest page than the client's read resolves'.
Add the two per-page witnesses that correlate the write side with the
read side:
- gsp_phys_gr_copyout_dump (booter, default 1): per-changed-page
OUT-PAGE-WIT va=<device VA> hpa=<guest PA written> first8/last8=<bytes
shipped>, bounded 12 pages per launch, plus one RELOC-CENSUS line per
launch (gaps vs late-resolved words vs late pages).
- sema_bigcopy_log (sema service, default 24, was a hardcoded per-boot
budget): the svc_copy large-copy proof already logs both operands'
resolved VA -> pa/aperture/provenance and the first 16 bytes each side;
raising the budget past the boot's head makes the failing clients'
H2D/D2H resolutions visible.
Discrimination test for the next cycle: same device VA on OUT-PAGE-WIT
(copy-out HPA) and sema-bigcopy (D2H source pa) with different PAs proves
the produced-but-misplaced class; equal PAs redirects to the readback
content/engine-input path.
(cherry picked from commit 78ea8b7807231ac418a17b85c29ed1845e1b0062)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Every GR launch now either completes or fails fast with a distinguishable
dmesg line; no host wait in the launch/consume path can run forever.
- gsp_phys_gr_wait_max_ms (new param, default 60000, clamp [10000,600000],
0 = built-in per-ladder bounds only): absolute per-launch submit budget.
Checked in the GR poll loop (after the release acks) and in the
retirement-retry loop; on expiry the submit fails out ('vgsp-gr:
GR-STALL-wait-max ... action=fail') or the retry loop stops
('GR-STALL-retry ... action=proceed'), so the chain of individually
bounded ladder phases can no longer hold the single-inflight work item
(and the client's launch wait) for the ladder sum.
- gsp_sema_service: pass-top svc_fwd_outcome_check() completes a forwarded
GR request that no hold owns. A GR push with no push-level release
records engages no release hold, so a backend failure (ladder timeout,
abort) left the QMD completion word unwritten and the client waited to
its own timeout (r219/r222, qmd_rel frozen at the previous value). The
check authors the word in software (full resolver, forward-only,
latched once) and logs 'GR-STALL-backend ... action=proceed'. It also
covers a *successful* submit whose release0 could not target the guest
page (engine completion landed on the module semaphore only) - the
work item records gr_rel_unobserved for that.
- Consume tail: a GR capture consumed without ever being forwarded (a
data-stall / busy-defer / no-capture hold fell through) now authors its
QMD completion release, latched per capture, and logs 'GR-STALL-hold'.
- r220 hold-storm fix (196 s re-hold): the m150/m156 data-stall hold
window is keyed to the GPFIFO slot index, and a pin-batch rewind
replays the same launch at different slots, so the window re-armed on
every pass; the absolute per-launch backstop could never fire because
its record fields were excluded from the pump snapshot write-back list
(the GRFIX NOTE record leak) and the fill gate defaulted off. The four
bound fields are now written back and svc_gr_hold_bound_expired is
effective on the default path (cap gsp_phys_gr_hold_max_ms, default
3000 ms, keyed to the launch push VA via svc_gr_launch_va so nothing
re-arms it across rewinds); the busy-defer hold site consults the same
bound. On expiry the entry is consumed with the software completion and
logs 'GR-STALL-hold-bound'.
- Hold expiry/failure edges log 'GR-STALL-fwd-hold' / 'GR-STALL-fwd-fail'
with waited_ms for held forwarded launches.
- Anti-degradation under pin churn/leftover clients: svc_rescue_scan_range
(up to 6 GB of guest-RAM reads on the service thread) is capped at 3 s
per range with RESCUE-SCAN-CAP logging; the ledger eviction burst
triggers a logged wholesale cache reset (LEDGER-RESET cap-and-degrade,
256 evictions); pending-ring overflow is counted and surfaced
(PENDING-OVF, pend_ovf in the drain line). New diag tokens:
gr_swrel/pend_ovf, ledger resets, rescue scan_capped, work_gr
stall_wait_max/stall_retry/wait_max_ms.
Green behavior: normal launches complete in ~0.2-11 s and never approach
the budget; the hold bound only fires where the old path could stall
forever. make kernel green; only pre-existing warnings.
(cherry picked from commit dc3650105f163631bb826ae3685a7d6e9025bfe8)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Trace the CE (copy-engine) data path that services every guest
cudaMemcpy: the guest driver chunks large copies into <=1 MB
LAUNCH_DMA pushes; svc_decode_stream makes one op each and the host
software copy svc_copy executes it (per-256 B bounce, per-page
translation, page-clipped chunks). The hardware-forwarding path only
carries <=4 KB single-page candidates, so it cannot move these.
Completeness verdict: the CE executor is complete-or-error at every
size - the loop copies exactly 'total' or returns an errno (never a
partial success); failures hold the entry (bounded opdef hold, loud
EXPIRED) or are dropped loudly at fwd-ownership. r219 confirms: all 96
audited big copies done==len, zero COPY-MISMATCH, and every CE failure
was a pre-loop resolution miss (rc=-2, no byte moved) - invisible to
the copy-level audit until now.
New instrumentation (all additive, no new locks/waits):
- CEAUDIT per >=32 KB CE op (copy and fill): op/src/dst/len/done/rc,
successes unconditional (4096/boot budget), rc!=0 repeats rate
limited; done=0x0 exposes the pre-loop resolution-miss retries.
- CEAUDIT-MISMATCH: the same true start/end 16-byte destination
readback as COPYAUDIT (moved to the real end sample), emitted next to
COPY-MISMATCH for ce-soft-copy.
- D2H authority visibility: ledger counters + table-fill diag line
(sema-svc: ledger n=/stores/evict/diverges/rewrites) and a
sema-ledger: EVICT warn on the silent home-replacement path.
D2H source verdict (documented in /tmp/opencode/agent-COPY.md): for
pages the launch owns the D2H follows the submit-time ledger landing =
the page the copy-out wrote (correct); it can silently return prior
content for unowned pages (fill-span cuts, work-item drops), on ledger
eviction, and after a failed launch whose copy-out never ran.
make kernel green; only the pre-existing %x warning at
gsp_phys_booter.c:14088.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Map every guest-data copy path and its caps (CE soft copy/fill/inline,
GR fill span, late/gap span, bounce copy-in/out, forwarded CE copy,
launch-page pools) and close the two silent partial-data holes:
- late/gap span cap (gsp_phys_booter.c): a word cut at the 16-page cap
was still rewritten to the partial run; probe one page past the cut
and log COPY-TRUNC when it is resolvable (r218 shape: 4 words x 16
pages = all four cut at the cap).
- GR fill span walk (gsp_sema_service.c): a mid-span resolution miss
after >=1 authored page now logs COPY-TRUNC (was GRFILL-DROP only,
launch proceeded with fill_hold=0).
Instrumentation (all additive, no new locks/waits, readback uses the
same accessor the write just used):
- COPYAUDIT path=... va/len/done for every >=32 KB copy (CE soft copy,
fill, GRCE inline, GR bounce/gap copy-in and copy-out).
- COPY-MISMATCH 16-byte start/end destination readback vs the written
bytes right after the copy (CE first/last chunk; fill and inline first
16 B; GR staging first+last 16 B of every page).
- COPY-TRUNC on every partial copy: CE loop returns, fill/grce chunk
failures, short GR staging/copy-out counts, reloc-fix-span,
gr-fill-span; want>done is rate-limited but never silent.
make kernel green; only the pre-existing %x warning at
gsp_phys_booter.c:14088.
|
| |
|
|
| |
pass (is_arg/hold semantics unchanged)
|
| |
|
|
| |
logging
|
| |
|
|
| |
landing resolver (RELOC-FIXED), param gsp_phys_gr_reloc_resolve default 1
|
| |
|
|
| |
record-destroyer root-cause notes
|
| |
|
|
| |
benign-unresolvable class undecidable (bounded, not bypassed)
|
| | |
|
| |
|
|
| |
relocation coverage counter, zero-write register dump (diagnostic)
|
| | |
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(param/program/banks) resolve through the landing table; zero-write diagnostics
Audit of the QMD relocation/mirror path vs the m180 landing table (r210
evidence: /tmp dmesg, 0 ledger DIVERGE, four multi-page clients still
wrong, 3 zero-write NO-STORES accepts):
The relocation itself does NOT pick PAs: gsp_phys_gr_rewrite_data_vars()
matches a device-VA-shaped 8-byte value against req->data[x].va (the
fill's gr_data list, which m180 fed through the ledger) and rewrites it to
req->data[x].reloc_va. The bounce page for reloc_va is the same page the
tracking wait and the copy-out use, so a *matched* pointer is ledger-
consistent by construction. The un-unified halves of that path are the
page the relocation READS and the page the m129 mirror PATCHES: both come
from pg->pa, resolved by svc_gr_page_note() with the independent sources
(pin/lat/map/walks/scan), never the landing table. The CE inline writes
that author those launches' param/program/bank pages go through
svc_grce_inline_write(), which (m180) records its landing and diverges to
it; svc_gr_page_note could still hand the win0 shadow copy and the mirror
a stale page. Consequences, both observed classes: relocation matches
nothing (the page holds another session's pointers) -> the engine runs
against unrelocated guest VAs it cannot walk -> grid aborts with zero
stores while the ack still lands; or the engine executes on a previous
client's parameter block -> plausible-but-wrong results invisible to the
ledger.
Fix (kernel/gsp_sema_service.c):
- svc_gr_page_note() consults svc_ledger_lookup() first for every
launch-internal page. On a hit it uses the recorded landing (and pins
it, so later resolves and the win0 copy path agree); when the
independent sources would answer differently it logs (bounded)
"sema-ledger: RELOC-DIVERGE va=... ledger_pa=... <src>_pa=... (ledger
wins)" - the requested relocation-vs-ledger witness. Non-hit pages keep
the historical svc_resolve_va path unchanged.
Diagnostics (kernel/gsp_phys_booter.c):
- the m135 below-0x100 guard now shape-checks first and logs
"bank N RELOC-SKIP @off: value is operand page ... (left guest-VA)".
r210 shows the same offset-0 bank pointer in EVERY client (green ones
included: the driver control block), so the guard is unchanged until a
dereferenced pointer in that range is proven; the log makes that
provable in one cycle.
- the zero-write accept now also dumps the deposit QMD words
(gsp_phys_gr_deposit_rb_log) so the next NO-STORES event shows whether
the engine consumed the header words, and the dual-bank sentinel line
is reworded to state facts only (the other bank is normally the
previous launch's legitimately-consumed deposit; the earlier "older
deposit consumed" reading was misleading). r210 classification: all
three NO-STORES events had THIS deposit fetched (sentinel overwritten)
and 0/968..0/967 stores even at +1 s, i.e. engine-side non-execution of
a fetched deposit - consistent with the stale-launch-page abort
mechanism above; the deposit-readback words will discriminate further.
Not changed (deliberate): the m135 guard semantics, the relocation
targets, the sentinel contract, retry/accept policy, and every semaphore
path (mm2/gate/rewind effects untouched).
make kernel -j green (srcversion DAB7E3AA75D364382DD6A3C; only the
pre-existing gsp_phys_rmcall_stage3 printf warning). Main tree untouched;
nothing pushed.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
device VA (kills the write/read translation split of the always-failing class)
Context (r209, fc2d1ae5): multikernel/memcpy_roundtrip/streams2/
victory_streams_events still fail with self-consistent, silent wrongness
while every launch reports OK. Evidence: streams2's bad data is exactly the
PREVIOUS client's h_in pattern (engine read stale inputs); the memcpy
window logs 16x "sema-ledger: DIVERGE read ... wrote=0x7f83000
read=0x8783000" (D2H resolved a different page than the write landed on);
the failing four's device operands are all sysmem-backed (ap=2,
pa=0x7800000/0x8800000 in every svc-gr cand[] line), i.e. exactly the
aperture the m171/m172 ledger never recorded.
Root cause class: a data-bearing guest VA had NO single authoritative PA.
Write paths (H2D svc_copy dst, fills, inline writes, engine write-back)
each landed wherever their own resolution said; read paths (GR launch
fill, D2H svc_copy src) resolved independently again. The ledger patched
up some of that, but (a) it recorded only VID-aperture copies >=64 KiB, so
the dominant sysmem case was invisible; (b) it was a 2048-entry RING -
one launch's submission record set (2505-4553 pages in r209) evicted the
H2D records of the same client before its own D2H read them, so the
readback fell back to its own (different) resolution; (c) several
resolution failures returned -ENOENT even when the landing was known; and
(d) copy paths deferred to the forwarded-request machinery skipped the
in-loop record entirely.
Fix - one landing authority, consulted everywhere:
1. The ledger is now a keyed hash table (16384 slots, 16-probe, one entry
per VA page, in-place rewrite). No cross-VA eviction: a client's whole
working set (8 MB in + 8 MB out + engine records) survives until its
readers have run. Rewrites are counted + logged (bounded) because a
write path moving an existing page is the write-side half of DIVERGE.
2. Every guest-visible write records its landing page:
- svc_copy dst_virt: SYS as well as VID (the sysmem half was the hole);
- svc_fill: SYS as well as VID;
- svc_grce_inline_write: unconditionally;
- small copies that are deferred to the forwarded path: recorded at
candidate capture (the software loop never runs for them);
- engine launch data pages: unchanged (submission-time record).
3. Every read/write consults it:
- GR fill: already diverted (m171); now actually has SYS records to
find;
- D2H svc_copy src: per-page divert (m171) + NEW fallback to the
recorded landing when the resolver fails (previously -ENOENT);
- svc_copy dst / svc_fill / svc_grce_inline_write: NEW - resolve,
then land on the recorded page when one exists (one page per VA,
readers already follow it), and fall back to it when resolution
fails. This also spares the m169 stale-walk hold for destinations
whose page is already known (no pin needed).
Effect on the failing class: the H2D/fill write landing becomes the page
the engine fill, the copy-out and the D2H all use; where a stale and a
fresh resolution disagree, the readers follow the landing and the other
translation is logged as DIVERGE (read/write) or rewrite - bounded, once.
streams2's c1 (engine read stale) is healed by the new SYS records;
streams2's c2 (readback of an engine output after later launches) is
healed by the hash table, which no longer evicts the first launch's
records.
Zero-write lottery diagnostics (secondary, booter; no behavior change on
the green path):
- content-hash settle rule in gsp_phys_gr_output_retire_wait: the
"settled" verdict now requires the per-page checksum FOLD to be stable
for ~300 ms, not just the changed-page COUNT (a continuing store ramp
with a constant count used to read as settled - the r209 "17/32
settled"/"2048/4552 settled" partial copy-outs);
- zero-write accept now logs both qmdwin banks' sentinel words
(ours intact + other bank overwritten = an older deposit was consumed,
i.e. mis-targeted deposit; ours overwritten = this deposit was fetched)
and re-scans the tracked pages once at +1 s: LATE-STORES (copy-out will
ship them) vs NO-STORES (dispatch lost or read-only kernel). The +1 s is
bounded to the rare zero-write accept path only.
The sentinel constant itself is left intact (it is load-bearing for the
m179 post-mortem classification); the dual-bank read answers the same
mis-targeting question without changing the engine-visible contract.
Untouched on purpose: mm2 publish arm/gate, GR fresh-put rewind, the
release machinery and every semaphore path (no regression surface there).
make kernel -j green (new srcversion); main tree and the running VM were
not touched; nothing pushed.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(fix r207/r208 victory_matmul second-client hang)
Regression window: r202 (m178) passed mail->matmul on the same host;
r207/r208 (8c634144 + 5de33b14 + 8ceea88d) hang the SECOND client of a
boot: victory_matmul stops after one GR push per channel and 44/51 CE
pushes (no pushes, no forwards, no launches for 242 s; the guest RM
thread idles on 1 Hz MC_SERVICE_INTERRUPTS). r208's bisect attempt did
not actually change the kernel (identical Module srcversion) and
reproduced the hang, so the rewind cannot be exonerated by it.
Mechanism: 8c634144 made svc_counter_publish target release_pa - the
channel's own completion word, i.e. the UVM tracking semaphore the
client polls. In a fresh session's first pass the pump consumes the
PREVIOUS session's residual ring and then publishes the residual GP_PUT
(109/112, also 96 on rank 16) into that word. The 8ceea88d completion
gate does not suppress it: the stale consumption satisfies last_get ==
put, which is exactly the gate's pass condition. The guest UVM then
observes a far-forward payload before the real release sequence; the
later downward observation is a >2*entries jump (UVM_GPU_SEMAPHORE_MAX_
JUMP), corrupting the 64-bit tracking counter and wedging the session's
first waits. The 5de33b14 rewind corrects the ring position but leaves
the polluted completion word and counter_last in place, so the live
ctx-init releases are skipped (rank 16's own pay=1/2 die against the
published 96).
Fix (worktree mm2-fix, tree unchanged otherwise):
- new per-channel gr_pub_armed: a GR channel's word may only receive a
GP_PUT publish once the channel is proven to belong to the current
session. Armed when the USERD already reads put <= 1 (pristine ring,
first client of a boot) or when the fresh-put-drop rewind proves the
session switch.
- svc_counter_publish returns before any write for an unarmed GR
channel (after the pending-release flush, so retry writes still run):
a residual put can no longer reach a fresh client's waited word, and
the pre-rewind race window disappears entirely.
- the GR rewind additionally voids the stale pass's publish state:
counter_last is reset and the word is re-armed to 0 so the re-decoded
live releases re-advance it from the client's own 1,2,... sequence
(idempotent when no stale publish happened).
Why the first client worked: a fresh boot's rings are pristine
(put=0/1), nothing stale to publish. Why post-hang clients worked: the
hung session's final put is 1, so the residual publish value equals the
fresh client's own first payload (harmless).
Build: make kernel -j green (only kernel/gsp_sema_service.c touched,
+87 lines; vgsp.ko srcversion 40EF888D2E67B00CF9F43E6, differs from the
running r207/r208 module 1418681BF898EE0C1071019).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
while entries are outstanding
r205 regression (my 8c634144 + 5de33b14 + m180): victory_mail fails in
13 s (was green on r202), victory_matmul hangs 242 s. Root cause: the
rank-pollution fix made svc_counter_publish target the channel's OWN
completion word (release_pa ?: counter_pa), but the publish writes the
ring GP_PUT - the client's SUBMISSION position - into that word. While
the pump still has entries to run the pushes are not complete, so the
publish short-circuited every client wait. Witness (r205, mail):
'[129589.831085] pub-adv h=0x5c000019 pa=0x1274beff0 val=108 cur=104'
(after 'walk ... put=108 get=104 last=104'): the waited word jumped to
the submitted 108 ~1 s BEFORE the engine ran ('[129590.802500] output
retirement wait: 1/968 pages changed after 372 ms (settled)'), the
release payloads for the remaining pushes were then forward-only-
skipped, the client's sync returned early and it dumped the stale
buffer -> STATUS FAIL. Pre-fix comparison (r202, green): the same
client's waited word (pa=...ff0) advanced ONLY by sequential release
payloads ('rel-apply ... pay=1..108 -> written'), while the publish
targeted the ascending slot (pa=...ec0) nobody polls - so the correct
completion timing came from the releases alone.
Fix: gate the publish on the pump's own consumed position - write the
GP_PUT only when nothing is outstanding (ch->last_get == put, the same
equality the entry loop consumes to; wrap-safe for the 1024-entry
rings). The releases remain the completion authority; the publish is
reduced to its designed role: the catch-up for entries the pump
consumed without decoding a release. The descending autoreg and the
release_pa ?: counter_pa target stay (they are what stops the busy
rank-16 channel's put from polluting rank 3's waited word - the evtest3
fix, still required: pre-fix r202 style would redirect that write back
onto 0x20400ffc0).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
residual ring
r203 battery forensics (rank 0, sema_dbg=1): hang and progress alternate
perfectly with the first USERD GP_PUT read of a fresh session. A session
whose first read is residual (106/107/108 = the previous session's final
put, GET==PUT) consumes the stale ring in its first pump pass, then never
decodes the new client's live entries: its ctx-init completion word (the
channel's descending counter slot) freezes at the first stale-decoded
payload (1), publishes stop (its USERD put reads never advance past 1),
and the client polls forever - multikernel (gen3), streams2 (gen5),
victory_mail (gen7) and smiley (gen9) hung; memcpy (gen4), streams_events
(gen6), victory_matmul (gen8) - each following a session whose final put
was 1 - proceeded. The GR ring/USERD live in fixed VRAM slots
(fb 0x140000 + rank*0x3000) that keep the dead session's state across
svc_session_flush; m176 reset learned counter state but not the ring
bookkeeping.
m178/m179's r48 resync (last_get=0 on the first put drop) applies only
when userd_wait_fresh is set, and the r73 rewind is CE-only
(engine_type != 0), so a GR channel whose fresh put drops below put_prev
keeps last_get ahead of the producer forever. Extend the r73 rewind shape
to GR: rewind last_get to slot 0 (the fresh session restarts its push
sequence at slot 0, and re-decoding slot 0 is idempotent - the
forward-only apply drops the duplicate payload), gated by the same
submissions < gpfifo_entries/2 freshness guard so a genuine wrap is
unaffected.
Live evidence (smiley, killed at 127105): counter page word 0x20400fff0
stuck at 1; ring put=107 get=107 at discovery; pub-adv 107->1; repeated
rel-apply pay=1 -> skip once per pass. Worktree build: make kernel green
(vgsp.ko linked).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
channel's own (descending) counter slot
The client zips GR channels in ascending handle order onto descending
16-byte counter slots: rank r's pushes release slot fec0+(19-r)*0x10,
which is the completion word its stream sync waits on (verified for all
20 ranks across r196/r199/r200). F1 (eea5ea93) registered the ASCENDING
slot fec0+r*0x10 as the publish target, so every channel's ring PUT was
written into the paired rank (19-r)'s waited word. The busy ctx-init
channel rank 16 (h=0x5c000059, PUT~94) then overwrote rank 3's word
0x20400ffc0 from 1 to 94 (witness r199/r200: 'pub-adv h=0x5c000059
pa=...fc0 val=94 cur=1' after rank-3's own release wrote 1); the rank-3
stream's own release (payload 1) was forward-only-skipped and its wait
never completed: evtest3 cudaStreamSynchronize(st[2]) hangs forever
(sync(0)/sync(1) pass because ranks 1/2 pair with the idle channels
18/17) and victory_streams_events hangs on its rank-3 stream (r198:
launches on h=1f/023/027, 302 s window). Restore the descending autoreg
documented at SVC_GR_COUNTER_TOTAL and make svc_counter_publish target
release_pa ?: counter_pa (the same selection svc_counter_release_apply
uses), so the PUT catch-up can never land in another channel's waited
word; release_pa keeps the push-named truth when the autoreg derive
fails.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
Kernel: authorship instrumentation + swrel gate for in-flight GR release VAs;
first-poll shortening + fast extended re-arm (params gr_retry_fast=1,
gr_first_poll_ms=2500); per-attempt win0 bank recompute + shadow recopy in the
outcome-retry loop; poll-break/deposit-readback diagnostics; C1-C10 sema
diagnostics (crel-entry, relx, opdef-clear, walk, pub-adv, hold-fallback,
busy-defer, CAND-OVERWRITE, inline recs); busy-defer hold instead of
software-release consume; counter fast path; QMD word 48/49 skip.
Tools: heartbeat thread, scanner-alive probe on reinstall, full periodic
re-assert, #nopid handling, per-client reset fix (debounce), dead-pid drop,
supervisor grace+threshold.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(1) svc_session_flush keeps svc.direct_base (per-boot constant; the pin-free
targeted scanner's gate/input) and a new gsp_sema_service_guest_reset()
hook wipes boot-scoped discovery state on guest RM reload;
(2) counter-page derive now takes the first READABLE source (vaspaces ->
discovered -> targeted) instead of failing -EIO on a walked-but-garbage
page, with one-shot diagnostics;
(3) new m177c UVM ring-space directory discovery: content-verified
(push_va -> push_pa) samples from the GR content discovery and CE push
pairing feed a second discovered-directory slot (walks the 0x2xxxxxxx
family), consulted by svc_walk_discovered;
(4) GR channels now deliver withheld opdef sets at entry consume (normal
tail and the stale-skip path) instead of waiting for the 10 s sweep;
(5) host tools hardened: semfix_uvmt QMP reconnect + pid-generation dedupe
+ stale-sem eviction + post-success pin latch + heartbeats; pin_gr_stream
accept timeout + guest-scanner reinstall backoff + restart loop +
subprocess timeouts; new pin_supervisor.sh (heartbeat watchdog, capped
restarts) wired into start_vm_keep.sh.
Evidence: r194-r196 cycles; pool resolves src=trusted pin-free; ring dir
0x10801c000 found 4/4; row: 'rel-apply ... written' with tools up.
|
