diff options
| author | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-10-23 17:37:24 -0400 |
|---|---|---|
| committer | Arthur Hanson Rasmusson <arthur@vgpu.io> | 2026-10-23 17:37:24 -0400 |
| commit | 032f17c044bd6b6ba8651e2d0b9ded6231832550 (patch) | |
| tree | 5ec55d5cb514cc3d697ed661500791069f921bc3 /tools/cuda-ab/README.md | |
| parent | 3ba89790de6bde0efec10c25be5e5254fe10014e (diff) | |
| download | open-iov-032f17c044bd6b6ba8651e2d0b9ded6231832550.tar.gz open-iov-032f17c044bd6b6ba8651e2d0b9ded6231832550.zip | |
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.
Diffstat (limited to 'tools/cuda-ab/README.md')
0 files changed, 0 insertions, 0 deletions
