fix(rccl): don't record doneEvent on a possibly-destroyed stream

[pvallem]

Motivation

PR #6130 moved the hipEventRecord(comm->doneEvent, ...) out of ncclLaunchKernel and
into ncclLaunchPrepare's stream-change branch, recording it on comm->lastStream
rather than on the stream being launched.

That turned lastStream from a compare-only value into a dereferenced handle. RCCL
never learns when the application destroys a stream — HIP has no destruction notification
and no safe validity probe — and lastStreamValid was only ever set true, never
cleared. So an application that destroys the stream a communicator last launched on, and
then issues a collective on a different stream, hits a use-after-destroy.

It surfaces two ways:

• ncclUnhandledCudaError from HIP failure: 'invalid resource handle'
• a segfault, when the freed handle has been recycled onto an unrelated live stream,
  producing a bogus ordering edge

This is not theoretical — it reproduces on a real QA workload (see Test Result).

Technical Details

Record doneEvent at launch time again, on launchStream, which is live by construction:
the application has just handed it to us. ncclLaunchPrepare now only waits on the event,
never records it.

lastStream + lastStreamValid collapse into a single uintptr_t lastStreamTag, produced
by ncclStreamTag(s) = (uintptr_t)s + 1 so 0 stays free as the "no launch yet" sentinel
and a prior launch on the default stream remains distinguishable (the case #5800 fixed).
One word also removes any torn-read question between a handle and a separate flag.

Comparing a stale tag is safe in both directions:

• handle recycled -> compares equal -> no edge installed. Safe: hipStreamDestroy
  defers handle reuse until the stream's work completes, so a recycled handle proves the
  prior kernel already finished.
• handle not recycled -> compares unequal -> wait on doneEvent, recorded on that
  stream while it was alive. The event is comm-owned and outlives the stream.

Typing it as uintptr_t makes the unsafe use unrepresentable rather than relying on a
comment to prevent recurrence.

JIRA ID

ROCM-29677

Test Plan

• Targeted repro: allreduce on stream A -> hipStreamDestroy(A) -> create decoy streams to
  force handle divergence -> allreduce on a new stream B, same comm.
• New unit test DoneEventOrdering.StreamDestroyedThenSwitch, run against both a patched
  and an unpatched RCCL to confirm it actually catches the defect.
• Real workload: AIGQA concurrent_collectives weekly (8x MI350X, 1000 iterations, 256 MB),
  which creates 6 streams per datatype phase and destroys them while the 6 communicators
  span all three phases.
• Version sweep across every ROCm install on the test node, HIP held constant and only
  librccl swapped via LD_PRELOAD.
• rccl-tests all_reduce_perf A/B for the performance impact.

Test Result

Targeted repro — same binary, same HIP (/opt/rocm/lib/libamdhip64.so.7), only librccl swapped:

RCCL	Result
stock 2.30.4 (has #6130)	invalid resource handle, exit 3
this branch	phase2 OK -- NO CRASH, exit 0

New unit test — validated in both directions:

library	StreamDestroyedThenSwitch
this branch	PASSED (23.1 s)
stock 2.30.4	FAILED — HIP failure: 'invalid resource handle'

Real workload, 4 runs per arm, alternating:

arm	result
this branch	0 failures / 4 runs — all 3 phases, ~675 s
stock 2.30.4	3 failures / 4 runs — rc 139/139/134, dying after 1 of 3 phases at ~250 s

Note the defect is intermittent at the workload level (it only fires when the replacement
stream draws a different handle than the destroyed one), which is why the deterministic
repro above uses decoy streams to force the condition.

Affected versions — the bug is present in every RCCL from 2.28.9 onward:

ROCm tree	RCCL	Verdict
therock-dist 10.0.0rc3	2.30.4	FAIL
rocm-7.14.0a20260612	2.29.7	FAIL
rocm-7.14.0-1417	2.28.9	FAIL
rocm-from-therock	2.28.9	FAIL
rocm-7.12.0	2.28.3	PASS
rocm-7.12.0-pass	2.28.3	PASS

Consistent with the dates: #6130 merged 2026-05-15; the 2.28.9 bump (46931872a0) landed
2026-05-18. ROCm 10.0.0rc3 ships this defect.

Performance Impact

This restores the per-launch hipEventRecord that #6130 removed, and that cost is real.
all_reduce_perf -b 8 -e 8K -f 8 -n 1000 -w 200, idle node, median of 5 alternating reps
per arm. All 20 runs exited 0 with #wrong = 0 on every size row; none discarded.

gpus	size	stock us	this branch us	delta us	delta/G us	change
8	8	38.36	57.07	18.71	2.34	+49%
8	64	37.87	62.62	24.75	3.09	+65%
8	512	38.91	63.09	24.18	3.02	+62%
8	4096	38.77	61.71	22.94	2.87	+59%
2	8	8.94	14.43	5.49	2.75	+61%
2	64	8.40	13.68	5.28	2.64	+63%
2	512	8.30	13.41	5.11	2.56	+62%
2	4096	8.17	13.78	5.61	2.81	+69%

rccl-tests reports one time per op while issuing -g N ncclAllReduce calls, so the per-op
delta scales with N; delta/G is the implied per-launch cost. It lands in the same
2.5-3.1 us band at both GPU counts, averaging 2.76 us per launch — matching the figure
#6130 cited.

This is a knowing tradeoff: a crash is being exchanged for a latency regression that
#6130 was created to remove. Recovering it is tracked as an immediate follow-up — fusing
the record into the launch via hipExtModuleLaunchKernel, which accepts both the packed
extra buffer this code needs and a stopEvent, exactly as RCCL did via
hipExtLaunchKernel's 8th argument before #3741. That restores the ordering edge at zero
extra host calls. Reviewers who care about the #6130 workloads (ROCM-21756, ROCM-24157)
should track that follow-up.

Submission Checklist

• Look over the contributing guidelines at
  https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

---

🤖 Generated with Claude Code

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/rccl/src/enqueue.cc, projects/rccl/src/include/comm.h, projects/rccl/src/init.cc, projects/rccl/test/DoneEventOrderingTests.cpp; no test file found
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?