Take into account thread-local mode in checks for valid capture

[iassiour]

Motivation

This PR addresses ROCm/hip#3876

Technical Details

In the case where one thread captures a stream in global capture mode, an action like hipEventQuery on a different thread for a different stream under thread-local capture mode currently returns hipErrorStreamCaptureUnsupported.
However In thread-local capture mode sequences in other threads should be ignored.

JIRA ID

Test Plan

Test Result

Submission Checklist

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

[Copilot]

Pull request overview

This PR refactors the StreamCaptureOngoing function to properly handle thread-local capture mode when validating stream captures, addressing hip-issue-3876. The changes aim to respect capture mode boundaries (Global, ThreadLocal, Relaxed) when determining if operations on a stream should be allowed or if captures should be invalidated.

• Restructured the capture validation logic to check both stream capture status and current thread's capture mode
• Removed a getStreamPerThread call from hipStreamBeginCapture_common that may have been redundant
• Added clearer code structure with early returns and a helper lambda for stream invalidation

Reviewed changes

Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.

File	Description
projects/clr/hipamd/src/hip_stream.cpp	Refactored StreamCaptureOngoing to properly handle thread-local vs global capture modes, with improved logic flow and clearer handling of different capture status/mode combinations
projects/clr/hipamd/src/hip_graph.cpp	Removed getStreamPerThread call from hipStreamBeginCapture_common, potentially for consistency with the spt/non-spt API pattern

Comments suppressed due to low confidence (1)

projects/clr/hipamd/src/hip_graph.cpp:1116

• Potential inconsistency: The getStreamPerThread(stream) call was removed, but other similar _common functions like hipStreamIsCapturing_common (line 1068) still have this call.

Without this call, if hipStreamBeginCapture (non-spt variant) is called with hipStreamPerThread as the stream parameter, the stream won't be converted to the actual per-thread stream. It will pass the nullptr/legacy check (line 1110) and then be incorrectly cast to hip::Stream* at line 1116.

Consider either:

1. Adding back getStreamPerThread(stream) at the beginning of this function for consistency and safety
2. Adding an explicit check for hipStreamPerThread and returning an appropriate error
3. Documenting why this function should never receive hipStreamPerThread and ensuring all call sites respect this

hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode,
                                        hipGraph_t graph = nullptr) {
  // capture cannot be initiated on legacy stream
  if (stream == nullptr || stream == hipStreamLegacy) {
    return hipErrorStreamCaptureUnsupported;
  }
  if (mode < hipStreamCaptureModeGlobal || mode > hipStreamCaptureModeRelaxed) {
    return hipErrorInvalidValue;
  }
  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.