vgsp: m175 periodic expiry sweep for withheld release sets (svc_opdef_flush is only reachable while the pump re-decodes the entry; when a withheld entry is parked and the pump moves on, r189 showed opdef=531/8/0 static forever with no EXPIRED/LOSS and the guest's st[2] counter froze - evtest3 cudaStreamSynchronize(st[2]) hang; svc_opdef_expire_sweep runs each service iteration: slots past SVC_OPDEF_HOLD_MS deliver forward-only to the live owner channel (SWEEP-EXPIRED) or count a LOSS (SWEEP-LOSS); static channel scratch to avoid the 26KB stack frame)

1 files changed, 65 insertions, 0 deletions

diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c
index 7b63c0f..94c3ead 100644
--- a/kernel/gsp_sema_service.c
+++ b/kernel/gsp_sema_service.c
@@ -12121,6 +12121,63 @@ static void svc_opdef_reset_all(void)
 }
 
 /*
+ * STALL FIX (r189 entry=82/83 va=0x20400fef0 pay=84/85): the 10 s deadline
+ * flush (svc_opdef_flush) is only reachable while the pump keeps RE-DECODING
+ * the entry. When a withheld entry stops being revisited (ring advanced, the
+ * channel went quiet, the entry was parked while the pump moved on), the
+ * deadline never runs and the withheld counter release strands forever -
+ * the guest's tracker word stops short (evtest3 cudaStreamSynchronize(st[2])
+ * hangs; opdef=531/8/0 static, no EXPIRED/LOSS lines). Sweep expired slots
+ * once per service iteration: deliver forward-only to the still-live owner
+ * channel, count a LOSS when the owner is gone.
+ */
+static void svc_opdef_expire_sweep(void)
+{
+	/*
+	 * struct svc_channel is >25 KB (per-slot op bitmaps); a stack copy
+	 * overflows the 16 KB kernel stack. Only the single service thread
+	 * calls this, so a static scratch is safe (same pattern as
+	 * svc_handle_chan_dump / svc_pump_channels).
+	 */
+	static struct svc_channel ch_scratch;
+	int i;
+
+	for (i = 0; i < SVC_OPDEF_SLOTS; i++) {
+		struct svc_opdef_ent *e = &svc_opdef[i];
+		struct svc_channel *local = NULL;
+		unsigned long flags;
+		int c;
+
+		if (!e->active || !e->n)
+			continue;
+		if (time_before(jiffies, e->since +
+				msecs_to_jiffies(SVC_OPDEF_HOLD_MS)))
+			continue;
+		spin_lock_irqsave(&svc.chan_lock, flags);
+		for (c = 0; c < SVC_MAX_CHANNELS; c++)
+			if (svc.chan[c].valid &&
+			    svc.chan[c].vgpu == e->vgpu &&
+			    svc.chan[c].h_object == e->h_object) {
+				ch_scratch = svc.chan[c];
+				local = &ch_scratch;
+				break;
+			}
+		spin_unlock_irqrestore(&svc.chan_lock, flags);
+		if (local) {
+			svc_fwd_rel_apply_all(e->vgpu, local, e->rel, e->n);
+			svc.opdef_expired++;
+			pr_info_ratelimited("sema-opdef: SWEEP-EXPIRED h=0x%x entry=%u n=%u applied forward-only (entry parked past the hold window)\n",
+					    e->h_object, e->entry_idx, e->n);
+		} else {
+			svc.opdef_drop_lost += e->n;
+			pr_info_ratelimited("sema-opdef: SWEEP-LOSS h=0x%x entry=%u n=%u owner channel gone (releases dropped)\n",
+					    e->h_object, e->entry_idx, e->n);
+		}
+		svc_opdef_clear(e);
+	}
+}
+
+/*
  * Proof line for a held record set resolved without a successful forwarded
  * copy (failure/timeout/abort): the releases below are applied directly, so
  * none is lost to a hardware copy that never ran. Called with svc_fwd_hold
@@ -17090,6 +17147,14 @@ static int svc_thread(void *arg)
 		}
 
 		/*
+		 * STALL FIX (r189): expire withheld release sets whose entry
+		 * stopped being re-decoded (the deadline flush is only
+		 * reachable from the entry path). Forward-only delivery to a
+		 * live owner; counted LOSS otherwise. Cheap: 8 slots.
+		 */
+		svc_opdef_expire_sweep();
+
+		/*
 		 * Service EVERY open mdev each pass, in list order. With one
 		 * mdev this is exactly the historical single-vgpu pass; with
 		 * two, both guests get a pump turn per iteration (hardware