Function hipStreamBeginCapture_common

hipamd/src/hip_graph.cpp:1100–1135

}

hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode,
                                        hipGraph_t graph = nullptr) {
  getStreamPerThread(stream);
  // 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);
  // It can be initiated if the stream is not already in capture mode
  if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
    return hipErrorIllegalState;
  }
  if (graph == nullptr) {
    s->SetCaptureGraph(new hip::Graph(s->GetDevice()));
  } else {
    s->SetCaptureGraph(reinterpret_cast<hip::Graph*>(graph));
  }
  s->SetCaptureId();
  s->SetCaptureMode(mode);
  s->SetOriginStream();
  if (mode != hipStreamCaptureModeRelaxed) {
    hip::tls.capture_streams_.push_back(s);
  }
  if (mode == hipStreamCaptureModeGlobal) {
    amd::ScopedLock lock(g_captureStreamsLock);
    g_captureStreams.push_back(s);
  }
  {
    amd::ScopedLock lock(g_streamSetLock);
    g_allCapturingStreams.insert(s);
  }
  return hipSuccess;
}

hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode) {
  HIP_INIT_API(hipStreamBeginCapture, stream, mode);