From e77efad422961e5266ef30ef5ce9f49ae1fd76c1 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Wed, 21 Oct 2026 11:39:25 -0400 Subject: gr: r410 wedge quiet window - bounded idle hold before retrying a lane whose previous launch also no-wrote (r416: 58-launch sticky wedge clears after ~50 s quiet); wedge counters in state; wedge_recover v2 unmask The r416 boot reproduced the sticky no-execute wedge inside 60 s: 58 consecutive launches accepted zero writes while the lane stayed busy (wedge_recover v1's deterministic verify masked every one of them), and the first launches were green again only after the lane sat idle ~52 s. Both r409 (~37 s idle heal) and r416 therefore show the wedge clears on idleness and persists under back-to-back retries (r154/r409), so the retry loop now holds a quiet window (gsp_phys_gr_wedge_quiet_ms, default 20000, doubling on the second quiet retry, 0 disables) before the normal fresh-bank + resched retry when the previous launch also no-wrote. The whole-launch stall budget is relaxed only on this bounded recovery path. NOWRITE lines now carry the consecutive streak and quiet/heal counters, and the state file shows nw_streak/wedge_quiet/wedge_heals. wedge_recover v2 mutates the inputs every iteration so a stale (never executed) launch can no longer match the expected values, and classifies the readback STALE-PREV(NOWRITE) vs OTHER/PARTIAL. --- tools/cuda-ab/run/wedge_recover | Bin 821024 -> 825120 bytes tools/cuda-ab/wedge_recover.cu | 49 +++++++++++++++++++++++++++++++--------- 2 files changed, 38 insertions(+), 11 deletions(-) (limited to 'tools/cuda-ab') diff --git a/tools/cuda-ab/run/wedge_recover b/tools/cuda-ab/run/wedge_recover index 2238fc1..3d9351c 100755 Binary files a/tools/cuda-ab/run/wedge_recover and b/tools/cuda-ab/run/wedge_recover differ diff --git a/tools/cuda-ab/wedge_recover.cu b/tools/cuda-ab/wedge_recover.cu index deba670..fd8f3f9 100644 --- a/tools/cuda-ab/wedge_recover.cu +++ b/tools/cuda-ab/wedge_recover.cu @@ -3,12 +3,17 @@ * wedge_recover.cu - reproduce the per-launch no-execute (wedge) and test * whether an idle gap heals a wedged context. * - * Each iteration launches a small vecadd, syncs, reads back with two separate - * D2H copies (two independent readback paths so a single copy flake is not - * counted), verifies, and prints one line. On a verification failure the - * launch is retried after a configurable sleep (default 20 s): if the wedge - * heals by idling, retried iterations pass; if it needs a fresh context they - * keep failing. Latency per iteration is printed so the fix knob is visible. + * v2 (session 10, r416 finding): the v1 loop reused one deterministic + * kernel+buffers, so a no-write launch was MASKED - the stale buffer held + * exactly the previous iteration's (identical) values and the verify passed. + * r416 ran 60/60 "OK" while 58 consecutive launches no-wrote underneath. + * + * v2 mutates the inputs every iteration (H2D of fresh ha/hb), so a stale + * readback cannot match the new expected values: a no-write launch must + * report BAD. On BAD the same iteration is retried after a sleep, and the + * readback is classified: STALE-PREV (byte-identical to the previous + * iteration's output = launch never executed) vs PARTIAL/OTHER (engine wrote + * something). Latency and per-iteration outcome are printed. * * Build: nvcc -O2 --cudart=static -o wedge_recover wedge_recover.cu */ @@ -41,7 +46,8 @@ int main(int argc, char **argv) { const int n = 1 << 15; size_t bytes = (size_t)n * sizeof(float); - float *ha, *hb, *hc, *da, *db, *dc; + float *ha, *hb, *hc, *hc_prev; + float *da, *db, *dc; int iters = argc > 1 ? atoi(argv[1]) : 40; int sleep_s = argc > 2 ? atoi(argv[2]) : 20; int fails = 0, heals = 0, i; @@ -50,12 +56,15 @@ int main(int argc, char **argv) ha = (float *)malloc(bytes); hb = (float *)malloc(bytes); hc = (float *)malloc(bytes); - if (!ha || !hb || !hc) + hc_prev = (float *)malloc(bytes); + if (!ha || !hb || !hc || !hc_prev) return 1; + /* iteration 0 baseline; mutated each iteration below (v2) */ for (i = 0; i < n; i++) { ha[i] = (float)(i % 977) * 0.25f; hb[i] = (float)(i % 251) * 0.5f; } + memset(hc_prev, 0, bytes); CHECK(cudaMalloc(&da, bytes)); CHECK(cudaMalloc(&db, bytes)); CHECK(cudaMalloc(&dc, bytes)); @@ -66,8 +75,21 @@ int main(int argc, char **argv) int attempt = 0; long long ms[2] = { 0, 0 }; + /* + * v2: fresh inputs every iteration so a stale (never + * executed) launch cannot match. Mutate after iteration 0. + */ + if (i > 0) { + ha[i % n] += 1.0f; + hb[(i * 7 + 3) % n] += 2.0f; + CHECK(cudaMemcpy(da, ha, bytes, + cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, + cudaMemcpyHostToDevice)); + } + for (;;) { - int bad = 0, j; + int bad = 0, stale_prev = 1, j; cudaEvent_t e0, e1; float ems = 0; @@ -87,6 +109,8 @@ int main(int argc, char **argv) if (hc[j] != expect) { bad++; + if (hc[j] != hc_prev[j]) + stale_prev = 0; } } ms[attempt & 1] = (long long)(ems * 1000.0f); @@ -95,11 +119,14 @@ int main(int argc, char **argv) heals++; printf("ITER %d OK attempt=%d launch_us=%lld\n", i, attempt, ms[attempt & 1]); + memcpy(hc_prev, hc, bytes); break; } fails++; - printf("ITER %d BAD bad=%d attempt=%d launch_us=%lld\n", - i, bad, attempt, ms[attempt & 1]); + printf("ITER %d BAD bad=%d/%d attempt=%d launch_us=%lld class=%s\n", + i, bad, n, attempt, ms[attempt & 1], + stale_prev ? "STALE-PREV(NOWRITE)" + : "OTHER/PARTIAL"); if (attempt >= 3) { printf("ITER %d GIVEUP\n", i); break; -- cgit v1.2.3