From 81c114202ceb152d9039c42b36edad6183eecab2 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Fri, 23 Oct 2026 10:44:50 -0400 Subject: vgsp-sema r541: hold a GR forward whose data page has an unexecuted device write Live root cause of the smiley/streams2/streams_events readback lottery (r539): smiley pass 2 verified a2 + b1 and pass 4 verified a4 + b2 - each even pass's launch read the PREVIOUS even pass's b page. The b H2D's release completed the client's synchronous cudaMemcpy while the copy op itself was still held (b's svc_copy executed at 294026.818; launch 2's engine read at ~294026.3), so the client legitimately submitted the launch against a shadow page the H2D had not written yet. Passes 1/3 looked 'correct' only because passes 2 and 3 share the same b formula. Fix, as recommended in CONT.100/94: a device-dst copy/fill that fails to execute this pass (rc=-ENOENT) records its destination VA page in svc_dstpend[] with a 10 s stamp. svc_gr_request, right after the fill and before any submit, checks the launch's filled data pages against the table (svc_gr_dstpend_overlap); a hit returns the busy-defer verdict, so the pump holds the entry in its bounded retry window and re-captures once the copy has run - the record clears on the copy's success. The 10 s window restores the historical forward when the copy never lands, so the guard can only delay a launch, never hang it. --- kernel/gsp_sema_service.c | 149 +++++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 147 insertions(+), 2 deletions(-) diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index d3ac29c..af2d056 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -1105,6 +1105,9 @@ struct sema_svc { u64 fwd_gr_fallback; u64 fwd_gr_refused; u64 fwd_gr_defer; /* m150: held for data-arg retry */ + u64 fwd_gr_dstpend; /* r541: held for an unexecuted + * device write on a data page */ + u32 fwd_gr_dstpend_log; u64 fwd_gr_timeout; u64 fwd_gr_fault; u64 fwd_gr_bridge; @@ -7844,6 +7847,84 @@ static void svc_opfail_clear(u64 va) svc_opfail[i].va = 0; } +/* + * r541: device-destination copies that decoded but have not executed yet. + * + * smiley live evidence (r539): passes 2 and 4 verified a + (previous b) - + * the client's cudaMemcpy(b) completed (its release landed) while the H2D + * op was still held, the client submitted the launch, the GR forward read + * the shadow page BEFORE the H2D executed (b's copy ran at 294026.818, + * launch 2's engine read at ~294026.3), and the engine computed with the + * PREVIOUS pass's b (bad=65024). The same shape is the streams2 / + * streams_events readback lottery: a device page whose writer has not run. + * + * The guard: when a device-dst copy (or fill) op fails to execute this + * pass, its destination VA page is recorded here with a stamp. A GR launch + * whose data pages overlap a fresh record is NOT forwarded (the pump's + * bounded busy-defer retries it; the record clears when the op finally + * executes or the SVC_DSTPEND_MS window expires, at which point the + * historical forward stands). Keyed by VA page: the page identity is what + * the launch operand and the copy share, and it works even when neither + * side resolves. + */ +#define SVC_DSTPEND_MAX 64 +#define SVC_DSTPEND_MS 10000u +static struct { + u64 va_page; + unsigned long since; +} svc_dstpend[SVC_DSTPEND_MAX]; +static u32 svc_dstpend_n; +static u32 svc_dstpend_log; + +static void svc_dstpend_note(u64 va) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + if (!page) + return; + for (i = 0; i < svc_dstpend_n; i++) { + if (svc_dstpend[i].va_page == page) { + svc_dstpend[i].since = jiffies; + return; + } + } + if (svc_dstpend_n < SVC_DSTPEND_MAX) + i = svc_dstpend_n++; + else + i = (u32)(jiffies % SVC_DSTPEND_MAX); + svc_dstpend[i].va_page = page; + svc_dstpend[i].since = jiffies; +} + +static void svc_dstpend_clear(u64 va) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + for (i = 0; i < svc_dstpend_n; i++) + if (svc_dstpend[i].va_page == page) + svc_dstpend[i].va_page = 0; +} + +static bool svc_dstpend_hit(u64 va) +{ + u64 page = va & ~(u64)(SEMA_PAGE_SIZE - 1); + u32 i; + + for (i = 0; i < svc_dstpend_n; i++) { + if (svc_dstpend[i].va_page != page) + continue; + if (time_after_eq(jiffies, svc_dstpend[i].since + + msecs_to_jiffies(SVC_DSTPEND_MS))) { + svc_dstpend[i].va_page = 0; + return false; + } + return true; + } + return false; +} + static bool svc_opfail_latched(u64 va) { u32 i; @@ -8134,9 +8215,24 @@ svc_wait_done: svc.memsets_done++; /* r507a: a success re-arms the failure window. */ svc_opfail_clear(m->dst); - } else if (rc == -ENOENT) + /* r541: the device page's writer has run; lift the + * launch-ordering guard for it. */ + svc_dstpend_clear(m->dst); + } else if (rc == -ENOENT) { + /* + * r541: a copy/fill that cannot run yet leaves its + * target page stale; a GR launch must not read it + * (see svc_dstpend_note). Record every unexecuted + * target - the VA-page key keeps the guard narrow. + */ + svc_dstpend_note(m->dst); + if (svc_dstpend_log < 64) { + svc_dstpend_log++; + pr_info("sema-svc: DSTPEND note va=0x%llx type=%u ap=%u (device write unexecuted; GR launches touching this page hold)\n", + m->dst, m->type, m->aperture); + } svc.memsets_nomap++; - else + } else svc.memsets_fail++; /* * ISSUE-0002 (r71/r72): a failed LARGE op leaves work owed on @@ -11102,6 +11198,14 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, st.line_length && st.have_out && st.have_in && !DMA_REMAP_ENABLE(dma) && st.offset_out >= SVC_VAS_FLOOR) { + /* + * r541: the denied device-dst copy has NOT + * run; whatever release edge later consumes + * this entry, a GR launch must not read the + * target page until the copy executes (see + * svc_dstpend_note). + */ + svc_dstpend_note(st.offset_out); /* * r314 (CONT.57 lead 2): capture the staging * source on the DENIED pass too. This is the @@ -20388,6 +20492,21 @@ static void svc_gr_req_bridge_fallback(struct nv_vgpu_device *vgpu) */ static void svc_handle_map_add(void); +/* + * r541: does any of the filled launch's data pages carry an unexecuted + * device write (see svc_dstpend_note)? Runs on the service thread right + * after the fill, so svc_fwd_req.gr_data holds this launch's records. + */ +static bool svc_gr_dstpend_overlap(void) +{ + u32 i; + + for (i = 0; i < svc_fwd_req.gr_n_data; i++) + if (svc_dstpend_hit(svc_fwd_req.gr_data[i].va)) + return true; + return false; +} + static bool svc_gr_request(struct nv_vgpu_device *vgpu, const struct svc_channel *ch) { @@ -20485,6 +20604,32 @@ static bool svc_gr_request(struct nv_vgpu_device *vgpu, svc_gr_bridge_fallback(vgpu); return false; } + /* + * r541: launch-ordering guard. The fill's data pages are what the + * engine will read; if any of them is the target of a device copy + * that decoded but has NOT executed yet (svc_dstpend_note - the + * smiley/streams2 class: the client's memcpy completed on its + * release while the H2D op was still held, and the launch read the + * PREVIOUS page), do not forward: return the busy-defer verdict so + * the pump holds the entry through its bounded retry window and + * re-captures once the copy has run (or the dstpend window expired, + * which restores the historical forward). Bounded by the same + * busy-defer + gr_hold machinery as every other busy refusal, so it + * can never hang the launch. + */ + if (svc_gr_dstpend_overlap()) { + clear_bit_unlock(SVC_FWD_INFLIGHT, &svc_fwd_state); + mutex_unlock(&svc_fwd_lock); + svc_gr_busy_defer = true; + svc.fwd_gr_refused++; + svc.fwd_gr_dstpend++; + if (svc.fwd_gr_dstpend_log < 32) { + svc.fwd_gr_dstpend_log++; + pr_info("vgsp-work: GR forward deferred h=0x%x push=0x%llx (data page has an unexecuted device write; bounded retry)\n", + ch->h_object, svc_gr_cand.push_va); + } + return false; + } /* * m150: the fill found high argument pointers that no source can * translate (the second-client zeros class). Do NOT forward a launch -- cgit v1.2.3