| Commit message (Collapse) | Author | Age |
|---|
| |
|
|
|
|
|
|
| |
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)
|
| |
|
|
|
|
|
|
| |
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)
|
| |
|
|
|
|
|
|
| |
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.
|
| |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
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.
|
