From 44d589d032c660a16b78758e6680868b04d5344c Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Sat, 17 Oct 2026 11:14:53 -0400 Subject: 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) --- kernel/gsp_sema_service.c | 65 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 65 insertions(+) (limited to 'kernel') 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 @@ -12120,6 +12120,63 @@ static void svc_opdef_reset_all(void) memset(svc_opdef, 0, sizeof(svc_opdef)); } +/* + * 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 @@ -17089,6 +17146,14 @@ static int svc_thread(void *arg) complete_all(&svc_sess_flush_done); } + /* + * 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 -- cgit v1.2.3