diff options
| -rw-r--r-- | kernel/gsp_sema_service.c | 67 |
1 files changed, 66 insertions, 1 deletions
diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 4ab450e..93aad52 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -815,6 +815,8 @@ struct sema_svc { u64 opdef_held; u64 opdef_applied; u64 opdef_expired; + u64 opdef_delivered_late; /* stall fix: set delivered at consume */ + u64 opdef_drop_lost; /* withheld releases dropped unapplied */ u64 vidmem_writes; /* semaphores written in vidmem */ u64 gpfifo_log; /* channel-state log budget */ u64 gpfifo_entry_log; /* per-entry log budget */ @@ -1349,6 +1351,7 @@ static bool svc_opdef_release(struct nv_vgpu_device *vgpu, static void svc_opdef_drop(struct nv_vgpu_device *vgpu, struct svc_channel *ch, u32 entry_idx); static void svc_opdef_reset_all(void); +static void svc_opdef_finish_ch(struct svc_channel *ch, u32 idx); void gsp_sema_service_channel_del(struct nv_vgpu_device *vgpu, u32 h_object) { @@ -11230,6 +11233,18 @@ static bool svc_entry_want_reached(struct nv_vgpu_device *vgpu, static void svc_replay_note(struct svc_channel *ch, u32 idx, u32 put) { (void)put; + /* + * STALL FIX (entry=82/pay=84 class, r186): the entry leaving the + * decode path must never strand a release set that svc_opdef_release() + * withheld for it while an operand was pending. svc_opdef_flush() is + * only reachable from the entry decode path; once the entry is + * consumed (e.g. via the entry_want shortcut) the withheld set sat + * until a reset/drop cleared it WITHOUT applying, and the guest's + * tracker word stopped short of the waiter's value (evtest3 + * cudaStreamSynchronize(st[2]) hung with starved=0). Deliver the set + * here, forward-only, before the slot is marked done. + */ + svc_opdef_finish_ch(ch, idx); if (idx < SVC_GPFIFO_MAX_BITS) set_bit(idx, ch->slot_done_epoch); } @@ -11845,6 +11860,40 @@ static void svc_opdef_apply(struct nv_vgpu_device *vgpu, } /* + * STALL FIX (r186 entry=82/pay=84): called from svc_replay_note() on every + * entry-consume path. If a release set was withheld for this entry while an + * operand was pending and the entry is now leaving the decode path (so the + * 10 s flush can never run for it), deliver the set forward-only and clear + * the slot. A drop without delivery is the stall class: the guest word + * stops short of the value its tracker waits for. + */ +static void svc_opdef_finish_ch(struct svc_channel *ch, u32 idx) +{ + u32 i; + + if (!ch || idx == SVC_ENTRY_NONE) + return; + for (i = 0; i < SVC_OPDEF_SLOTS; i++) { + struct svc_opdef_ent *e = &svc_opdef[i]; + + if (!e->active || e->h_object != ch->h_object || + e->vgpu != ch->vgpu || e->entry_idx != idx) + continue; + if (e->n) { + svc_fwd_rel_apply_all(e->vgpu, ch, e->rel, e->n); + svc.opdef_delivered_late++; + if (svc_opdef_log < 64) { + svc_opdef_log++; + pr_info("sema-opdef: LATE-FLUSH h=0x%x entry=%u n=%u (entry consumed before deadline; releases delivered)\n", + ch->h_object, idx, e->n); + } + } + svc_opdef_clear(e); + return; + } +} + +/* * Release-site gate, called from svc_decode_stream just before a decoded * release would be applied inline. Returns true when the caller must NOT * apply its own release now: @@ -11994,6 +12043,16 @@ static void svc_opdef_drop_stale(struct nv_vgpu_device *vgpu, /* A slot of THIS pump's guest whose channel is gone from the pass snapshot: * drop it. Slots of other guests are never touched (their channels are not in * this snapshot by construction). */ +static void svc_opdef_drop_counted(struct svc_opdef_ent *e) +{ + if (e->n) { + svc.opdef_drop_lost += e->n; + pr_info_ratelimited("sema-opdef: LOSS entry h=0x%x idx=%u dropped with %u undelivered release(s)\n", + e->h_object, e->entry_idx, e->n); + } + svc_opdef_clear(e); +} + static void svc_opdef_drop_gone(struct nv_vgpu_device *vgpu, const struct svc_channel *local, int n) { @@ -12009,12 +12068,18 @@ static void svc_opdef_drop_gone(struct nv_vgpu_device *vgpu, local[k].h_object == e->h_object) break; if (k == n) - svc_opdef_clear(e); + svc_opdef_drop_counted(e); } } static void svc_opdef_reset_all(void) { + int i; + + /* STALL FIX: a non-applying clear is a silent release loss - count it. */ + for (i = 0; i < SVC_OPDEF_SLOTS; i++) + if (svc_opdef[i].active && svc_opdef[i].n) + svc_opdef_drop_counted(&svc_opdef[i]); memset(svc_opdef, 0, sizeof(svc_opdef)); } |