diff options
| -rw-r--r-- | debug/issue/README.md | 74 | ||||
| -rw-r--r-- | debug/issue/issue-number-1.md | 128 | ||||
| -rw-r--r-- | debug/issue/issue-number-2.md | 92 | ||||
| -rwxr-xr-x | tools/cuda-ab/ab.sh | 1 | ||||
| -rw-r--r-- | tools/cuda-ab/smiley.cu | 230 |
5 files changed, 525 insertions, 0 deletions
diff --git a/debug/issue/README.md b/debug/issue/README.md new file mode 100644 index 0000000..2c08afc --- /dev/null +++ b/debug/issue/README.md @@ -0,0 +1,74 @@ +# Issue register — process + +**Policy: do NOT create issues by default.** An issue here is a *nasty bug*: +something that blocks mission-critical behaviour, is non-deterministic or not +reproducible on demand, corrupts state, or spans subsystems such that a +single-session fix is unrealistic. Ordinary fixes, TODOs and feature work do +not get an issue — they live in the handoff and the code comments. + +File an issue when **any** of these holds: + +1. Reproducibility / liveness is variable (works sometimes, stalls others) + and the variance itself is the bug. +2. The failure blocks a 1.0 acceptance criterion (guest CUDA, multi-VM, + performance target, ...). +3. Cross-subsystem state (host module + guest driver + physical GPU + + userspace tools) is involved and the ownership of the fix is unclear. +4. The investigation has already burned more than one session without a + confirmed root cause. + +## How to file + +* Directory: `open-iov/debug/issue/`; solved ones move to + `open-iov/debug/issue/solved/` (never delete, never renumber). +* Filename: `issue-number-N.md` with the next free integer `N` + (`issue-number-1.md`, `issue-number-2.md`, ...). Create + `issue-number-N+1.md` as a copy of the template below — do not create + placeholder issues. +* Keep it **public-safe** (this directory is tracked): describe components, + behaviour and evidence (same rules as every tracked document). +* Reference evidence by artifact: session/log file names, counter lines, + commit hashes. Every claim needs its evidence pointer. + +## Template + +```markdown +# ISSUE-00NN: <one-line title> +* Status: OPEN | FIX IMPLEMENTED (VALIDATION PENDING) | SOLVED +* Filed: <date>, session <label> | Priority: <top/p1/p2> | Component: <kernel/guest/...> + +## Symptom +What is observed, by whom, how often. Include the exact observable signature. + +## Impact +Which acceptance criterion is blocked; what work is blocked behind it. + +## Evidence +Logs/counters/commits that establish the facts. Include the one-line +discriminators between healthy and failing runs. + +## Root cause +The causal chain established so far, with confidence and what is still +hypothesis (label uncertainty explicitly). + +## Fix / fix plan +Code changes with commit hashes; for open items, the ordered plan and the +files/functions to touch. + +## Validation plan +Acceptance criteria (numbers of consecutive runs, counters that must stay +zero, suites to run) and how the fix is rolled back if it regresses. + +## History +Dated one-liners: what was tried, what it taught, what was rejected. +``` + +## Lifecycle + +* OPEN: filed, root cause partial. +* FIX IMPLEMENTED (VALIDATION PENDING): code landed (committed) but the + acceptance runs have not completed on a live stack. +* SOLVED: acceptance criteria met with evidence; move the file to `solved/` + and add a `## Resolution` section (what fixed it, the validating runs, + residual risks). If the same failure returns, file a **new** number and + cross-reference — do not reopen the moved file. diff --git a/debug/issue/issue-number-1.md b/debug/issue/issue-number-1.md new file mode 100644 index 0000000..f23721f --- /dev/null +++ b/debug/issue/issue-number-1.md @@ -0,0 +1,128 @@ +# ISSUE-0001: Guest CUDA is not reproducibly green — the client stalls at DEVICE (GR context init never submitted) + +* Status: FIX IMPLEMENTED (VALIDATION PENDING) +* Filed: 2026-10-08, session `smiley-m1`/deep-debug | Priority: **top (blocks 1.0)** | Component: kernel (`vgsp.ko` sema service + RM emulation) with a userspace-watcher race as the variance source + +## Symptom + +`victory_mail` (and every other stock-CUDA client) run in the guest does one +of two things on a cold boot: + +* **Green**: prints `DEVICE …`, then `MAIL …`, `VADD n=32768 bad=0`, `STATUS OK` + within ~15–190 ms of engine work; host forwards CE and GR pushes + (`work_gr req>0`, `GUEST GR LAUNCH OK`). +* **Stalled**: prints `DEVICE …` and then nothing. The guest RM spins in + `NV_ESC_RM_CONTROL` (`0x2a`, `NV2080_CTRL_CMD_MC_SERVICE_INTERRUPTS` + `0x20801702`) at ~1 Hz forever. The GR rings are **never written by the + guest**; the host sees `work_gr req=0`; FB scans find only CE ring entries. + +The split is per-boot. Historically (m104, m141-r3, vmkeep6/7/14) roughly +30–50 % of boots were green on a *fresh* host; after ~8 h of host uptime and +~20 module claim/release cycles the green rate goes to ~0 % and the failure +can degrade further (the guest agent never comes up; smiley r1–r3). + +**This is not a regression against any code change: m141-r3 (the accepted +victory, commit `895ab8b8`) was itself a lucky boot of the same intermittent +class; the same recipe/commit stalls on other boots.** + +## Impact + +* Blocks the 1.0 acceptance: "stock guest CUDA works flawlessly, every time". +* Forces `GRFWD_RETRIES`/`GRFWD_STALL_SECS` lottery runs (minutes per boot) + and hides new regressions under the stall noise. +* Blocks everything behind it: smiley/multi-kernel validation, P2 kernels, + performance numbers, multi-mdev, two guests. + +## Evidence + +One-line discriminators (host dmesg counters, `sema-svc:` periodic stats): + +| Counter | Green (m141-r3) | Stalled (p13-r4) | Meaning | +|---|---|---|---| +| `pushes` / `applied` | 655 / 655 | **90 / 89, then frozen** | one completion never lands | +| `pushes_nomap` | 0 | **1** | one release page had no translation | +| `gr_rel` / `pub_calls` | 219 / 2801 | **0 / 0** | guest never starts GR | +| `ring_auto` | 26 | **6** (CE only) | GR rings never created/pinned | +| `work_gr req` | 2+ | **0** | no GR submission reaches the host | +| `injected mapping` lines | 10 / 40 | **0** | watcher variance | + +Other confirmed facts: + +* `tools/guest-diag/guest_kprobes.sh` armed all 6 kprobes (`ok`×6) and the + semfix watcher translated CE sem words in stalled boots (p13-r4) — so the + stall is **not** the earlier "kprobes never armed" class (`6eee9ac3` fix + was real but is not this issue). +* The guest's own wait is a tracker completion: wedged stacks end in + `uvm_tracker_wait_for_entry ← uvm_push_end_and_wait ← channel_pool_add ← + uvm_channel_manager_create` (repeat-client instance) — i.e. the guest is + waiting for a CE push completion to become visible in a semaphore/waiter + page. +* `svc_apply_pending()` keeps a release that cannot translate ("not + translatable yet: keep") and retries every pass; nothing else ever moves it. +* Related (separate correctness bug, fixed under this issue): the + `PROMOTE_CTX` (0x2080012b) reply. See "Fixes". + +## Root cause + +The completion (release) of the client's last CE push in the boot-critical +window cannot be written back because its semaphore/waiter page has no +VA→PA translation yet. The release is deferred into `svc_pending[]` and waits +for a translation that arrives only from the userspace watcher pipeline +(`pin_gr_stream`/`semfix_uvmt`) or from the kernel's *partial* fallback set. +The watcher pipeline is racy across boots (kprobe arming vs the client's +first push, scanner vs pool allocation, kcore scans vs tracker windows), so +on some boots every fallback misses; the guest's tracker wait never wakes; +the client never creates its GR channels; nothing ever reaches the GR rings. + +Kernel fallback coverage was incomplete exactly for the pool/waiter page: the +pool resolver (`svc_pool_pa_resolve`) had no walk of the **driver-published +VAS page directories** (`svc_walk_vaspaces`) — the deterministic, watcher-free +source the general resolver (`svc_lookup_trusted`) has always used. + +Confidence: high for the mechanism (the deferral + the missing walk + the +observed freeze at `applied=89/90`). The exact *reason the walk itself misses +on some boots* (directory registered later than the first push, or the +watcher path succeeding first on greens) remains boot-timing dependent; the +fix below removes the dependence rather than chasing the timing. + +## Fix + +1. **`kernel/gsp_sema_service.c` — `svc_pool_pa_resolve`:** add + `svc_walk_vaspaces()` (driver-published directories, read-back proven) + between the self-discovered walk and the targeted rm_mem scan. Kernel-native + translation of the deferred release page no longer depends on the watcher. +2. **`kernel/gsp_sema_service.c` — `svc_pending` starvation telemetry:** + per-entry `tries` counter + bounded logs at 4/64/1024 attempts naming + `va`, `payload` and the consequence. A future starvation is visible in + dmesg instead of silent. +3. **`kernel/gsp_rm_emulation.c` — PROMOTE_CTX reply (supersedes + `fa9296af`):** the reply must keep `bInitialize` set. Per + `kernel_graphics_object.c` `kgrobjPromoteContext_IMPL`: + `if (!params.promoteEntry[i].bInitialize) continue; + kgrctxMarkCtxBufferInitialized(...)` — the flag is the server's + "initialized" ACK; clearing it (fa9296af) made every ctx buffer re-promote + forever. This was a correctness bug in the stall's vicinity, not the stall + itself (the stall predates and outlives it). + +Durable direction (if a stall variant survives the fix): fold the watcher's +job into the kernel entirely (complete `svc_pb_discover`/`svc_pd` coverage of +the CE push/pool windows) so the watchers are accelerators only. + +## Validation plan + +1. Fresh host boot (reboot approved 2026-10-08): accepted recipe, plain + runner — expect `victory_mail` green. +2. **10/10 consecutive cold boots** green within seconds; zero stall + captures; `pending release starved` never logs; `pushes_nomap` returns 0 + on every green. +3. Re-run the smiley workload (`tools/cuda-ab/smiley.cu`) per green VM. +4. Rollback: revert commits 1–3 independently; rebuild; the previous + behaviour (lottery) returns without new failure modes. + +## History + +* 2026-10-02..07: stall hunted through m1xx series; greens m104, m141-r3; + kprobe arming fixed (`6eee9ac3`); plan to kernel-native discovery recorded. +* 2026-10-08: campaigns p1–p13 + smiley r1–r3 all stall; evidence shows the + missing completion + missing pool-page translation; `fa9296af` PROMOTE_CTX + change tested and disproven (and is itself corrected here). diff --git a/debug/issue/issue-number-2.md b/debug/issue/issue-number-2.md new file mode 100644 index 0000000..e04e08f --- /dev/null +++ b/debug/issue/issue-number-2.md @@ -0,0 +1,92 @@ +# ISSUE-0002: Multiple CUDA kernels/workloads inside one VM — sequential, concurrent and simultaneous must all work + +* Status: OPEN (root cause established for the sequential case; concurrent/simultaneous untested) +* Filed: 2026-10-08 | Priority: **p1 (blocks 1.0 "multiple kernels", follows ISSUE-0001)** | Component: kernel per-client/session state + guest UVM/RM interaction + +## Symptom + +* **Sequential (known):** a *second* CUDA process in the same boot wedges: + spins at ~100 % CPU with no output (not even `DEVICE`), host traffic for it + is zero. The first client in the same boot is green (10–34 s). Until fixed, + operations are "one client per boot, reboot for a rerun". +* **Concurrent:** not yet tested (blocked by the sequential wedge). +* **Simultaneous (one process, multiple streams/kernels at once):** not yet + tested; expected to hit the same per-channel/per-client state class. + +## Impact + +* Blocks 1.0 acceptance "arbitrary CUDA kernels flawlessly" beyond a single + sequential run, all realistic workloads (frameworks spawn many processes), + and the P2 kernel suite (`tools/cuda-ab/p2/run_guest.sh`) which runs + multiple programs per boot by design. + +## Evidence + +* Wedged guest stack (vmkeep10 run C, pid 1447): + `uvm_ioctl → uvm_api_register_gpu → uvm_va_space_register_gpu → + uvm_gpu_retain_by_uuid+0x44 [nvidia_uvm]` (main thread), second thread in + `do_poll`. +* Source (guest UVM, `kernel-open/nvidia-uvm/uvm_gpu.c:3129-3163`): + `uvm_gpu_retain_by_uuid` takes `g_uvm_global.global_lock`; + `uvm_gpu_release_locked` (last-release path) holds the same lock while + calling `nvUvmInterfaceUnregisterGpu` — an RM call to the host. If that + exchange never completes host-side, the lock is held forever and every + later client wedges at registration. +* Host view (vmkeep10): run A full RM traffic (channel `0xcaf000d4` + registered), then idle; run B issues **zero** RM/sema traffic; each client + run re-runs the GSP boot handshake (`GSP_INIT_DONE` per run start), so the + host must fully re-arm per session. Around run A's exit: + `sema-svc: completion hint rc=-22` twice on the fresh channel. +* Host counters freeze for the second run (`counter pub_calls=2790 + gr_rel=216`, `drain held=2`, `faked=0`). +* Evidence pack: `backups/evidence-20261002-victory/analysis-2026-10-08-repeat-stall/` + (`run1/2/3.log`, `dmesg-w-vmkeep10.log`, semfix/pin logs). +* An m146/m147 experiment (re-arm per-rank one-shot latches on channel + registration; derive CE rank from ring VA; re-arm on GSP session reset via + `gsp_sema_service_session_reset()` from `gsp_rm_emu_reset`) was implemented + and **reverted** — it did not fix run B and temporarily stopped first-run + greens from reproducing. WIP diff: + `…/analysis-2026-10-08-repeat-stall/m146-m147-session-reset-WIP.patch`. + +## Root cause + +Established for the sequential case: per-client state in the host sema/GR +service is not reset when a new guest client context appears (the guest +re-runs the GSP boot handshake per client). The old client's channel/rank/ +ring/pool state makes the host's completion path go quiet for the new client +(`faked=0`, `drain held=2`, zero new applies), so the new client's UVM +registration RM exchange never completes — and because the guest UVM holds +`global_lock` across that exchange (`uvm_gpu_release_locked` → +`nvUvmInterfaceUnregisterGpu`), the whole guest UVM wedges for every later +client too. Concurrent/simultaneous cases are expected to share this class +(per-channel one-shot latches, single per-vgpu sema context). + +## Fix / fix plan + +1. Reset/rebind host per-client/per-channel sema+GR state when a new client + context appears (new GSP session / new channel registration): re-arm + one-shot rank latches, clear stale ring/userd publications, bound any + drain hold so a miss degrades instead of freezing (the `drain held` path). +2. Per-vgpu sema context + P0 isolation work: ordered hunks in + `backups/evidence-20261002-victory/analysis-2026-10-07/sa-p2-patchnotes.md` + (P0-1..P0-7, stages A1–A6, gates G1/G3, `users/idle` anti-UAF drain). +3. Graduation criteria: sequential 3+ clients in one boot green; two + concurrent processes (different streams) green; one process with N + simultaneous streams/kernels green; then the P2 suite in a single boot. + +## Validation plan + +* Sequential: 3 consecutive `victory_mail` runs in one VM, all green. +* Concurrent: two `victory_mail` instances launched together, both green. +* Simultaneous: multi-stream kernel test (P2 suite addition) green, + oracle-matched against bare metal. +* Rollback: patch reverts independently (see WIP diff for the reverted + experiment). + +## History + +* 2026-10-08 00:05: first-client greens became deterministic (`6eee9ac3`); + second-client wedge discovered the same night (vmkeep9–14). +* 2026-10-08 01:00: root-cause chain assembled; m146/m147 experiment + reverted as ineffective; kernel-native per-client reset identified as the + durable path. diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh index 8ae05ee..5a290e4 100755 --- a/tools/cuda-ab/ab.sh +++ b/tools/cuda-ab/ab.sh @@ -86,6 +86,7 @@ host) run_host ;; guest) run_guest ;; compare) compare ;; all) run_host; run_guest; compare ;; +push) push_bin ;; victory) WORKLOAD=victory_mail SRC="$HERE/$WORKLOAD.cu" diff --git a/tools/cuda-ab/smiley.cu b/tools/cuda-ab/smiley.cu new file mode 100644 index 0000000..42770cc --- /dev/null +++ b/tools/cuda-ab/smiley.cu @@ -0,0 +1,230 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * smiley.cu - vectorAdd draws a smiley face. + * + * Four vectorAdd passes execute on the GPU; each computes one scalar field + * + * c[i] = a[i] + b[i] + * + * over the same n = 32768 elements laid out as a 256x128 sample grid: + * + * pass 0 a=u^2 b=v^2 -> face disk field + * pass 1 a=(u+xe)^2 b=(v+ye)^2 -> left-eye disk field + * pass 2 a=(u-xe)^2 b=(v+ye)^2 -> right-eye disk field + * pass 3 a=u^2 b=(v-ym)^2 -> mouth annulus field + * + * with u,v centre-relative, aspect-corrected coordinates. The host copies + * the four fields back and thresholds them to render a 64x32 ASCII smiley + * (4x4 samples per character cell). The artwork therefore comes out of GPU + * addition results only - no GPU-side drawing, no CPU-side math on pixels + * beyond the thresholds. + * + * Output lines are stable and greppable (matches tools/cuda-ab/ab.sh): + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * MALLOC bytes=<n> rc=<0|err> + * VADD passes=4 n=<count> bad=<mismatches> checksum=<sum> + * RESULT cells=<W>x<H> inked=<cells> checksum=<sum> + * STATUS OK|FAIL + * + * Build with a static CUDA runtime (--cudart=static) so the guest only needs + * the driver's libcuda.so, not a CUDA toolkit installation: + * + * nvcc -O2 --cudart=static -o smiley smiley.cu + * + * Then, from the guest: /tmp/smiley + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cuda_runtime.h> + +#define CHECK(x) \ + do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("CUDA_ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("STATUS FAIL\n"); \ + return 1; \ + } \ + } while (0) + +__global__ void vadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +/* sample grid */ +#define GW 256 +#define GH 128 +#define GN (GW * GH) + +/* character cell = 4x4 samples */ +#define CW 64 +#define CH 32 +#define CELL 4 + +/* feature geometry, in aspect-corrected grid units (u ~ +/-127, v ~ +/-127) */ +#define FACE_R2 10000.0f /* face disk: u^2+v^2 <= 100^2 */ +#define EYE_X 36.0f +#define EYE_Y 30.0f +#define EYE_R2 144.0f /* eyes: r=12 */ +#define MOUTH_Y 24.0f +#define MOUTH_IN 2025.0f /* 45^2 */ +#define MOUTH_OUT 3025.0f /* 55^2 */ +#define MOUTH_V 40.0f + +int main(void) +{ + cudaDeviceProp prop; + float *h_a, *h_b, *h_out, *fields; + float *d_a, *d_b, *d_c; + size_t bytes = (size_t)GN * sizeof(float); + double checksum = 0.0; + int bad = 0, pass, i; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem); + + { + cudaError_t e_a = cudaMalloc(&d_a, bytes); + cudaError_t e_b = e_a == cudaSuccess ? + cudaMalloc(&d_b, bytes) : e_a; + cudaError_t e_c = e_b == cudaSuccess ? + cudaMalloc(&d_c, bytes) : e_b; + + printf("MALLOC bytes=%zu rc=%d\n", bytes, (int)e_a); + if (e_a != cudaSuccess) { + size_t fm = 0, tm = 0; + + cudaMemGetInfo(&fm, &tm); + printf("MEMINFO free=%llu total=%llu\n", + (unsigned long long)fm, + (unsigned long long)tm); + } + CHECK(e_a); + CHECK(e_b); + CHECK(e_c); + } + + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_out = (float *)malloc(bytes); + fields = (float *)malloc((size_t)4 * bytes); + if (!h_a || !h_b || !h_out || !fields) + return 1; + + for (pass = 0; pass < 4; pass++) { + for (i = 0; i < GN; i++) { + int x = i % GW, y = i / GW; + float u = (float)x - 127.5f; + float v = ((float)y - 63.5f) * 2.0f; + float pa, pb; + + switch (pass) { + case 0: /* face */ + pa = u * u; pb = v * v; + break; + case 1: /* left eye */ + pa = (u + EYE_X) * (u + EYE_X); + pb = (v + EYE_Y) * (v + EYE_Y); + break; + case 2: /* right eye */ + pa = (u - EYE_X) * (u - EYE_X); + pb = (v + EYE_Y) * (v + EYE_Y); + break; + default: /* mouth */ + pa = u * u; + pb = (v - MOUTH_Y) * (v - MOUTH_Y); + break; + } + h_a[i] = pa; + h_b[i] = pb; + } + + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + vadd<<<(GN + 255) / 256, 256>>>(d_a, d_b, d_c, GN); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_c, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < GN; i++) { + if (h_out[i] != h_a[i] + h_b[i]) + bad++; + checksum += (double)h_out[i]; + } + for (i = 0; i < GN; i++) + fields[(size_t)pass * GN + i] = h_out[i]; + } + + printf("VADD passes=4 n=%d bad=%d checksum=%.0f\n", GN, bad, checksum); + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + { + static char art[CH][CW + 1]; + int inked = 0, cy, cx; + + for (cy = 0; cy < CH; cy++) { + for (cx = 0; cx < CW; cx++) { + int hit = 0, sy, sx; + const float *f0 = fields; + const float *f1 = fields + GN; + const float *f2 = fields + 2 * GN; + const float *f3 = fields + 3 * GN; + + for (sy = 0; sy < CELL; sy++) { + int y = cy * CELL + sy; + float v = ((float)y - 63.5f) * 2.0f; + + for (sx = 0; sx < CELL; sx++) { + int x = cx * CELL + sx; + int idx = y * GW + x; + int inside; + + inside = f0[idx] <= FACE_R2; + if (f1[idx] <= EYE_R2 || + f2[idx] <= EYE_R2) + inside = 0; + if (f3[idx] >= MOUTH_IN && + f3[idx] <= MOUTH_OUT && + v >= MOUTH_V) + inside = 0; + hit += inside; + } + } + art[cy][cx] = hit >= CELL * CELL / 2 ? '#' : + ' '; + inked += art[cy][cx] == '#'; + } + art[cy][CW] = 0; + } + + printf("RESULT cells=%dx%d inked=%d checksum=%.0f\n", + CW, CH, inked, checksum); + for (cy = 0; cy < CH; cy++) + printf("%s\n", art[cy]); + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + free(h_a); + free(h_b); + free(h_out); + free(fields); + return 0; +} |