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.

1 files changed, 29 insertions, 0 deletions

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
@@ -11180,6 +11180,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
 			 * raw push and decode its method stream so the true