vgsp-sema: fix rank-3 completion-word pollution - publish/autoreg target the channel's own (descending) counter slot

The client zips GR channels in ascending handle order onto descending
16-byte counter slots: rank r's pushes release slot fec0+(19-r)*0x10,
which is the completion word its stream sync waits on (verified for all
20 ranks across r196/r199/r200). F1 (eea5ea93) registered the ASCENDING
slot fec0+r*0x10 as the publish target, so every channel's ring PUT was
written into the paired rank (19-r)'s waited word. The busy ctx-init
channel rank 16 (h=0x5c000059, PUT~94) then overwrote rank 3's word
0x20400ffc0 from 1 to 94 (witness r199/r200: 'pub-adv h=0x5c000059
pa=...fc0 val=94 cur=1' after rank-3's own release wrote 1); the rank-3
stream's own release (payload 1) was forward-only-skipped and its wait
never completed: evtest3 cudaStreamSynchronize(st[2]) hangs forever
(sync(0)/sync(1) pass because ranks 1/2 pair with the idle channels
18/17) and victory_streams_events hangs on its rank-3 stream (r198:
launches on h=1f/023/027, 302 s window). Restore the descending autoreg
documented at SVC_GR_COUNTER_TOTAL and make svc_counter_publish target
release_pa ?: counter_pa (the same selection svc_counter_release_apply
uses), so the PUT catch-up can never land in another channel's waited
word; release_pa keeps the push-named truth when the autoreg derive
fails.

1 files changed, 50 insertions, 24 deletions

diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c
index 01f94cf..dd7a5e1 100644
--- a/kernel/gsp_sema_service.c
+++ b/kernel/gsp_sema_service.c
@@ -6089,14 +6089,22 @@ static int svc_counter_derive(struct nv_vgpu_device *vgpu, u64 va, u64 *pa_out)
 
 /*
  * Auto-register a GR channel's completion-counter slot from the observed
- * layout, without waiting for a readable push: the counter records form a
- * 16-byte-strided array in the client's counter page, allocated in channel
- * order from the low end (SA-COUNTER FORENSICS-grv7-counter: the ascending
- * slot fec0 + r*0x10 is the polled completion word; the push's REPORT_SEM
- * release writes a DIFFERENT (descending) family and is tracked separately
- * in release_pa). The slot PA comes from a host-verified pin or from the
- * in-kernel structural derivation (svc_counter_derive), so the write target
- * is a page the client actually polls.
+ * layout, without waiting for a readable push: the client zips its GR
+ * channels in ascending handle order onto DESCENDING 16-byte slots, so the
+ * rank-r channel's own completion word - the slot its pushes' REPORT_SEM
+ * release names, verified for all 20 ranks - is
+ * 0x20400fec0 + (SVC_GR_COUNTER_TOTAL-1-r)*0x10. The ascending formula
+ * (F1/eea5ea93) registered the completion word of the PAIRED rank (19-r)
+ * instead, so svc_counter_publish wrote each channel's ring PUT into
+ * another channel's waited word: the busy ctx-init channel rank 16
+ * (h=0x5c000059, PUT~94) overwrote rank 3's completion word 0x20400ffc0
+ * from 1 to 94 (witness `pub-adv h=0x5c000059 pa=...fc0 val=94 cur=1`),
+ * after which the rank-3 stream's own release (payload 1) was
+ * forward-only-skipped and its sync never completed (evtest3
+ * cudaStreamSynchronize(st[2]); victory_streams_events). The slot PA comes
+ * from a host-verified pin or from the in-kernel structural derivation
+ * (svc_counter_derive), so the write target is a page the client actually
+ * polls.
  */
 static void svc_counter_autoreg(struct nv_vgpu_device *vgpu,
 				struct svc_channel *ch, int n_gr)
@@ -6109,7 +6117,8 @@ static void svc_counter_autoreg(struct nv_vgpu_device *vgpu,
 		return;
 	if (ch->gr_index >= n_gr)
 		return;
-	va = 0x20400fec0ULL + (u64)ch->gr_index * 0x10ULL;
+	va = 0x20400fec0ULL +
+	     (u64)(SVC_GR_COUNTER_TOTAL - 1 - ch->gr_index) * 0x10ULL;
 	if (ch->counter_pa) {
 		/*
 		 * FORENSICS-v2 anomaly 6: a slot derived in an earlier pass
@@ -6238,12 +6247,14 @@ static void svc_counter_autoreg(struct nv_vgpu_device *vgpu,
 /*
  * Learn a GR channel's REPORT_SEM release slot from a release VA the push
  * carried (0x20400fXX observed). This is the slot the release write
- * targets (release_pa); it is deliberately NOT the completion-counter slot
- * (counter_pa), which the publish advances and which the client polls
- * (SA-COUNTER FORENSICS-grv7-counter F1: green wrote both families of the
- * counter page). An injected pin is the authoritative source; on a miss the
- * slot is derived in-kernel (svc_counter_derive) so the release is authored
- * immediately rather than waiting for the external pin watcher.
+ * targets (release_pa) and - per the descending layout the autoreg above
+ * documents - the channel's own polled completion word, the same slot the
+ * autoreg registers as counter_pa; keeping it in its own field preserves
+ * the push-named truth even when the registration derive fails, and
+ * svc_counter_publish prefers it. An injected pin is the authoritative
+ * source; on a miss the slot is derived in-kernel (svc_counter_derive) so
+ * the release is authored immediately rather than waiting for the external
+ * pin watcher.
  */
 static void svc_counter_learn(struct nv_vgpu_device *vgpu, u32 h, u64 sem_va)
 {
@@ -11468,15 +11479,31 @@ static void svc_counter_publish(struct nv_vgpu_device *vgpu,
 {
 	u32 val;
 	u32 cur = 0;
+	/*
+	 * STALL FIX (rank-3 completion-word pollution): the publish must
+	 * target the channel's OWN completion word - the descending slot its
+	 * pushes release, which the autoreg above registers and which the
+	 * decoded releases write (release_pa). Writing the ring PUT into the
+	 * ascending rank slot (F1/eea5ea93) made every channel's publish land
+	 * in the paired rank (19-r)'s waited word: the busy ctx-init channel
+	 * rank 16 (h=0x5c000059, PUT~94) overwrote rank 3's word 0x20400ffc0
+	 * from 1 to 94 (`pub-adv h=0x5c000059 pa=...fc0 val=94 cur=1`), the
+	 * rank-3 stream's own release (payload 1) was then forward-only-
+	 * skipped, and its exact-value sync wait never completed (evtest3
+	 * cudaStreamSynchronize(st[2]); victory_streams_events). release_pa
+	 * is the push-named truth and wins when known; counter_pa stays the
+	 * fallback for a channel whose pushes are unreadable.
+	 */
+	u64 pa = ch->release_pa ? ch->release_pa : ch->counter_pa;
 
-	if (!ch->counter_pa)
+	if (!pa)
 		return;
 	svc.counter_pub_calls++;
 	if (gsp_sema_dbg && !ch->engine_type &&
 	    svc.counter_pub_dbg < 96) {
 		svc.counter_pub_dbg++;
 		pr_info("sema-chan: pub-info h=0x%x rank=%u pa=0x%llx put=%u last=%u\n",
-			ch->h_object, ch->gr_index, ch->counter_pa, put,
+			ch->h_object, ch->gr_index, pa, put,
 			ch->counter_last);
 	}
 	svc_counter_flush_pending(vgpu, ch);
@@ -11489,8 +11516,7 @@ static void svc_counter_publish(struct nv_vgpu_device *vgpu,
 	 * advanced the slot past the put position the client last submitted;
 	 * the client polls this u32, so never write it backwards.
 	 */
-	if (svc_read_ap(vgpu, ch->counter_pa, SVC_APERTURE_SYS, &cur,
-			sizeof(cur)))
+	if (svc_read_ap(vgpu, pa, SVC_APERTURE_SYS, &cur, sizeof(cur)))
 		return;
 	if (val <= cur) {
 		if (cur > ch->counter_last)
@@ -11504,16 +11530,15 @@ static void svc_counter_publish(struct nv_vgpu_device *vgpu,
 	 */
 	if (gsp_sema_dbg)
 		pr_info("sema-diag: pub-adv h=0x%x pa=0x%llx val=%u cur=%u last=%u\n",
-			ch->h_object, ch->counter_pa, val, cur,
+			ch->h_object, pa, val, cur,
 			ch->counter_last);
-	if (svc_write_ap(vgpu, ch->counter_pa, SVC_APERTURE_SYS,
-			 &val, sizeof(val)))
+	if (svc_write_ap(vgpu, pa, SVC_APERTURE_SYS, &val, sizeof(val)))
 		return;
 	ch->counter_last = val;
 	if (gsp_sema_dbg && svc.counter_pub_dbg < 48) {
 		svc.counter_pub_dbg++;
 		pr_info("sema-chan: counter publish h=0x%x pa=0x%llx val=%u (was %u)\n",
-			ch->h_object, ch->counter_pa, val, cur);
+			ch->h_object, pa, val, cur);
 	}
 }
 
@@ -15126,7 +15151,8 @@ static void svc_pump_channels(struct nv_vgpu_device *vgpu)
 			 * svc_counter_autoreg). Skip while a relay hold owns
 			 * the channel, exactly like the end-of-pass call.
 			 */
-			if (ch->counter_pa && ch->userd_pa &&
+			if ((ch->release_pa || ch->counter_pa) &&
+			    ch->userd_pa &&
 			    !(svc_fwd_hold.active &&
 			      svc_fwd_hold.vgpu == vgpu &&
 			      svc_fwd_hold.ch_handle == ch->h_object) &&