| Commit message (Collapse) | Author | Age |
|---|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
structurally (the last per-draw GREEN blocker)
The r583 boot's fail-closed refusal traced to a fill that never executes:
the CE channel-setup push's 8 KB zero fill to out=0x204820000 decoded fine
(type=0 op, dma=0x586, line=8192, value=0) but its EXECUTION resolved the
window VA through the generic chain (pin -> ledger -> translate), which
answered SYS with a stale landing pa=0x1300b0000; the r369 below-floor
guard then refused it ("LOWFILL-REFUSED ... below-floor fill resolved
outside vidmem"), the fill never ran, its DSTPEND note stayed live
("DSTPEND note va=0x204820000 type=0", dma-trace denied every pass), and
every GR launch touching that page deferred to the r556 cap and (since
r583) took the fail-closed no-write path.
The vidmem window is offset-mapped by construction (SVC_VIDMEM_VA_BASE
0x200080000 rationale: the client's vidmem VAs are its window base plus
the vidmem offset into the emulated BAR1 buffer; the same convention
svc_lookup already applies). svc_fill_1d now resolves a VID-aperture fill
whose destination is a window VA structurally - addr = VA - window base,
aperture VID - so the existing contiguous shadow-write block executes it.
Witnessed once per boot ("sema-lowfill-window va=... -> fb offset ...").
No guest page table can express the window offset; the structural rule is
faster AND safer than the stale-looking generic chain.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
has an unexecuted device write
The r581 boot's saxpy red draw (bad=524121 csum=874874129) shows the
residual hole of the r541 dstpend guard: GR-DEC "push=0x20043c198
n_data=2505 dropped=8 partial=1 dstpend=1 copypend=0 defer_exp=0" and on
the final dump defer_exp=1 - the launch overlapped a page whose device
write had never executed (live dstpend records 9.5-85 s old), the guard
deferred through its bounded retry, the r556 15 s absolute cap expired,
and the historical path forwarded the launch anyway: the engine read the
known-stale operand and produced the ~50%-wrong result.
r583: at cap expiry with the overlap still live the forward is REFUSED.
Fail-closed: the candidate's QMD completion is authored in software
(forward-only; rel_done latched so the consume tail cannot double-author)
and the candidate is consumed without an engine launch - the client
completes, the engine never touches a page whose writer did not run, and
the event is witnessed loudly ("GR forward REFUSED at defer cap ...").
This cannot hang the guest: the refusal consumes instead of holding.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
latency root cause)
The r581 boot's guest scanner died mid-boot with ValueError in main():
"invalid literal for int() with base 16:
'0x20400f000\\n0x121000000\\n0x12000ab4c'". A guest-agent transport-escaped
append merged demand VAs onto one line with literal backslash-n, and the
old parser crashed on it. The crash stopped ALL new pins; the host only
re-asserted known pairs, so the UVM pool VA 0x121000000 never arrived
while 185 other injections flowed - the pool-pin latency that blocked
every validation boot.
* pin_gr_stream._demand_send/_demand_loop now write one "echo <va> >>"
per VA (no newlines, no backslashes in the transported command; immune
to transport escaping).
* guest_pin_stream's demand parser extracts every 0x token with re and
skips malformed ones instead of raising out of main().
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
rescued content
r577's foreign-release gate made unreadable entries take the !rel_seen path,
where each entry pays the full ce_grace_ms hold before the exact-slot
fallback delivers its pool step ('skipping stale entry ... after 8000 ms
grace' every ~8 s on the r577/r580 boots; a client's init crawls for
minutes). When the rescue decoded a release but the gate refused it as a
tracking slot, the content is provably stale/foreign: no work may run from
it (rescued work is never executed anyway) and the 8 s wait buys only the
small chance that the real page lands inside the window. Backdate
defer_since so the consume path delivers the pool step immediately; the
translation grace still applies to entries whose rescue decoded nothing
(their content may still be the real push).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
fast-path hint
Pool resolution (thread: the starve wedge surviving the r577 fix): the
boot-scoped pool cache survives session flushes, but a pool page MOVES
between sessions and the old page often stays readable, so the read-back
proof cannot tell - the resolver wrote releases to the stale cache PA while
the client's tracker read its current page. Live evidence (r579 boot):
pool-wr src=bootcache pa=0x11e6d9000 for va=0x121000000 while the injected
pin/current page was pa=0x134730000; the client spun in uvm_tracker_wait
for minutes. svc_pool_pa_resolve now tries the structural walk (the
client's own page tables) BEFORE trusting the cache: a differing,
read-proven walk page wins (src=walk-over-boot) and the cached page is kept
as the fan-out secondary.
Event burst (r579 boot: events_dropped=724, the 'no postable event
(events_n=8 dropped=724)' sentinel every 0.5 s): apply the 250 ms per-event
dedup to the per-completion shotgun (8 records per completion burned the
48/s budget), and hoist the completion hint above the pump fast path so it
runs under load too (a dropped wake is re-posted within 0.5 s).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
registration word
THE 2nd-client event-wait root cause (live guest Oops): the NV01_EVENT
registration word we echoed (e.g. 0x3c00000c) is an OVERLOADED field -
cl0005.h: INDEX=bits[15:0], SUBDEVICE=bits[23:16], FLAGS=bits[31:24] - but
the guest's _kgspRpcPostEvent passes the payload's notifyIndex straight to
gpuNotifySubDeviceEvent_IMPL, which indexes its notifier array with it.
Posting 0x3c00000c produced a wild pointer (base + 0x3c00000c*4 + 0x1508)
and a kernel Oops on every dispatched event:
BUG: unable to handle page fault for address: ffffd04d4235f808
RIP: gpuNotifySubDeviceEvent_IMPL+0x21f [nvidia] RAX: 000000003c00000c
note: victory_matmul[12151] exited with irqs disabled
Fixing recursive fault but reboot is needed!
The Oops killed the dispatch before osNotifyEvent/nv_post_event (hence the
r576/r577 experiments: records consumed, nv_post_event=0, independent of
sequence/bNotifyList/hClient), dropped every event, and tainted the guest
kernel - clients that needed an event wake (the 2nd-client device-state
spins) wedged. Mask the posted notifyIndex to bits[15:0] (the INDEX field).
Also loadable probe r578: gsp_rm_ev_func (event records carry a
caller-chosen RPC function for reachability tests).
|
| |
|
|
|
|
|
|
|
|
|
|
| |
The r577 A/B (sequence, bNotifyList, hClient) all yielded nv_post_event=0
with every record consumed by the guest: the drop is content-independent.
This knob posts event records with a caller-chosen RPC function instead of
POST_EVENT 0x1003; posting an event whose handler has an unmistakable guest
side effect (UCODE_LIBOS_PRINT 0x100C trips NV_ASSERT_FAILED on an
unsupported ucode; OS_ERROR_LOG 0x1006 logs an errString) proves whether
the guest dispatches our status-queue event records at all - dispatch vs
consumed-without-dispatch decides the next probe (inside _kgspRpcPostEvent
or earlier in the drain).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
slot; event-record A/B knobs; GR-DEC dump deepened
r575 starved-release wedge (live): an unreadable UVM CE entry's rescued
stale page decoded a release to a foreign VA (0x12006c004, pay=60); the
rescue paths adopted it as the channel's tracking slot (last_rel_va) and
every later recovery release targeted it, while the client's real pool
(0x121000000) stopped at 25 - the client spun in uvm_tracker_wait for 34
(exact-slot poke delivered it and the client completed green). For CE
channels only a pool-shaped release target (0x121000000..0x122000000) may
become last_rel_va; a foreign rescued release is refused (witnessed,
bounded) and the pool step is delivered through the exact-slot fallback /
deferred-release machinery instead. The fake's read-before-write keeps the
delivery monotonic even when a legit push carries multiple releases.
Also: gsp_rm_ev_{seq,nlist,cl} runtime knobs for the guest-RM event drop
A/B (the r576 experiment: guest consumes every POST_EVENT record and never
enters nv_post_event), and the GR-DEC launch dump now covers 256 data
pages (the mail VADD fill carries n_data=969).
New: tools/cuda-ab/evhold.cu - deterministic CUDA event holder/wait probe
(context + cudaEventSynchronize + a hold window) for the event experiments.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
The r575 atomics wedge: a UVM tracker wait for payload 34 on its channel
tracking semaphore stalled at 25 while the pump consumed the channel's
remaining ring entries via the unreadable-entry grace/fake ladder and the
client's wait never cleared (a manual exact-slot delivery of 34 via
poke_guest unblocked it and the client completed green). The recovery
release chosen by svc_fake_release is the suspect: its pool target is
last_rel_va, learned from readable pushes, and a foreign release VA
appearing there would send every later recovery to the wrong semaphore.
Log every authored/queued recovery with its chosen target and payload
(bounded, gsp_sema_dbg) so the next draw names the divergence.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(gsp_phys_gr_dec_log, gsp_rm_ev_log)
The GR launch-decision dump logs, for the first N draws whose fill saw
dropped pages / a partial argument set or whose decision point carries a
dstpend/copy-pending verdict: the decision verdicts, the launch's resolved
data pages (VA -> guest PA), the live deferred-write records and the pending
CE opdef entries - the operand evidence needed to explain a red draw of the
readback lottery (the wrong-but-similar page hypothesis).
The event-post witness logs every EVENT_POST_EVENT record actually sent with
its full payload (hClient/hEvent/notifyIndex/data) and send result, plus a
sentinel when post_all_events found valid events but posted none while none
were dedup-skipped. This is the evidence needed when a client parks waiting
for a completion wake-up that never arrives (the 2nd-client device-state
spin): it shows whether the post was generated, for whom, and whether the
status-queue write accepted it or dropped it (burst cap / ring full).
Both are inert by default (budget 0).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
The page-level dstpend guard is blind when the launch's argument base was
dropped (GRFILL-DROP): the launch's data set then lacks the buffer's pages.
r571 mail: the VADD's copy-in ran at 319818.68 while the b H2D executed at
319820.16 (its base dropped, 9 pages unresolved) - the engine computed
c = a + 0 again, same signature as r567/r568. r570 passed only because the
page-level guard happened to catch it there.
svc_copy_pending(): true when any active opdef slot belongs to a CE
channel (a withheld CE release = a device write that has not executed;
GR channels excluded so a launch's own withheld release cannot defer
itself). A launch that is partially resolved (gr_arg_partial or
gr_dropped > 0) now defers while any such write is pending, bounded by the
same busy-defer + absolute 15 s cap machinery. Fully-resolved launches are
unaffected (the page-level guard already covers them).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
stays owed
The r569 mail failure's missing half: the b H2D's COPY op needs have_in
(the staging upload), which was not resolved on the early passes, so the
push decoded NO op; the release was still applied (the client's
cudaMemcpy(b) returned at 318935) and the slot was marked complete, while
the copy itself only executed at 318952 (17 s later) - the launch's copy-in
(318945.7) read b = 0 and the engine computed c = a + 0 (readback src16 =
a's values). The r541 guard had nothing to defer on (no op, no dstpend).
A trusted, allow_ops, transfer-shaped push that decodes no op with a
device destination now:
- notes the destination via svc_dstpend_note (the launch guard defers
until the copy runs or the 10 s window expires), and
- sets svc.op_owed (the slot stays incomplete, so a later fresh pass runs
the copy instead of replay-suppressing it forever).
No hold is armed, so a push whose trust never arrives cannot wedge the
pump; the guard provides the ordering, exactly like the r541 design.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
launch guard
Root cause (r567/r568, deterministic): the mail client's b H2D COPY op was
COLLECTED but never attempted while its entry was held for another operand;
the VADD launch's copy-in ran 0.6 s BEFORE the copy executed, so the engine
read b = 0 and computed c = a + 0 (the readback page held a's values
{0.0,0.5,1.0,1.5} vs the correct {0.0,0.75,1.5,2.25}). The r541 guard only
saw copies that FAILED with -ENOENT (svc_dstpend_note at the failure site),
so a merely PENDING copy was invisible to it.
svc_dstpend_note_ops() runs after each svc_run_ops() call in
svc_decode_stream whenever svc.op_tr_pending is set: every collected
device-dst copy whose per-op done bit is not set (i.e. did not execute this
pass or earlier) gets its destination noted. The launch guard then defers
until the copy runs (the copy's success path clears the record), bounded by
the existing dstpend TTL + busy-defer + 15 s cap machinery.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
The UVM pool reuses the same launch push VAs for every client of a boot.
svc_gr_cap_hist[8] matches on the push VA ALONE (no content hash) and was
boot-scoped, so a previous client's capture could classify a later
client's fresh launch as already captured and skip it (LAUNCH-SKIP
(replay; already captured=1) from a dead session's record). svc_gr_seen
(hash-keyed) and svc_inline_once (content-keyed) share the same lifetime
and the same class of risk.
Reset all three in svc_session_flush; they refill within the session from
the live client's own pushes, so same-session dedup is unchanged.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
r565 evidence: the fill dropped 4 argument bases (GRFILL-DROP), RDW showed
nrecs=2 where healthy launches run 968-969, the launch was still submitted
and the backend timed out after ~19 s (GR forward failed rc=-110); the
tail then authored the software release and the client kept spinning on a
completion for a kernel that never ran - 20 s of backend occupancy plus a
false completion, and a plausible wedge source for the following clients.
gsp_phys_gr_submit now refuses n_data < GSP_PHYS_GR_MIN_DATA_RECS (16)
with a loud line, failing fast into the same software-completion fallback
without touching the engine.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
| |
cumulative counters
r566 evidence: the dump fired 4x with done=8929 and opdef=215/7/5 frozen
while the ring was EMPTY (every channel last_get == put - the clients were
waiting, no entry pending) and the forced flush had nothing to deliver:
is a lifetime counter (215 holds applied/expired long ago),
not proof of a current hold. A real stall has either a pending translation
(svc_pending_n) or an ACTIVE opdef slot with releases recorded
(e->active && e->n); the dump (and with it the self-heal arm) now requires
one of those. The r564-class true stalls (active slots holding releases)
still dump and still heal.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
opdef flush + per-channel state)
r564 evidence: the STALL-DUMPs fired (4x, done=12700 held=727 frozen) but
ZERO STALL-HEAL lines - the force-consume budget never engaged because the
walk breaks BEFORE the instrumented hold sites (the cheap-wait unread
window at the loop top, or a stale GR hold/nocap index). The self-heal now
acts on the state every break site consults:
- svc_opdef_force_flush(): deliver every withheld release forward-only,
record the delivered marks (so a re-decode cannot re-hold), clear the
slots - the same edge the 3 s bound would run, at the watchdog's
discretion;
- per channel (under chan_lock): unread_va/unread_until = 0 (the
cheap-wait break), gr_hold_idx/gr_hold_until/gr_nocap_idx/gr_nocap_va
reset;
- the force-consume budget stays as the belt for the hold sites.
One loud STALL-HEAL line per firing.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
progress)
Every remaining frozen-pump shape (r562: STALL-DUMP no entry consumed for
20 s, held=99 opdef=150/11/3 frozen, pool_dual=0 conflicts=1) has the same
end state: some hold condition keeps winning and the walk breaks on the
same entry forever, whichever hole was closed last (r548..r558 each closed
one). Add the single backstop that cannot be out-run: when the r539 stall
watchdog fires (20 s without consuming a GPFIFO entry while withheld work
exists), arm a budget of 4 forced consumes; the next hold decisions that
would park an entry are bypassed (opdef hold, no-capture hold, busy-defer
re-hold, data-stalled hold), the entry is consumed with its withheld
releases applied forward-only and the GR QMD completion authored by the
existing tail, and the client's wait advances.
Bounded by the dump cadence (<=4/boot) and the per-arming budget, so a
healthy boot never sees it; each bypass logs one STALL-HEAL line. This
turns 'the client stalls forever' into 'the client completes with the
best-effort data the guards would have fallen back to anyway'.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
walk)
r559 soak2 evidence (victory_saxpy, 167 s at 100% CPU, module alive and
printing stats but consuming no GPFIFO entries): the guest pin stream
flooded the injection queue (svc_handle_map_add: 1867-6101 callbacks
suppressed per burst, several bursts per 10 s) and the unbounded for(;;)
drain spent every service pass processing injections, so the ring walk
never ran - work_guard counters froze (deferred=513 applied=50 unchanged
for minutes) and the client spun.
Process at most 64 queued requests per svc_handle_map_add() call. The
drain runs again on every service pass and from svc_gr_request/
svc_run_ops, so nothing is lost; under a flood the ring walk now gets its
cycles every pass. Bounded batches are the same fairness rule the ring
walk itself already uses.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
| |
r557 STALL-DUMP evidence (victory_matmul): "no GPFIFO entry consumed for
20s" repeated while opdef=3880/16/64 and held=155 climbed - the pump
froze on one held entry whose operand never resolved, re-holding it every
pass. The opdef hold was the only pump hold path without the absolute
per-launch-VA bound (svc_gr_hold_bound_expired, the r509 one-shot used by
every GR hold path); it could therefore extend indefinitely through
fresh slot allocations even though each individual slot window is 3 s.
Consult the same absolute bound before holding: once it fires the entry
flushes its withheld releases and is consumed, so no operand-pending entry
can freeze the pump's entry loop.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r553/r555 boots still show clients walking the individual defer/hold
bounds for 400+ s: victory_reduce 430 s (same push VA re-deferred every
pass, r554's 64-slot latch), victory_saxpy 254 s at write time with
work_guard deferred climbing by the thousand. Each guard is bounded on
its own, but their bounds interact and re-arm across re-decode passes.
Arm an ABSOLUTE per-push-VA deadline on first sighting (64-slot table);
once 15 s have passed, svc_gr_defer_expired() returns true and every defer
site forwards best-effort: the r541 dstpend guard, the r543 arg-partial
guard and the m150 data-stalled hold. 15 s sits far above the 0.5-2 s pin
latency the deferrals exist to cover and far below every client-side
timeout (>= 60 s), so the pin-late cases keep their window and no launch
can be parked longer than the cap, whichever guard keeps winning.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
r553 soak evidence (victory_reduce, 430 s stall, 0 hangs): the same launch
push VA (0x20043be64) was deferred on every pass for 7+ minutes ("partial
argument set: 7-8 pages unresolved"). A client's UVM pool allocates ~8-10
channels and each command buffer reuses two launch push slots, so more
than 8 distinct push VAs rotate through the guard; with 8 latch slots each
new VA evicted another and the one-deferral-per-VA guarantee thrashed -
the second sighting of any VA found its record gone and deferred again,
so a partial argument set that never resolves held the launch forever.
64 slots cover a full client with margin (1 KB of static state).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
r551 soak evidence (victory_strided, same boot as a green battery): the
client's init copies keep failing -ENOENT and are retried by the pump;
every retry re-noted its destination page and re-freshened the r541
dstpend record, so svc_gr_dstpend_overlap() never lapsed and the launch
guard deferred the launch forever: "GR busy-defer timeout -> hold
8000 ms" every ~4 s for 427 s, zero GR LAUNCH submitted, client at 100%
CPU.
Stamp a page's record once per failure episode: an existing record inside
its 10 s window is not re-stamped (the window stays absolute from the first
unexecuted write); an expired record is re-stamped for a genuinely new
episode. The launch then forwards best-effort at the bound exactly as the
guard's own comment promises ("bounded by the same busy-defer + gr_hold
machinery ... it can never hang the launch").
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
651 ms)
Measured A/B on real boots (same module, same client set, only
gsp_phys_gr_settle_ms varied):
- r549 (default 300): 32 launches, min 862 / p50 1119 / p90 2363 ms;
retire wait 680-901 ms per launch ("(settled)", 1..64 of 1000+
tracked pages changed).
- r551 (100): 22 launches, min 476 / p50 651.5 / p90 1010 / max 1428 ms;
retire wait 231-247 ms. Battery 12/15, 0 hangs (unchanged profile;
the 3 fails are the streams2/streams_events/memset2d readback class).
Why the wait was the dominant cost: a read-mostly kernel writes far fewer
pages than the launch's tracked set, so the "changed >= n_all" shortcut
never fires and every launch settles through the content-stable rule,
which paid the full window plus scan overhead. The stability verdict is
still hash-gated (m180); only the minimum stable duration shrinks, and
the battery correctness profile is unchanged. 300 remains available
through the module parameter.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
re-hold cycle)
r550 evidence (victory_strided, reused CE channel 0x5c000059, init sequence
with ~40 device-aperture fills whose walk fails -ENOENT): the sweep
delivered each entry's withheld counter release every ~3 s and cleared the
slot, the pump re-decoded the still-parked entry, svc_opdef_release
recorded a FRESH slot and held it again - SWEEP-EXPIRED repeated six times
for the same entry (entries 41..43 at ~18-21 s each) before
OPDEF-EXPIRED consumed it; the client span at 100% CPU for 12+ min and the
battery timed out (r549's 3 s bound alone only shrank, did not remove, the
cycle).
Remember every delivered release by (channel, entry, va, payload) in a
bounded 256-slot / 60 s ring:
- svc_opdef_release: a re-decode of an already-delivered release claims
it (rel_seen bookkeeping stays correct) and does NOT re-record or hold;
- svc_opdef_hold: an entry with any recent delivery is consumed, never
held again;
- svc_opdef_flush / expire_sweep record what they deliver;
- the ring is per-session (svc_session_flush), because handles, ring slots
and payload sequences are reused by the next client.
Releases are monotonic completion signals and every delivery path is
forward-only, so re-applying one is a no-op - the change cannot double
anything, it only removes the re-hold.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
opfail window)
r548 evidence (matmul, reused channel handle 0x5c000059 after the previous
client exited): the new client's init sequence carries ~65 small
device-destination ops whose operands no resolver covers (memcopy/memset
rc=-2, SVC_VAS_FLOOR destinations like 0x2000c0100). Each one armed the
opdef hold, expired after exactly 10 s (OPDEF-EXPIRED idx 7..18 at ~10.5 s
spacing) and only then released the counter increment the client waits on:
the pump walked one entry per 10 s while the client spun at 100 % CPU
(12+ min, battery timeout).
The two bounds were both 10000 ms. Ops that DO resolve show 0.5-2 s pin
latency (SVC_OPDEF_HOLD_MS comment), so 3000 ms keeps the m163 guarantee
(a failed device copy still gets its retry once the pins land) and the
op-failure latch now fires on the first 3 s retry instead of a second
10 s window. Both are pure bounds - no path loses its retry.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
saxpy stall)
Live evidence (r547, 16+ min stuck victory_saxpy): both launches were
captured, deferred by the r541 (unexecuted device write) and r543
(partial argument set) guards, and left as unoffered candidates - then
every later pass re-decoded the entries as replays (LAUNCH-SKIP;
already captured=1) and the consume path dropped them with no GR LAUNCH
and no GR-COMPLETE-SW for the client, whose cudaDeviceSynchronize waited
forever (client-side strace: RM-control poll loop, libcuda condvar
500 ms retimer; host log: consume idx=105/106 every ~2 s, releases never
authored for rel=0x204827f7c).
Three coordinated changes:
- svc_gr_arg_defer_ok: per-VA latch table (8 slots) instead of a single
slot. Two interleaved launches made the single slot ping-pong (each
call saw the other VA and re-deferred both forever); each VA now gets
exactly one deferral inside the 30 s window and the next sighting
proceeds, so the guard delays without ever hanging.
- svc_decode_stream: a fresh decode pass no longer wipes a valid,
never-offered candidate (a deferred forward is not stale). The
capture path already rotates to a free candidate-ring slot.
- pump no-capture hold: fires for a REPLAY whose push VA still has an
unoffered candidate in the ring (owed forward), forcing a fresh
capture + request per bounded window; and when the window expires
with the forward still owed, svc_gr_owed_resolve authors the QMD
completion release from that candidate and marks it consumed, so an
entry is never consumed without a completion (no late double-run).
New counter gr_owed in the sema-svc drain line; hold log carries owed=.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
| |
glsi+0x33b80's state machine
Batch-gdb entry trace (r543 boot): engine-init -> 1aba0 -> 22f30 (device
lookup, miss=0xee00002) / 2c360 (object lookup, miss=0xee00003) / 33b80
(device flags + state machine) -> 301b0 (state, 3a6f0 x7, 2ac10 x2) ->
EGL_BAD_ALLOC. All three lookups were entered, so the failure lands inside
33b80's state machine; the busy path (30ef2 in 30b20) sets bit 0x20 on the
per-engine state record and returns 0xee00001 once, and eglcore rejects the
engine without retry. 22f30 is called ~1300x (enumeration), so it is not a
lookup miss. Also records the inconclusive Xorg DMA/push-buffer test on a
boot that had a stalled client.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(bounded)
The r542 matmul failure was the NOWRITE wedge: the fill resolved 969 pages
but 2 of 13 high argument bases resolved 0 pages; with
gsp_phys_gr_fill_hold=0 the fill is log-only for dropped pages and only the
aggregate all-arguments-unresolved rule holds, so the launch forwarded, the
engine fetched the QMD (qmd_rel=0x600df00d, engine fetch OK) and wrote
nothing (retirement wait 0/1032, NOWRITE-AMBIGUOUS fe_esr=0x80000000,
trap=0x80010250); the client read zeros (MATMUL bad=65536 csum=0).
Fix: the fill now records gr_arg_partial when an ARGUMENT candidate page
was dropped while other arguments resolved. svc_gr_request defers such a
launch once per push VA (SVC_GR_ARG_DEFER_MS window, svc_gr_arg_defer_ok
latch) through the existing bounded busy-defer retry, so the pins can land
before the engine reads the operands. A genuinely junk pointer word pays
the window once per push, not per retry; after the window the historical
forward stands, so the guard can only delay a launch, never hang it.
|
| |
|
|
| |
was never pushed - the battery misreported it as a fail; probes/egl_frame/frame_producer now ride along)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Live root cause of the smiley/streams2/streams_events readback lottery
(r539): smiley pass 2 verified a2 + b1 and pass 4 verified a4 + b2 - each
even pass's launch read the PREVIOUS even pass's b page. The b H2D's
release completed the client's synchronous cudaMemcpy while the copy op
itself was still held (b's svc_copy executed at 294026.818; launch 2's
engine read at ~294026.3), so the client legitimately submitted the launch
against a shadow page the H2D had not written yet. Passes 1/3 looked
'correct' only because passes 2 and 3 share the same b formula.
Fix, as recommended in CONT.100/94: a device-dst copy/fill that fails to
execute this pass (rc=-ENOENT) records its destination VA page in
svc_dstpend[] with a 10 s stamp. svc_gr_request, right after the fill and
before any submit, checks the launch's filled data pages against the
table (svc_gr_dstpend_overlap); a hit returns the busy-defer verdict, so
the pump holds the entry in its bounded retry window and re-captures once
the copy has run - the record clears on the copy's success. The 10 s
window restores the historical forward when the copy never lands, so the
guard can only delay a launch, never hang it.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
The r539 boot showed the 5-minute device-init stall class live: a GRCE
channel push (eng=0) carrying an H2D copy is decoded before the pin stream
reaches its operands, so allow_ops=0 and the copy is denied; the r308 path
sets op_tr_pending and holds the entry. That path is capped at 4 tries per
slot ('so a slot whose trust never arrives cannot wedge the channel'), but
once the cap is hit op_tr_pending stays 0 and svc_decode_stream's tail
marks the slot COMPLETE - so every later pass classifies the entry as a
replay and suppresses the copy forever. The pin arrives seconds later; the
copy still never runs and the client spins at DEVICE until UVM's own
retry. Evidence (r539, /tmp/opencode/r539-live.log): held idx 89/91..94 on
h=0x5c000059 with 'denied ... replay=1 trusted=1' and 'sema-pend: hold
... idx=89' for minutes while the client printed DEVICE.
Fix: new svc.op_owed flag. When the pend cap blocks the retry, the pass is
marked owed so slot_ops_complete stays 0; the next fresh pass runs the
copy as soon as the translation/injection exists. No hold is armed for an
owed pass, so an operand that never translates still cannot wedge the
pump (the denied path's liveness guarantee is preserved - only the
freshness is).
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
watchdog
The init-stall boots (r530-r538) share one shape: the pump loops but the
entry loop stops consuming, and a client spins in uvm channel_pool_add ->
uvm_push_end_and_wait while its pool tracker word is frozen. The r538 live
boot showed the cause live: va=0x121000000 flapped between the injected pin
(the guest kernel's tracker->page->rm_mem chain) and the host GMMU walk on
every resolve (pool-boot conflict ping-pong, ~50/s), and the release write
follows whichever source answered that pass - a release written to the
losing page strands the client that polls the other one.
Fix: svc_pool_boot_store keeps BOTH most-recent distinct pages (primary +
secondary, each with its own proof stamp). svc_pool_boot_fanout applies
every successful pool release to the secondary as well (forward-only,
read-back proven, 10 s freshness so a freed/reused page is never written).
Wired into all three svc_sem_write success paths and svc_fake_release's
recovery write. The r536 walk-over-pin preference stays as the primary
choice; the pin is now recorded as the secondary there.
Also: STALL-DUMP - when the entry loop consumes no entry for 20 s while
withheld work exists (pending releases or opdef slots), dump the pending
list, opdef slots, pool candidates and counters once (bounded 4/boot), so
the next stall is named in one place instead of N bounded logs.
|
| |
|
|
|
|
|
|
|
|
|
|
| |
The VAS candidate verification needs the client's ring to be found by the
RAM content scan; on a boot where that never anchors the ring the
candidate stays parked forever and every launch operand walk fails rc=-2
(observed: pd_found=0, vas[0] holding an internal pdb, walks stuck at
9555 successes with 114 K attempts while the client spun for minutes).
After 15 s parked, register the candidate unverified with a loud log: it
came from the client's own RM publish and the newest-first walk order
keeps an older stale directory from winning, so the bounded wait's
registration is the lesser evil.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(moved-but-readable pool)
The pool resolver trusted an injected pin after a read-back proof. That
proof does not catch a pool page that MOVED: the old page is still
readable, so the release write lands on the old page while the client
polls the new one. Observed: 'pending release starved va=0x12100000x
tries=64' for ~45 s with the stale pin (pa=0x120b30000) readable, until
the fresh pin (pa=0x130e9d000) arrived; the client's tracker never woke
and its init hung (multikernel battery timeout, copies stuck rc=-2).
The client's own page tables are the GPU/UVM ground truth for where the
tracker reads, so when a structural walk resolves a different page than
the pin (and it is not a degenerate identity), the walk wins - both
sources are read-back proven at the exact offset.
|
| |
|
|
|
|
|
|
|
|
|
|
| |
A CUDA device buffer is a PROT_NONE reservation (often hundreds of MB).
Its pages are never present in the pagemap, so scanning it yields no
pairs - but it consumed the VMA_MAX_PAGES budget, and the scan never
reached the client's present staging buffers that sit ABOVE the device
reservation. Observed: a client whose device range started below its
staging buffer got no high-VA pins at all; its H2D copies failed rc=-2
forever and the client read back stale data (soak run: 9/9 green then
streams2 bad=32737). Skipping '---' VMAs restores the pins (9 fresh
high-VA injections in the first run; streams2 bad=0).
|
| |
|
|
|
|
|
|
|
| |
(shipped-binary evidence)
Records the eglcore->glsi chain that rejects the engine candidate
(0xee00001 out of the glsi engine-init), the ruled-out hypotheses (RM
controls answer NV_OK; the engine passes the walk's own filters; stable
under gdb no-ASLR) and the exact breakpoints/offsets for the next trace.
|
| |
|
|
|
|
|
|
|
|
|
|
|
| |
starving the pump
The full-RAM discovery scan read one 4 KB page per vfio_dma_rw call:
1.5 M guest reads over the 6 GB range, ~10-11 s per rebuild, and the
rebuild runs every 10 s - so the pump got roughly one pass per rebuild.
A client whose resolutions were dropped by a session flush then waited
minutes for the scan to re-find its pages (the observed post-flush
stall: data-VA copies returning -ENOENT with the pump counters frozen).
Reading 64 KB chunks amortizes the iommu walk + pin over 16 pages; the
per-page detection logic is unchanged.
|
| |
|
|
|
|
| |
(re-compact, not wipe)"
This reverts commit 1a49c7eb69c03f58bc51939828de222a2ac42fea.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
wipe)
The flush fires when the channel table momentarily empties - a live
client's UVM channel cycling triggers it - and wiping the map left the
client's already-resolved data VAs unresolvable until a rebuild scan
re-found them, which it may not. Evidence: a mid-battery boot where
multikernel's operands went NO ENTRY 40+ s after a flush (the rebuild
restored 14-31 K entries without them); the pump stalled for minutes and
the next clients read back whole-page garbage. The map is a cache with a
merge-on-rebuild contract, so re-compact it (dedup + sort) and keep it as
the merge base; reused-VA correctness stays with the ledger landing, the
pin-wins rule and the fresh scan.
Also: push_battery.sh now pushes inplace_scale (the battery list had it,
the push list did not).
|
| |
|
|
| |
advance even when the r391f once-only gate suppresses the write. A re-decode pass that skipped the FIRST run of a split upload left the cursor at the first run's destination, so the SECOND run's data landed at offset 0 (the deterministic tgemm X[0..1]<-X[2..3] 8-byte overwrite, present in every pass); the gate suppresses the write, not the method stream's positional semantics
|
| |
|
|
| |
runtime root alloc answers NV_ERR_INVALID_CLIENT (0x23) because the server's internal client is adopted, so allocate the class-0x2081 child directly under the internal subdevice with a candidate handle list
|
| |
|
|
| |
the live session (root -> device 0xcaf00000 -> subdevice 0xcaf00004 -> class-0x2081 child 0xcaf00099) and send the query as that client, so the server dispatches from a genuine BINAPI object
|
| |
|
|
| |
32-byte request header (hObject at +8, status at +16) - the backend alloc itself succeeded (rpc_result=0) in r522
|
| |
|
|
| |
the SUBDEVICE (0xbeef0004 is the class-0x2080 child), so mirror that on the backend - the device parent answered NV_ERR_INVALID_OBJECT_PARENT (0x36)
|
| |
|
|
| |
live session (GSP_RM_ALLOC as the internal client/device parent) and send the query to it - the subdevice attempt answered NV_ERR_NOT_SUPPORTED because the server dispatches by the object's class
|
| |
|
|
| |
forwarding only 0x20810107 keeps the P3 whitelist off so the boot is perturbed by exactly the one control under study
|
| |
|
|
| |
reserved-interface query to the live GSP session (internal subdevice translation) and log the server's own status/bytes as evidence; scoped to the observed 2-byte shape, fail-open to the emulated generic reply
|
| |
|
|
|
|
|
|
|
|
|
| |
binary was a copy of streams_events_probe)
The run/tgemm_probe artifact had the same md5 as streams_events_probe, so
every 'tgemm probe' invocation so far ran the streams probe instead. The
rebuilt binary reproduces the transformer class on the r518 boot:
'PASS name=k pass=2 bad=256 xbad=2 first=(0,0)' plus the XDIFF rows that
show the input's first two floats shifted by one pair (X[0..1] read as
X[2..3]) - the deterministic first-page corruption.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
MakeCurrent wall
The guest driver dups memory/VASpace objects into the GL client at
EGL MakeCurrent time; the generic-SUCCESS default replied with an EMPTY
payload and the driver, expecting rpc_dup_object_v03_00
(NVOS55_PARAMETERS {hClient, hParent, hObject, hClientSrc, hObjectSrc,
flags, status}) echoed back, failed the call and tore the whole EGL
session down. r510 guest witness: 'RPC UNKNOWN(21) seq=0 len=60 ->
generic SUCCESS' immediately followed by the session's FREE cascade and
eglMakeCurrent -> EGL_BAD_ALLOC (0x3003); EGL device enumeration and
eglInitialize were already green (n=3 devices, NVIDIA 1.5). The handler
echoes the params with status=NV_OK (the emulation keeps no per-object
state for these classes).
|
