| Commit message (Collapse) | Author | Age |
|---|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
freeze)
The guest PMA scrubber (in-guest CeUtils CE channel, 4096 GPFIFO entries)
waits on its channel buffer's finish payload (offset 0x6c004) after the
third CUDA session's pages must be scrubbed; nothing in the emulation
authored it, so cudaMalloc froze in scrubCheckAndWaitForSize.
- channel_add: detect the 4096-entry scrubber channel and stash its USERD
PA/ring geometry (the channel itself is left untouched: no rank change,
no pump change - the colliding numeric-handle FREE may still unregister
it, the geometry survives).
- sema_scrubbuf_add debug param: register the finish/host-semaphore word
PAs the guest scanner resolves (vmallocinfo + kcore walk), because the
RM kernel VA space collides numerically with the client CE window in the
global VA map and cannot be resolved through it.
- svc_scrub_pass (service thread, per pass): from the channel's own USERD
GPPut, track submitted scrub blocks across wraps, then author USERD
GP_GET = put, host semaphore = put and finish payload = submitted.
Forward-only, idempotent, zero writes while the scrubber idles.
Live-validated mechanism (r137 VM): authoring the finish word at the kcore-
resolved PA let the frozen victory_matmul leave scrubCheckAndWaitForSize
and proceed to its launch.
|
| |
|
|
|
|
|
| |
default
smiley (cudaMallocManaged-class, known-failing until the param-offset
fix) moved last so a baseline pass exercises the other 13 first.
|
| |
|
|
|
|
|
|
| |
victory_saxpy, victory_reduce (shared+atomics), victory_matmul (tiled
shared), victory_strided, victory_memset2d (pitched), victory_atomics,
victory_streams_events (3 streams + events), victory_math (mixed
int/float/double). Each verifies via readback and prints STATUS OK|FAIL.
soak.sh now runs the full 14-client set (seq or sustained).
|
| |
|
|
| |
smiley push wiring, near-native perf harness (offline readiness; guest validation pending)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
(multiple kernels per VM)
- debug/issue/README.md: when to file (nasty bugs only; default is no
issues), template, numbering, lifecycle incl. solved/ move.
- issue-number-1.md: guest CUDA is not reproducibly green (top priority);
evidence discriminators, root cause (starved pool release translation),
the three fixes landed, 10/10-boot validation plan.
- issue-number-2.md: multiple kernels in one VM (sequential/concurrent/
simultaneous) must work; second-client wedge root-cause chain + plan.
- tools/cuda-ab/smiley.cu: 4-pass vectorAdd that renders a 64x32 ASCII
smiley from GPU-computed distance fields; ab.sh gains a generic 'push'
mode so any WORKLOAD binary can be delivered to the guest.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
ring+userd) + GR ring userd-derived candidates
guest_pin_stream.py now resolves translations entirely from guest truth via
/proc/kcore + a manual page-table walk (no QMP races): the CE semaphore pool
and pushbuffer memory from the UVM rm_mem structures, and each channel's
GPFIFO ring + GPPUT from the tracking semaphore's embedding uvm_channel_t
(channel_info located by its gpFifoGpuVa value). This removes the watcher
pipeline's flakiness on the cudaMalloc stall class (missing pool/pushbuffer
pins).
gsp_sema_service.c: svc_gr_ring_autopin tries the guest-registered USERD
memdesc base as a ring candidate (the fixed FB convention only matches when
the guest's VRAM allocator happens to reproduce the old layout), plus an
optional (gsp_sema_dbg) one-shot FB scan + raw-word diagnostics when every
candidate fails.
tools/cuda-ab/p2: self-checking arbitrary-kernel acceptance suite (10
programs: sequential launches, shapes, params, smem/GEMM, streams, malloc
churn, printf+error, FMA benchmark, mixed, managed probe) with host oracle
expectations and a guest runner.
|
| |
|
|
|
|
|
|
| |
mail)
The run/ binary was the Oct 4 build (4 MB buffers, __device__ kMail);
ab.sh push_bin only builds when the binary is absent, so the guest shipped
the stale build. This is the rebuilt binary matching today's .cu.
|
| |
|
|
|
|
|
|
|
| |
Two changes for the native GR path: vadd buffers shrink to 1<<15 elements
(128 KB each) so all data fits the engine-visible 2 MB window, and
mail_kernel receives the message as a kernel-parameter buffer instead of a
__device__ constant (the c[0x4] module-data base is driver-populated and is
zero in the forwarded launch). The guest binary must be rebuilt in
tools/cuda-ab/run (ab.sh push_bin only builds when the binary is absent).
|
| |
|
|
|
|
|
|
|
|
|
|
| |
the demo
All UPDATE_PDE_2 installs (post- or pre-commit) are invisible to the SM
walker; the only window it honours is the commit-time main small-PT page
(VA 0x100000+, the proven win0/const window). Relocate each candidate data
pointer's first 64 pages into slots 0x110..0x1f0 of that window (fresh PTE
lines, direct leaf writes), rewrite bank pointers per page record, and put
the mail page at the reserved slot 0x1f0 (c[0x4][0x0]=0x1f0000). Shrink
victory_mail's buffers (n=1<<15, 128 KB each) to fit the 2 MB window.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
The no-claim emulation path now runs a stock CUDA program end to end:
device properties, cudaMalloc, a GPU-authored string kernel read back
byte-for-byte, two 4 MB host-to-device copies, a 1 M-element vectorAdd
executed on the physical GPU through the host bridge, and the
device-to-host copy-back verified element by element. Three consecutive
supervised cycles report MAIL + VADD bad=0 checksum=392344090.500000 +
STATUS OK with identical input/output hashes.
Two data-path correctness fixes make the transfers sound:
- Once-only op execution per GPFIFO entry. The pump replays recently
consumed ring entries when translation batches are injected so missed
releases can land; re-executing a CE copy is not idempotent because
its source is a client staging buffer that may have been refilled.
Slots now record their raw entry words and a replayed entry matching
that identity suppresses operations while still processing releases.
- Per-page translation of virtual operands. svc_copy resolved only the
first page of a VA range and walked the PA linearly; client staging
ranges are frequently split (observed 195 + 61 pages), so the tail of
alternating chunks was read from unrelated memory. Multi-page copies
now re-resolve each 4 KB page through the trusted resolver.
Also carries the earlier uncommitted emulation fixes that got the client
to the launch (QMD address-window gate, counter registration
persistence, 512 MiB emulated FB, bridge slot hygiene) and bounded
diagnostics (sema-bigop, sema-bigcopy, bridge-in/out hashes, daemon
INCHECK/OUTCHECK/dumps).
|
| |
|
|
|
|
|
|
| |
- kernel: gsp_cuda_bridge (QMD launch capture, debugfs job/out/done, param+output resolution cached from inline writes), sema drain fixes (never hold/starve a channel), GR ring autopin + counter autoreg, per-page push reads, counter-derive gate; module builds clean
- tools/cuda-bridge: CUDA driver-API daemon executing mail/vadd kernels on the physical GPU, CUBB2/CBD2 v2 protocol, layout check + self-tests
- tools/guest-diag + tools/cuda-ab: scanner/pin fixes, victory push-only mode
- debug/risc-v-gsp-rm-vm: Phase S simulator wave (GSP device model, boot ladder, triage/re-demo harnesses, docs 30-41), GPGPU-sim scaffolding
- STATE-CURRENT/docs updates
|
|
|
The progress service now decodes push buffers as method streams and performs
the work the driver pushes, not just the tracking semaphore:
- C5B5 method decoder: SET_REMAP_CONST_A/B + COMPONENTS + OFFSET_OUT +
LINE_LENGTH_IN + LAUNCH_DMA = memset (into vidmem or guest RAM);
OFFSET_IN + OFFSET_OUT + LAUNCH_DMA = copy; SET_SEMAPHORE_A/B/PAYLOAD +
LAUNCH_DMA = release. Contiguous method runs (count 2/4) are applied in
full - OFFSET_IN/OUT arrive as one run.
- Pushes span page boundaries, so pages are decoded as method streams
without a start marker. Collected operations execute only when the
push's release is reached and still pending (payload ahead of the
semaphore), so recycled buffers replaying old pushes are skipped and
their stale memsets cannot overwrite newer state.
- VA->PA map is first-wins: stale copies of freed allocations and
unrelated adjacent u64 pairs produced bogus entries that hijacked a
good mapping and sent semaphore writes to the wrong page.
- dump_guest module parameter (hex dump of guest PA, kthread via
completion; kthread_stop before the thread runs loses the work).
With this the guest no longer stalls in uvm_tracker_wait and no longer
oopses in the UVM path; the failure moved to a userspace segfault inside
libcuda (NULL-0x38) after the driver's page allocations.
Also add tools/cuda-ab: run one deterministic CUDA workload on the host
(bare metal oracle) and in the guest, then compare the stable result
lines. The binary links the CUDA runtime statically so the guest only
needs the driver's libcuda.so.
|
