open-iov.git/tools/cuda-ab, branch master

cuda-ab: soak.sh client list env-overridable (SOAK_CLIENTS); smiley last by default

2026-10-12T03:44:24+00:00

smiley (cudaMallocManaged-class, known-failing until the param-offset fix) moved last so a baseline pass exercises the other 13 first.

cuda-ab: victory-series coverage battery (8 new clients) + soak runner list

2026-10-12T03:44:06+00:00

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).

tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), smiley push wiring, near-native perf harness (offline readiness; guest validation pending)

2026-10-08T15:49:48+00:00

debug/issue: issue register + ISSUE-0001 (reproducibility) + ISSUE-0002 (multiple kernels per VM)

2026-10-08T15:07:42+00:00

- 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.

vgsp+guest-diag: kcore-truth pin resolver (pool/pushbuffer/channel ring+userd) + GR ring userd-derived candidates

2026-10-08T08:15:24+00:00

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.

cuda-ab: rebuild run/victory_mail for the native path (n=1<<15, param-based mail)

2026-10-07T20:12:24+00:00

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.

cuda-ab: victory_mail native-path sizing + param-based mail message

2026-10-07T19:49:31+00:00

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).

vgsp: m110 - relocate data into the engine-walkable main-PT window; shrink the demo

2026-10-07T19:25:31+00:00

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.

vgsp: end-to-end CUDA acceptance in the guest

2026-10-04T20:19:31+00:00

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).

vgsp: host CUDA bridge + sema release recovery; greeneye simulator wave

2026-10-03T18:05:38+00:00

- 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