From ff69bcc903e42d1e7e0c8032632106108dda0d13 Mon Sep 17 00:00:00 2001 From: "Godavarthy Surya, Anusha" Date: Mon, 28 Apr 2025 15:09:11 +0530 Subject: [PATCH] SWDEV-469422 - Avoid using of hipStream_t in internal methods (#69) Change-Id: Ifd5362f371c846a88241927383cb95cf046548ef [ROCm/clr commit: fb92683d8617c44115658b602fa2e0158c122b76] --- projects/clr/hipamd/src/hip_event.cpp | 20 +++++-------- projects/clr/hipamd/src/hip_event.hpp | 10 +++---- projects/clr/hipamd/src/hip_event_ipc.cpp | 29 +++++++++---------- projects/clr/hipamd/src/hip_graph.cpp | 3 +- .../clr/hipamd/src/hip_graph_internal.cpp | 4 +-- .../clr/hipamd/src/hip_graph_internal.hpp | 5 ++-- projects/clr/hipamd/src/hip_mempool.cpp | 2 +- projects/clr/hipamd/src/hip_mempool_impl.cpp | 2 +- projects/clr/hipamd/src/hip_module.cpp | 2 +- projects/clr/hipamd/src/hip_stream.cpp | 4 +-- 10 files changed, 35 insertions(+), 46 deletions(-) diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 0fcb468976..62ed7c1dae 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -183,20 +183,17 @@ hipError_t Event::streamWaitCommand(amd::Command*& command, hip::Stream* stream) return hipSuccess; } // ================================================================================================ -hipError_t Event::streamWait(hipStream_t stream, uint flags) { - // Get the stream without any resolution - constexpr bool kWait = false; - hip::Stream* hip_stream = hip::getStream(stream, kWait); +hipError_t Event::streamWait(hip::Stream* stream, uint flags) { // Access to event_ object must be lock protected amd::ScopedLock lock(lock_); - if ((event_ == nullptr) || (event_->command().queue() == hip_stream) || ready()) { + if ((event_ == nullptr) || (event_->command().queue() == stream) || ready()) { return hipSuccess; } if (!event_->notifyCmdQueue()) { return hipErrorLaunchOutOfResources; } amd::Command* command; - hipError_t status = streamWaitCommand(command, hip_stream); + hipError_t status = streamWaitCommand(command, stream); if (status != hipSuccess) { return status; } @@ -226,7 +223,7 @@ hipError_t Event::recordCommand(amd::Command*& command, amd::HostQueue* stream, } // ================================================================================================ -hipError_t Event::enqueueRecordCommand(hipStream_t stream, amd::Command* command) { +hipError_t Event::enqueueRecordCommand(hip::Stream* stream, amd::Command* command) { command->enqueue(); if (event_ == &command->event()) { return hipSuccess; @@ -240,18 +237,15 @@ hipError_t Event::enqueueRecordCommand(hipStream_t stream, amd::Command* command } // ================================================================================================ -hipError_t Event::addMarker(hipStream_t stream, amd::Command* command, +hipError_t Event::addMarker(hip::Stream* hip_stream, amd::Command* command, bool batch_flush) { - // Skip wait as we should not be resolving stream in this sub - constexpr bool kWait = false; - hip::Stream* hip_stream = hip::getStream(stream, kWait); // Keep the lock always at the beginning of this to avoid a race. SWDEV-277847 amd::ScopedLock lock(lock_); hipError_t status = recordCommand(command, hip_stream, 0, batch_flush); if (status != hipSuccess) { return hipSuccess; } - status = enqueueRecordCommand(stream, command); + status = enqueueRecordCommand(hip_stream, command); return status; } @@ -425,7 +419,7 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream, unsigned if (g_devices[e->deviceId()]->devices()[0] != &hip_stream->device()) { return hipErrorInvalidHandle; } - status = e->addMarker(stream, nullptr, !hip::Event::kBatchFlush); + status = e->addMarker(hip_stream, nullptr, !hip::Event::kBatchFlush); } return status; } diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp index 7740b5ea37..456bf3d0a6 100644 --- a/projects/clr/hipamd/src/hip_event.hpp +++ b/projects/clr/hipamd/src/hip_event.hpp @@ -121,12 +121,12 @@ class Event { hipError_t elapsedTime(Event& eStop, float& ms); virtual hipError_t streamWaitCommand(amd::Command*& command, hip::Stream* stream); - virtual hipError_t streamWait(hipStream_t stream, uint flags); + virtual hipError_t streamWait(hip::Stream* stream, uint flags); virtual hipError_t recordCommand(amd::Command*& command, amd::HostQueue* stream, uint32_t flags = 0, bool batch_flush = true); - virtual hipError_t enqueueRecordCommand(hipStream_t stream, amd::Command* command); - hipError_t addMarker(hipStream_t stream, amd::Command* command, + virtual hipError_t enqueueRecordCommand(hip::Stream* stream, amd::Command* command); + hipError_t addMarker(hip::Stream* stream, amd::Command* command, bool batch_flush = true); void BindCommand(amd::Command& command) { @@ -209,11 +209,11 @@ class IPCEvent : public Event { hipError_t synchronize(); hipError_t query(); - hipError_t streamWait(hipStream_t stream, uint flags); + hipError_t streamWait(hip::Stream* stream, uint flags); hipError_t recordCommand(amd::Command*& command, amd::HostQueue* queue, uint32_t flags = 0, bool batch_flush = true) override; - hipError_t enqueueRecordCommand(hipStream_t stream, amd::Command* command); + hipError_t enqueueRecordCommand(hip::Stream* stream, amd::Command* command); }; diff --git a/projects/clr/hipamd/src/hip_event_ipc.cpp b/projects/clr/hipamd/src/hip_event_ipc.cpp index 52849e874b..9504a923d3 100644 --- a/projects/clr/hipamd/src/hip_event_ipc.cpp +++ b/projects/clr/hipamd/src/hip_event_ipc.cpp @@ -104,13 +104,11 @@ hipError_t IPCEvent::synchronize() { } // ================================================================================================ -hipError_t IPCEvent::streamWait(hipStream_t stream, uint flags) { - +hipError_t IPCEvent::streamWait(hip::Stream* stream, uint flags) { int offset = ipc_evt_.ipc_shmem_->read_index; - hipError_t status = ihipStreamOperation(stream, ROCCLR_COMMAND_STREAM_WAIT_VALUE, - &(ipc_evt_.ipc_shmem_->signal[offset]), - 0, - 1, 1, sizeof(uint32_t)); + hipError_t status = + ihipStreamOperation(reinterpret_cast(stream), ROCCLR_COMMAND_STREAM_WAIT_VALUE, + &(ipc_evt_.ipc_shmem_->signal[offset]), 0, 1, 1, sizeof(uint32_t)); return status; } @@ -122,7 +120,7 @@ hipError_t IPCEvent::recordCommand(amd::Command*& command, amd::HostQueue* strea } // ================================================================================================ -hipError_t IPCEvent::enqueueRecordCommand(hipStream_t stream, amd::Command* command) { +hipError_t IPCEvent::enqueueRecordCommand(hip::Stream* stream, amd::Command* command) { amd::Event& tEvent = command->event(); createIpcEventShmemIfNeeded(); @@ -136,15 +134,14 @@ hipError_t IPCEvent::enqueueRecordCommand(hipStream_t stream, amd::Command* comm ipc_evt_.ipc_shmem_->owners_device_id = deviceId(); command->enqueue(); - // device writes 0 to signal after the hipEventRecord command is completed - // the signal value is checked by WaitThenDecrementSignal cb - hipError_t status = ihipStreamOperation(stream, ROCCLR_COMMAND_STREAM_WRITE_VALUE, - &(ipc_evt_.ipc_shmem_->signal[offset]), - 0, - 0, 0, sizeof(uint32_t)); - if (status != hipSuccess) { - return status; - } + // device writes 0 to signal after the hipEventRecord command is completed + // the signal value is checked by WaitThenDecrementSignal cb + hipError_t status = ihipStreamOperation( + reinterpret_cast(stream), ROCCLR_COMMAND_STREAM_WRITE_VALUE, + &(ipc_evt_.ipc_shmem_->signal[offset]), 0, 0, 0, sizeof(uint32_t)); + if (status != hipSuccess) { + return status; + } // Update read index to indicate new signal. int expected = write_index - 1; diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 1241cf6ddc..c36e7c1d19 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -1597,7 +1597,8 @@ hipError_t ihipGraphLaunch(hip::GraphExec* graphExec, hipStream_t stream) { if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } - return graphExec->Run(stream); + hip::Stream* launch_stream = hip::getStream(stream); + return graphExec->Run(launch_stream); } hipError_t hipGraphLaunch_common(hip::GraphExec* graphExec, hipStream_t stream) { diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index 0ad1144b2f..b30a77ecb7 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -682,11 +682,9 @@ bool Graph::RunNodes( } // ================================================================================================ -hipError_t GraphExec::Run(hipStream_t graph_launch_stream) { +hipError_t GraphExec::Run(hip::Stream* launch_stream) { hipError_t status = hipSuccess; - hip::Stream* launch_stream = hip::getStream(graph_launch_stream); - if (flags_ & hipGraphInstantiateFlagAutoFreeOnLaunch) { if (!topoOrder_.empty()) { topoOrder_[0]->GetParentGraph()->FreeAllMemory(launch_stream); diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 97c32532b2..79f3522872 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -776,7 +776,7 @@ class GraphExec : public amd::ReferenceCountedObject, public Graph { uint64_t GetFlags() const { return flags_; } hipError_t Init(); hipError_t CreateStreams(uint32_t num_streams); - hipError_t Run(hipStream_t stream); + hipError_t Run(hip::Stream* stream); // Capture GPU Packets from graph commands hipError_t CaptureAQLPackets(); hipError_t UpdateAQLPacket(hip::GraphNode* node); @@ -2118,8 +2118,7 @@ class GraphEventRecordNode : public GraphNode { if (!commands_.empty()) { hip::Event* e = reinterpret_cast(event_); // command release during enqueueRecordCommand - hipError_t status = e->enqueueRecordCommand( - reinterpret_cast(stream), commands_[0]); + hipError_t status = e->enqueueRecordCommand(stream, commands_[0]); if (status != hipSuccess) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Enqueue event record command failed for node %p - status %d", this, diff --git a/projects/clr/hipamd/src/hip_mempool.cpp b/projects/clr/hipamd/src/hip_mempool.cpp index 567928a844..230622f945 100644 --- a/projects/clr/hipamd/src/hip_mempool.cpp +++ b/projects/clr/hipamd/src/hip_mempool.cpp @@ -202,7 +202,7 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) { event = new hip::Event(0); if (event != nullptr) { if (hipSuccess != - event->addMarker(reinterpret_cast(hip_stream), nullptr)) { + event->addMarker(hip_stream, nullptr)) { delete event; event = nullptr; } else { diff --git a/projects/clr/hipamd/src/hip_mempool_impl.cpp b/projects/clr/hipamd/src/hip_mempool_impl.cpp index 511e1a8a2c..d979ba7bec 100644 --- a/projects/clr/hipamd/src/hip_mempool_impl.cpp +++ b/projects/clr/hipamd/src/hip_mempool_impl.cpp @@ -294,7 +294,7 @@ bool MemoryPool::FreeMemory(amd::Memory* memory, Stream* stream, Event* event) { // Add a marker to the stream to trace availability of this memory Event* e = new hip::Event(0); if (e != nullptr) { - if (hipSuccess == e->addMarker(reinterpret_cast(stream), nullptr)) { + if (hipSuccess == e->addMarker(stream, nullptr)) { ts.SetEvent(e); // Make sure runtime sends a notification auto result = e->ready(); diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index 7d048ec67a..c78f23c7db 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -465,7 +465,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, if (startEvent != nullptr) { hip::Event* eStart = reinterpret_cast(startEvent); - status = eStart->addMarker(hStream, nullptr); + status = eStart->addMarker(hip_stream, nullptr); if (status != hipSuccess) { return status; } diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 4764fcffca..b0ea00ec74 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -451,7 +451,7 @@ hipError_t hipStreamWaitEvent_common(hipStream_t stream, hipEvent_t event, unsig if (event == nullptr || !hip::isValid(stream)) { return hipErrorInvalidHandle; } - hip::Stream* waitStream = reinterpret_cast(stream); + hip::Stream* waitStream = hip::getStream(stream); hip::Event* e = reinterpret_cast(event); auto eventStreamHandle = reinterpret_cast(e->GetCaptureStream()); // the stream associated with the device might have been destroyed @@ -491,7 +491,7 @@ hipError_t hipStreamWaitEvent_common(hipStream_t stream, hipEvent_t event, unsig eventStream->GetDevice()->AddSafeStream(eventStream, waitStream); } } - status = e->streamWait(stream, flags); + status = e->streamWait(waitStream, flags); } return status; }