AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-469422 - Avoid using of hipStream_t in internal methods (#69)

Browse Source 
Change-Id: Ifd5362f371c846a88241927383cb95cf046548ef

[ROCm/clr commit: fb92683d86]
This commit is contained in:
Godavarthy Surya, Anusha
2025-04-28 15:09:11 +05:30
committed by GitHub
parent 0eb2e5e8f2
commit ff69bcc903
10 changed files with 35 additions and 46 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_event.cpp
+7 -13
View File
@@ -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;
}
projects/clr/hipamd/src/hip_event.hpp
+5 -5
View File
@@ -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);
};
projects/clr/hipamd/src/hip_event_ipc.cpp
+13 -16
View File
@@ -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<hipStream_t>(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<hipStream_t>(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;
projects/clr/hipamd/src/hip_graph.cpp
+2 -1
View File
@@ -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) {
projects/clr/hipamd/src/hip_graph_internal.cpp
+1 -3
View File
@@ -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);
projects/clr/hipamd/src/hip_graph_internal.hpp
+2 -3
View File
@@ -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<hip::Event*>(event_);
// command release during enqueueRecordCommand
hipError_t status = e->enqueueRecordCommand(
reinterpret_cast<hipStream_t>(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,
projects/clr/hipamd/src/hip_mempool.cpp
+1 -1
View File
@@ -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<hipStream_t>(hip_stream), nullptr)) {
event->addMarker(hip_stream, nullptr)) {
delete event;
event = nullptr;
} else {
projects/clr/hipamd/src/hip_mempool_impl.cpp
+1 -1
View File
@@ -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<hipStream_t>(stream), nullptr)) {
if (hipSuccess == e->addMarker(stream, nullptr)) {
ts.SetEvent(e);
// Make sure runtime sends a notification
auto result = e->ready();
projects/clr/hipamd/src/hip_module.cpp
+1 -1
View File
@@ -465,7 +465,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
if (startEvent != nullptr) {
hip::Event* eStart = reinterpret_cast<hip::Event*>(startEvent);
status = eStart->addMarker(hStream, nullptr);
status = eStart->addMarker(hip_stream, nullptr);
if (status != hipSuccess) {
return status;
}
projects/clr/hipamd/src/hip_stream.cpp
+2 -2
View File
@@ -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<hip::Stream*>(stream);
hip::Stream* waitStream = hip::getStream(stream);
hip::Event* e = reinterpret_cast<hip::Event*>(event);
auto eventStreamHandle = reinterpret_cast<hipStream_t>(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;
}