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.

5 files changed, 165 insertions, 17 deletions

diff --git a/kernel/gsp_phys_booter.c b/kernel/gsp_phys_booter.c
index b5898b6..b966ac3 100644
--- a/kernel/gsp_phys_booter.c
+++ b/kernel/gsp_phys_booter.c
@@ -16303,6 +16303,17 @@ static u64 gsp_phys_gr_sem_wedged;
  */
 static u64 gsp_phys_gr_stall_wait_max;
 static u64 gsp_phys_gr_stall_retry;
+/*
+ * r410 wedge quiet window state. nw_streak counts consecutive accepted
+ * zero-write launches with tracked data (reset by any launch that wrote);
+ * the sticky wedge is the state where this grows without bound while the
+ * lane is kept busy. quiet counts the quiet holds issued; heals counts the
+ * quiet retries after which the engine wrote. Single-inflight submit path,
+ * so plain statics.
+ */
+static u32 gsp_phys_gr_nw_streak;
+static u64 gsp_phys_gr_wedge_quiet;
+static u64 gsp_phys_gr_wedge_heals;
 /* m86: program-page override blob (see gsp-phys-gr-override). */
 static u8 gsp_phys_gr_override[GSP_PHYS_PAGE_SIZE];
 static u32 gsp_phys_gr_override_len;
@@ -20500,8 +20511,33 @@ gr_page_unmapped:
 		    gsp_phys_rmcall_gr_nowrite_retry_value()) {
 			u32 nw_retries =
 				(u32)gsp_phys_gr_nowrite_retries_value();
+			/*
+			 * r410: the sticky per-lane wedge (previous launch
+			 * also no-wrote while the lane stays busy) is only
+			 * cleared by idleness (r409 ~37 s, r416 ~52 s), so a
+			 * back-to-back retry is futile and keeps it alive.
+			 * When the streak says the lane is wedged, hold a
+			 * quiet window with NO submissions before the retry
+			 * (base on the first retry, doubled on the second),
+			 * then let the normal fresh-bank + resched retry run.
+			 * The quiet path relaxes the whole-launch budget
+			 * (its own hold is the bound: 2 holds max) - the
+			 * alternative is accepting every launch silently.
+			 */
+			u32 quiet_used = 0;
 
 			while (wrote == 0 && nw_attempts < nw_retries) {
+				u32 quiet_ms = 0;
+
+				if (READ_ONCE(gsp_phys_gr_nw_streak) > 0 &&
+				    nw_attempts <= 1) {
+					quiet_ms =
+						gsp_phys_gr_wedge_quiet_ms_value();
+					if (quiet_ms && nw_attempts)
+						quiet_ms <<= 1;
+					if (quiet_ms > 240000u)
+						quiet_ms = 240000u;
+				}
 				/*
 				 * STALLFREE: the same whole-launch budget caps
 				 * the retirement retry chain. When it is crossed
@@ -20509,9 +20545,13 @@ gr_page_unmapped:
 				 * its current outcome (the zero-write accept
 				 * below restores the release words and logs
 				 * loudly), so a chain of slow dud retries can
-				 * never hold the client past the budget.
+				 * never hold the client past the budget. The
+				 * r410 quiet path is exempt: it is entered only
+				 * for a launch whose previous launch also
+				 * no-wrote, and its own holds are bounded.
 				 */
-				if (gsp_phys_gr_wait_budget_over(start)) {
+				if (!quiet_ms &&
+				    gsp_phys_gr_wait_budget_over(start)) {
 					gsp_phys_gr_stall_retry++;
 					gsp_phys_gr_stall_log("retry",
 							      req->push_va, start,
@@ -20522,6 +20562,17 @@ gr_page_unmapped:
 				pr_info("vgsp-gr: outcome retry %u/%u (wrote=%u tracked=%u): refill + re-publish\n",
 					nw_attempts, nw_retries,
 					wrote, n_tracked);
+				if (quiet_ms) {
+					gsp_phys_gr_wedge_quiet++;
+					pr_warn("vgsp-gr: WEDGE-QUIET hold %u ms before outcome retry %u (push=0x%llx; %u consecutive zero-write launch(es); no submissions in between)\n",
+						quiet_ms, nw_attempts,
+						req->push_va,
+						READ_ONCE(gsp_phys_gr_nw_streak));
+					msleep(quiet_ms);
+					quiet_used = quiet_ms;
+					pr_warn("vgsp-gr: WEDGE-QUIET hold done (push=0x%llx); lanes re-armed by the retry below\n",
+						req->push_va);
+				}
 				/*
 				 * m156 (E5): on the first zero-write retry, push the
 				 * proven stage-4 GR semaphore block before re-publishing.
@@ -20778,6 +20829,20 @@ gr_page_unmapped:
 					req, bounce_csum_in, &n_tracked);
 				pr_info("vgsp-gr: outcome retry %u done: engine wrote %u page(s)\n",
 					nw_attempts, wrote);
+				/* r410: the quiet-window verdict. */
+				if (quiet_used) {
+					if (wrote) {
+						gsp_phys_gr_wedge_heals++;
+						pr_warn("vgsp-gr: WEDGE-QUIET healed: engine wrote %u page(s) after the %u ms hold (streak was %u)\n",
+							wrote, quiet_used,
+							READ_ONCE(
+							gsp_phys_gr_nw_streak));
+					} else {
+						pr_warn("vgsp-gr: WEDGE-QUIET no heal: %u ms quiet and retry %u still wrote 0 page(s)\n",
+							quiet_used, nw_attempts);
+					}
+					quiet_used = 0;
+				}
 			}
 		}
 		/*
@@ -20808,11 +20873,14 @@ gr_page_unmapped:
 			 * may also be a legitimately read-only kernel
 			 * (AMBIGUOUS).
 			 */
-			pr_warn("vgsp-gr: NOWRITE-%s h=0x%x push=0x%llx n_data=%u tracked=%u wrote=0 retries=%u qmd_rel=0x%x (engine fetch %s)\n",
+			pr_warn("vgsp-gr: NOWRITE-%s h=0x%x push=0x%llx n_data=%u tracked=%u wrote=0 retries=%u qmd_rel=0x%x (engine fetch %s) consec=%u quiet=%llu heals=%llu\n",
 				fetched ? "AMBIGUOUS" : "CONFIRMED",
 				req->h_object, req->push_va, req->n_data,
 				n_tracked, nw_attempts, want_qmd,
-				fetched ? "OK" : "NONE");
+				fetched ? "OK" : "NONE",
+				READ_ONCE(gsp_phys_gr_nw_streak) + 1u,
+				gsp_phys_gr_wedge_quiet,
+				gsp_phys_gr_wedge_heals);
 
 			/*
 			 * m180 (lottery discriminator, agent-order H1/H2).
@@ -20928,6 +20996,17 @@ gr_page_unmapped:
 				pr_warn("vgsp-gr: no-write accept: release words restored (qmd=0x%08x ch=0x%08x); client's own check decides\n",
 					want_qmd, want_ch);
 		}
+		/*
+		 * r410: streak bookkeeping for the quiet-window trigger. A
+		 * launch with tracked data that wrote nothing extends the
+		 * zero-write streak; any launch that wrote resets it. Runs
+		 * even when the retry loop is disabled so the streak always
+		 * reflects the accepted-outcome history.
+		 */
+		if (n_tracked)
+			WRITE_ONCE(gsp_phys_gr_nw_streak,
+				   wrote ? 0u :
+				   READ_ONCE(gsp_phys_gr_nw_streak) + 1u);
 	}
 	/*
 	 * m114: the kernel ran against the pre-commit bounce window; copy the
@@ -22399,7 +22478,7 @@ static int gsp_phys_state_show(struct seq_file *m, void *v)
 		gsp_phys_backend_gr_counters(&b_sub, &b_ok, &b_refused,
 					     &b_timeout, &b_pool);
 		seq_printf(m,
-			   "work_gr: req=%llu fallback=%llu ok=%llu refused=%llu timeout=%llu fault=%llu submitted=%llu backend_ok=%llu backend_refused=%llu backend_timeout=%llu pool_exhausted=%llu forward_param=%u exec_ok=%u sem_wedged=%llu stall_wait_max=%llu stall_retry=%llu wait_max_ms=%u\n",
+			   "work_gr: req=%llu fallback=%llu ok=%llu refused=%llu timeout=%llu fault=%llu submitted=%llu backend_ok=%llu backend_refused=%llu backend_timeout=%llu pool_exhausted=%llu forward_param=%u exec_ok=%u sem_wedged=%llu stall_wait_max=%llu stall_retry=%llu wait_max_ms=%u nw_streak=%u wedge_quiet=%llu wedge_heals=%llu\n",
 			   gr_req, gr_fallback, gr_ok, gr_refused, gr_timeout,
 			   gr_fault, b_sub, b_ok, b_refused, b_timeout, b_pool,
 			   gsp_phys_gr_forward_enabled(),
@@ -22407,7 +22486,10 @@ static int gsp_phys_state_show(struct seq_file *m, void *v)
 			   READ_ONCE(gsp_phys_gr_sem_wedged),
 			   READ_ONCE(gsp_phys_gr_stall_wait_max),
 			   READ_ONCE(gsp_phys_gr_stall_retry),
-			   gsp_phys_gr_wait_max_ms_value());
+			   gsp_phys_gr_wait_max_ms_value(),
+			   READ_ONCE(gsp_phys_gr_nw_streak),
+			   READ_ONCE(gsp_phys_gr_wedge_quiet),
+			   READ_ONCE(gsp_phys_gr_wedge_heals));
 	}
 	{
 		/*
diff --git a/kernel/gsp_phys_gpu.c b/kernel/gsp_phys_gpu.c
index e896664..6bd5aa0 100644
--- a/kernel/gsp_phys_gpu.c
+++ b/kernel/gsp_phys_gpu.c
@@ -1124,6 +1124,35 @@ int gsp_phys_gr_nowrite_retries_value(void)
 EXPORT_SYMBOL_GPL(gsp_phys_gr_nowrite_retries_value);
 
 /*
+ * r410 wedge quiet window (gsp_phys_gr_wedge_quiet_ms; default 20000, clamped
+ * [0,120000]; 0 disables). The sticky per-lane no-execute wedge (100% of
+ * launches no-write while kept busy) is proven to CLEAR when the lane stays
+ * quiet for tens of seconds (r409: a ~37 s client gap healed a wedged run;
+ * r416: 58 consecutive no-write launches during a 126 s client run, first
+ * launches green again after the lane sat ~52 s idle). Immediate retries and
+ * re-arms do not heal it (r409/r154). So when a zero-write outcome follows a
+ * previous accepted zero-write launch, the retry loop first holds this many ms
+ * with NO submissions to the engine, then re-arms (lane enable + schedules)
+ * and re-publishes once. The value doubles per further quiet retry. The
+ * quiet window is recovery-only: a launch that wrote anything never enters it.
+ */
+static unsigned int gsp_phys_gr_wedge_quiet_ms_param = 20000u;
+module_param_named(gsp_phys_gr_wedge_quiet_ms,
+		   gsp_phys_gr_wedge_quiet_ms_param, uint, 0644);
+MODULE_PARM_DESC(gsp_phys_gr_wedge_quiet_ms,
+		 "Quiet-window ms before retrying a launch on a lane whose previous launch also no-wrote (0=disable, default 20000)");
+
+unsigned int gsp_phys_gr_wedge_quiet_ms_value(void)
+{
+	unsigned int v = gsp_phys_gr_wedge_quiet_ms_param;
+
+	if (v > 120000u)
+		v = 120000u;
+	return v;
+}
+EXPORT_SYMBOL_GPL(gsp_phys_gr_wedge_quiet_ms_value);
+
+/*
  * m69 window-0 shadow pool (gsp_phys_rmcall_gr_shadow_win0; default 0). When
  * set, the forwarded GR launch copies the QMD's program/param pages into the
  * window-0 shadow pool (module pages at VAs 0x108000..0x10f000, PTEs planted
diff --git a/kernel/gsp_phys_gpu.h b/kernel/gsp_phys_gpu.h
index c664e12..e02d6d8 100644
--- a/kernel/gsp_phys_gpu.h
+++ b/kernel/gsp_phys_gpu.h
@@ -399,6 +399,16 @@ int gsp_phys_gr_nowrite_retries_value(void);
  * and the sema service authors the QMD completion release in software. */
 unsigned int gsp_phys_gr_wait_max_ms_value(void);
 
+/* r410 wedge quiet window: when a launch's zero-write outcome follows a
+ * previous accepted zero-write launch (the sticky per-lane no-execute wedge),
+ * hold the retry for this many ms with NO submissions to the engine before
+ * re-publishing. The wedge is proven to clear when the lane stays quiet for
+ * tens of seconds (r409 ~37 s, r416 ~52 s); immediate retries keep it alive.
+ * Base value for the first quiet retry, doubled per further attempt, clamped
+ * [0,120000]; 0 disables the quiet path (retries run back-to-back as before).
+ * Default 20000. */
+unsigned int gsp_phys_gr_wedge_quiet_ms_value(void);
+
 /* m73: use host-visible sysmem chunks for V7 instance/RAMFC (m62 shape). */
 int gsp_phys_rmcall_gr_inst_sys_value(void);
 
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;