vgsp: m174 deliver withheld COUNTER releases at the fwd-ownership handoff (svc_opdef_drop cleared the entry's withheld set blindly - the r186/r187 entry=82 va=0x20400fef0 pay=84 record never rode the forwarded candidate, so st[2]'s completion counter froze right after the launches and evtest3 cudaStreamSynchronize(st[2]) hung with no LATE-FLUSH/LOSS/EXPIRED; counter records now apply forward-only before the clear, reductions skipped, remaining drops counted+logged)

1 files changed, 39 insertions, 2 deletions

diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c
index 93aad52..7b63c0f 100644
--- a/kernel/gsp_sema_service.c
+++ b/kernel/gsp_sema_service.c
@@ -12014,8 +12014,45 @@ static void svc_opdef_drop(struct nv_vgpu_device *vgpu, struct svc_channel *ch,
 {
 	struct svc_opdef_ent *e = svc_opdef_find(vgpu, ch->h_object, entry_idx);
 
-	if (e)
-		svc_opdef_clear(e);
+	if (!e)
+		return;
+	/*
+	 * STALL FIX (r186/r187 entry=82, va=0x20400fef0 pay=84): when the
+	 * forwarding machinery takes ownership of an entry, the withheld set
+	 * may NOT ride the forwarded candidate's rel[] (svc_opdef_release()
+	 * returned before the release was appended), so a blind clear strands
+	 * the guest's tracker word with NO LATE-FLUSH/LOSS/EXPIRED ever
+	 * logged (crel freezes right after the launches; evtest3
+	 * cudaStreamSynchronize(st[2]) hangs). Counter releases are
+	 * monotonic completion signals: apply them forward-only here - a
+	 * record at/ahead of the word is a no-op, so the fwd edge cannot
+	 * double-apply and an additive reduction is skipped.
+	 */
+	{
+		u8 i;
+		u32 delivered = 0;
+
+		for (i = 0; i < e->n; i++) {
+			if (!svc_counter_va(e->rel[i].va))
+				continue;
+			if (e->rel[i].kind == SVC_FWD_REL_906F &&
+			    e->rel[i].op == HOST_SEM_OP_REDUCTION)
+				continue;	/* additive: not idempotent */
+			svc_fwd_rel_apply_all(vgpu, ch, &e->rel[i], 1);
+			svc.opdef_delivered_late++;
+			delivered++;
+			pr_info_ratelimited("sema-opdef: DROP-APPLY h=0x%x entry=%u va=0x%llx pay=%u (counter release delivered at fwd-ownership handoff)\n",
+					    ch->h_object, entry_idx,
+					    e->rel[i].va, e->rel[i].payload);
+		}
+		if (delivered < e->n) {
+			svc.opdef_drop_lost += (u32)(e->n - delivered);
+			pr_info_ratelimited("sema-opdef: LOSS h=0x%x entry=%u %u of %u withheld release(s) dropped at fwd-ownership handoff\n",
+					    ch->h_object, entry_idx,
+					    (u32)(e->n - delivered), e->n);
+		}
+	}
+	svc_opdef_clear(e);
 }
 
 /*