| Commit message (Collapse) | Author | Age |
|---|
| |
|
|
|
|
|
|
|
|
|
|
|
|
| |
The committed w256 kernel mixed element and block strides
(bx*1024 + tx*4) and could not cover the buffer. Replace with the
well-formed index i = (blockIdx.y*16 + threadIdx.y)*256 +
blockIdx.x*16 + threadIdx.x, writing out[i..i+3] = i..i+3.
Verified on the emulated vGPU: w512 bad=0; w256 bad=0 with all 16
blocks zero-bad=256/full-bad=0; span sweep firstbad=-1, incl. the
page-15 tail (16376..16383). Two consecutive boots green. This
isolates the earlier w256 bad=12288 to the test kernel itself and
retires blk512 as an index-math control; A-read (arow) failures
remain tracked via ctaid_probe PROBE2.
|
| |
|
|
|
|
|
|
|
|
| |
w512<<<8,512>>> writes 16384 ints out[i]=i perfectly (bad=0 across
repeated runs); w256<<<(16,4),(16,16)>>> with the same write does
bad=12288 deterministically: rows 0..15 correct, rows >= 16 wrong. This is
the transformer/row-band failure without any probe machinery, and the host
scan reports the engine wrote 31/31 tracked pages for the failing launch,
i.e. the by>=1 stores exist but do not reach the pages the client reads.
Build: nvcc -O2 --cudart=static -arch=sm_75 -o blk512 blk512.cu
|
| |
|
|
|
|
|
|
| |
The r235 A-TRACE lines could not be attributed to buffers. The probe now
prints ATAG lines with the host buffers' VAs (h_A, h_out) next to the
device VAs (A, out, out2) and one ATAG line per phase (RAWA/PROBE1/
PROBE2/PROBE2B/RAWA2 naming src/dst/buf) right before its copies, so every
host trace event maps 1:1 to a client buffer and phase.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
PROBE2B repeat
r229 first run: PROBE1 (coordinate map) perfect; PROBE2 (flat A read)
wrong at every element - band 0 read A's own values at wrong offsets
(16382/16383 = A's last elements), bands >=1 read float-typed data. Before
concluding engine-side, the probe now records ground truth in-band:
- RAWA: A uploaded then copied straight back D2H with no kernel, first8/
last8 printed (the buffer itself must hold i);
- RAWA2: same readback after both phases (input-corruption check);
- every band prints its raw first 16 values;
- PROBE2B: the same arow_read runs a second time into a fresh output
buffer, so a first-attempt-only failure is the engine's per-launch
lottery rather than addressing.
The probe's own math re-checked: A is int32 everywhere, expected is
row*256+col, and PROBE1's exact pass over the same output pipeline proves
the harness itself is sound. Committed binary rebuilt static sm_75.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
transformer_mt
The vGPU GR engine's failing chains are correct for the by=0 band and
replicate one tile for rows >= 16 with the QMD geometry and the input recs
proven clean, so the remaining hypothesis is engine-side CTA-coordinate
behaviour. ctaid_probe reads it directly in two phases on the exact
(16,4,1)/(16,16,1) geometry:
Phase 1 coord_map writes by*10000+bx*100+ty*10+tx per element (int32
64x256) and reports per-band row0 plus a bad count - a collapsed or
wrong CTAID.Y shows instantly.
Phase 2 arow_read writes A[(by*16+ty)*256+(bx*16+tx)] with A[i]=i, so a
correct engine reads each band's own rows; a collapsed row register reads
one band's rows everywhere.
transformer_mt prints its per-chain geometry (GEO lines) for the host-side
GRID-WIT comparison. Runtime <10 s. Built with
nvcc -O2 --cudart=static -arch=sm_75 (deterministic static runtime).
|
| |
|
|
| |
4-layer FFN, fp32, CPU-reference checked; multi-stream and 5-launch loop variants)
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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.
|
