From 032f17c044bd6b6ba8651e2d0b9ded6231832550 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Fri, 23 Oct 2026 17:37:24 -0400 Subject: vgsp-sema r567: a transfer-shaped push with no op notes its destination and stays owed The r569 mail failure's missing half: the b H2D's COPY op needs have_in (the staging upload), which was not resolved on the early passes, so the push decoded NO op; the release was still applied (the client's cudaMemcpy(b) returned at 318935) and the slot was marked complete, while the copy itself only executed at 318952 (17 s later) - the launch's copy-in (318945.7) read b = 0 and the engine computed c = a + 0 (readback src16 = a's values). The r541 guard had nothing to defer on (no op, no dstpend). A trusted, allow_ops, transfer-shaped push that decodes no op with a device destination now: - notes the destination via svc_dstpend_note (the launch guard defers until the copy runs or the 10 s window expires), and - sets svc.op_owed (the slot stays incomplete, so a later fresh pass runs the copy instead of replay-suppressing it forever). No hold is armed, so a push whose trust never arrives cannot wedge the pump; the guard provides the ordering, exactly like the r541 design. --- kernel/gsp_sema_service.c | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) (limited to 'kernel') diff --git a/kernel/gsp_sema_service.c b/kernel/gsp_sema_service.c index 80c858b..0fc37cd 100644 --- a/kernel/gsp_sema_service.c +++ b/kernel/gsp_sema_service.c @@ -11179,6 +11179,35 @@ static void svc_decode_stream(struct nv_vgpu_device *vgpu, const u8 *buf, } } + /* + * r567: a trusted, transfer-shaped push that decoded + * NO op this pass (its operands are not resolved yet: + * the r569 mail b-copy needed have_in, which the + * staging upload had not produced) must (a) keep its + * slot incomplete so a later pass runs the copy + * instead of replay-suppressing it forever (op_owed, + * the r540 rule), and (b) make its stale destination + * visible to the r541 launch guard, which defers any + * launch reading that page until the copy runs or the + * dstpend window expires. The release still applies + * (no hold is armed: a push whose trust never arrives + * must not wedge the pump) - the guard, not the + * release, provides the ordering. + */ + if (allow_ops && transfer != 0 && trusted && ch && + n_ops == n_ops_before && st.have_out && + st.offset_out >= SVC_VAS_FLOOR) { + svc_dstpend_note(st.offset_out); + svc.op_owed = 1; + if (svc.op_owed_log < 64) { + svc.op_owed_log++; + pr_info("sema-pend: OWED-PUSH h=0x%x push=0x%llx out=0x%llx line=%u (transfer-shaped, no op; dst noted for the launch guard)\n", + ch->h_object, push_va, + st.offset_out, + st.line_length); + } + } + /* * r365 FILL-RAW: the fill class that still decodes no * op (window VA, no remap const decoded) - dump the -- cgit v1.2.3