vgsp-sema: fix rank-3 completion-word pollution - publish/autoreg target the channel's own (descending) counter slot

The client zips GR channels in ascending handle order onto descending
16-byte counter slots: rank r's pushes release slot fec0+(19-r)*0x10,
which is the completion word its stream sync waits on (verified for all
20 ranks across r196/r199/r200). F1 (eea5ea93) registered the ASCENDING
slot fec0+r*0x10 as the publish target, so every channel's ring PUT was
written into the paired rank (19-r)'s waited word. The busy ctx-init
channel rank 16 (h=0x5c000059, PUT~94) then overwrote rank 3's word
0x20400ffc0 from 1 to 94 (witness r199/r200: 'pub-adv h=0x5c000059
pa=...fc0 val=94 cur=1' after rank-3's own release wrote 1); the rank-3
stream's own release (payload 1) was forward-only-skipped and its wait
never completed: evtest3 cudaStreamSynchronize(st[2]) hangs forever
(sync(0)/sync(1) pass because ranks 1/2 pair with the idle channels
18/17) and victory_streams_events hangs on its rank-3 stream (r198:
launches on h=1f/023/027, 302 s window). Restore the descending autoreg
documented at SVC_GR_COUNTER_TOTAL and make svc_counter_publish target
release_pa ?: counter_pa (the same selection svc_counter_release_apply
uses), so the PUT catch-up can never land in another channel's waited
word; release_pa keeps the push-named truth when the autoreg derive
fails.

0 files changed, 0 insertions, 0 deletions