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.

2 files changed, 38 insertions, 11 deletions

diff --git a/tools/cuda-ab/run/wedge_recover b/tools/cuda-ab/run/wedge_recover
index 2238fc1..3d9351c 100755
--- a/tools/cuda-ab/run/wedge_recover
+++ b/tools/cuda-ab/run/wedge_recover
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;