diff --git a/projects/clr/hipamd/src/hip_event.hpp b/projects/clr/hipamd/src/hip_event.hpp index 2bd43b9658..b7c888d392 100644 --- a/projects/clr/hipamd/src/hip_event.hpp +++ b/projects/clr/hipamd/src/hip_event.hpp @@ -210,8 +210,6 @@ class IPCEvent : public Event { hipError_t synchronize(); hipError_t query(); - hipError_t streamWaitCommand(amd::Command*& command, hip::Stream* stream); - hipError_t enqueueStreamWaitCommand(hipStream_t stream, amd::Command* command); hipError_t streamWait(hipStream_t stream, uint flags); hipError_t recordCommand(amd::Command*& command, amd::HostQueue* queue, diff --git a/projects/clr/hipamd/src/hip_event_ipc.cpp b/projects/clr/hipamd/src/hip_event_ipc.cpp index 3840ad82f0..0a63870b8b 100644 --- a/projects/clr/hipamd/src/hip_event_ipc.cpp +++ b/projects/clr/hipamd/src/hip_event_ipc.cpp @@ -101,42 +101,14 @@ hipError_t IPCEvent::synchronize() { return hipSuccess; } -hipError_t IPCEvent::streamWaitCommand(amd::Command*& command, hip::Stream* stream) { - command = new amd::Marker(*stream, false); - if (command == NULL) { - return hipErrorOutOfMemory; - } - return hipSuccess; -} - -hipError_t IPCEvent::enqueueStreamWaitCommand(hipStream_t stream, amd::Command* command) { - auto t{new CallbackData{ipc_evt_.ipc_shmem_->read_index, ipc_evt_.ipc_shmem_}}; - StreamCallback* cbo = new StreamAddCallback( - stream, reinterpret_cast(WaitThenDecrementSignal), t); - if (!command->setCallback(CL_COMPLETE, ihipStreamCallback, cbo)) { - command->release(); - return hipErrorInvalidHandle; - } - command->enqueue(); - command->release(); - command->awaitCompletion(); - return hipSuccess; -} - hipError_t IPCEvent::streamWait(hipStream_t stream, uint flags) { - hip::Stream* hip_stream = hip::getStream(stream); - amd::ScopedLock lock(lock_); - if(query() != hipSuccess) { - amd::Command* command; - hipError_t status = streamWaitCommand(command, hip_stream); - if (status != hipSuccess) { - return status; - } - status = enqueueStreamWaitCommand(stream, command); - return status; - } - return hipSuccess; + 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)); + return status; } // ================================================================================================