From 4f37e552ad3584ff1c159c2d07963d0735ce3caf Mon Sep 17 00:00:00 2001 From: foreman Date: Mon, 6 May 2019 17:43:06 -0400 Subject: [PATCH] P4 to Git Change 1778971 by cpaquot@cpaquot-ocl-lc-lnx on 2019/05/06 17:03:12 SWDEV-187125 - [HIP] Protect hip::Event with lock and add threadId to logs. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#9 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.hpp#3 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#30 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#24 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#18 edit [ROCm/clr commit: 84d257cf3bb6a867a8a695048ce3148658ad7368] --- projects/clr/hipamd/api/hip/hip_event.cpp | 146 +++++++++++++------ projects/clr/hipamd/api/hip/hip_event.hpp | 14 +- projects/clr/hipamd/api/hip/hip_internal.hpp | 2 +- projects/clr/hipamd/api/hip/hip_module.cpp | 16 +- projects/clr/hipamd/api/hip/hip_stream.cpp | 17 +-- 5 files changed, 121 insertions(+), 74 deletions(-) diff --git a/projects/clr/hipamd/api/hip/hip_event.cpp b/projects/clr/hipamd/api/hip/hip_event.cpp index 317e0edc90..aa61443d2f 100644 --- a/projects/clr/hipamd/api/hip/hip_event.cpp +++ b/projects/clr/hipamd/api/hip/hip_event.cpp @@ -24,6 +24,98 @@ THE SOFTWARE. #include "hip_event.hpp" +namespace hip { + +bool Event::ready() { + event_->notifyCmdQueue(); + + return (event_->status() == CL_COMPLETE); +} + +hipError_t Event::query() { + amd::ScopedLock lock(lock_); + + if (event_ == nullptr) { + return hipErrorInvalidResourceHandle; + } + + return ready() ? hipSuccess : hipErrorNotReady; +} + +hipError_t Event::synchronize() { + amd::ScopedLock lock(lock_); + + if (event_ == nullptr) { + return hipErrorInvalidResourceHandle; + } + + event_->awaitCompletion(); + + return hipSuccess; +} + +hipError_t Event::elapsedTime(Event& eStop, float& ms) { + amd::ScopedLock startLock(lock_); + amd::ScopedLock stopLock(eStop.lock_); + + if (event_ == nullptr || + eStop.event_ == nullptr) { + return hipErrorInvalidResourceHandle; + } + + if ((flags | eStop.flags) & hipEventDisableTiming) { + return hipErrorInvalidResourceHandle; + } + + if (!ready() || !eStop.ready()) { + return hipErrorNotReady; + } + + ms = static_cast(static_cast(eStop.event_->profilingInfo().submitted_ - + event_->profilingInfo().submitted_))/1000000.f; + + return hipSuccess; +} + +hipError_t Event::streamWait(hipStream_t stream, uint flags) { + amd::HostQueue* hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); + + if (stream_ == hostQueue) return hipSuccess; + + amd::ScopedLock lock(lock_); + + cl_event clEvent = as_cl(event_); + + amd::Command::EventWaitList eventWaitList; + cl_int err = amd::clSetEventWaitList(eventWaitList, *hostQueue, 1, &clEvent); + if (err != CL_SUCCESS) { + return hipErrorUnknown; + } + + amd::Command* command = new amd::Marker(*hostQueue, true, eventWaitList); + if (command == NULL) { + return hipErrorOutOfMemory; + } + command->enqueue(); + command->release(); + + return hipSuccess; +} + +void Event::addMarker(amd::HostQueue* queue, amd::Command* command) { + amd::ScopedLock lock(lock_); + + stream_ = queue; + + if (event_ != nullptr) { + event_->release(); + } + + event_ = &command->event(); +} + +} + hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { if (event == nullptr) { return hipErrorInvalidValue; @@ -58,13 +150,7 @@ hipError_t ihipEventQuery(hipEvent_t event) { hip::Event* e = reinterpret_cast(event); - if (e->event_ == nullptr) { - return hipErrorInvalidResourceHandle; - } - - e->event_->notifyCmdQueue(); - - return (e->event_->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady; + return e->query(); } hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { @@ -98,31 +184,14 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { HIP_RETURN(hipErrorInvalidResourceHandle); } - hip::Event* eStart = reinterpret_cast(start); - hip::Event* eStop = reinterpret_cast(stop); - - if (eStart->event_ == nullptr || - eStop->event_ == nullptr) { - HIP_RETURN(hipErrorInvalidResourceHandle); - } - - if ((eStart->flags | eStop->flags) & hipEventDisableTiming) { - HIP_RETURN(hipErrorInvalidResourceHandle); - } - - if (ihipEventQuery(start) == hipErrorNotReady || - ihipEventQuery(stop) == hipErrorNotReady) { - HIP_RETURN(hipErrorNotReady); - } - if (ms == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - *ms = static_cast(static_cast(eStop->event_->profilingInfo().submitted_ - - eStart->event_->profilingInfo().submitted_))/1000000.f; + hip::Event* eStart = reinterpret_cast(start); + hip::Event* eStop = reinterpret_cast(stop); - HIP_RETURN(hipSuccess); + return HIP_RETURN(eStart->elapsedTime(*eStop, *ms)); } hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { @@ -134,24 +203,21 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { hip::Event* e = reinterpret_cast(event); + amd::HostQueue* queue; if (stream == nullptr) { - e->stream_ = hip::getNullStream(); + queue = hip::getNullStream(); } else { - e->stream_ = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } - amd::Command* command = e->stream_->getLastQueuedCommand(true); + amd::Command* command = queue->getLastQueuedCommand(true); if (command == nullptr) { - command = new amd::Marker(*e->stream_, true); + command = new amd::Marker(*queue, true); command->enqueue(); } - if (e->event_ != nullptr) { - e->event_->release(); - } - - e->event_ = &command->event(); + e->addMarker(queue, command); HIP_RETURN(hipSuccess); } @@ -165,13 +231,7 @@ hipError_t hipEventSynchronize(hipEvent_t event) { hip::Event* e = reinterpret_cast(event); - if (e->event_ == nullptr) { - HIP_RETURN(hipErrorInvalidResourceHandle); - } - - e->event_->awaitCompletion(); - - HIP_RETURN(hipSuccess); + HIP_RETURN(e->synchronize()); } hipError_t hipEventQuery(hipEvent_t event) { diff --git a/projects/clr/hipamd/api/hip/hip_event.hpp b/projects/clr/hipamd/api/hip/hip_event.hpp index 19f93a5c27..953665ed5b 100644 --- a/projects/clr/hipamd/api/hip/hip_event.hpp +++ b/projects/clr/hipamd/api/hip/hip_event.hpp @@ -24,6 +24,7 @@ THE SOFTWARE. #define HIP_EVENT_H #include "hip_internal.hpp" +#include "thread/monitor.hpp" namespace hip { @@ -38,7 +39,7 @@ public: class Event { public: - Event(unsigned int flags) : flags(flags), stream_(nullptr), event_(nullptr) {} + Event(unsigned int flags) : flags(flags), lock_("hipEvent_t"), stream_(nullptr), event_(nullptr) {} ~Event() { if (event_ != nullptr) { event_->release(); @@ -46,8 +47,19 @@ public: } unsigned int flags; + hipError_t query(); + hipError_t synchronize(); + hipError_t elapsedTime(Event& stop, float& ms); + hipError_t streamWait(hipStream_t stream, uint flags); + + void addMarker(amd::HostQueue* queue, amd::Command* command); + +private: + amd::Monitor lock_; amd::HostQueue* stream_; amd::Event* event_; + + bool ready(); }; }; diff --git a/projects/clr/hipamd/api/hip/hip_internal.hpp b/projects/clr/hipamd/api/hip/hip_internal.hpp index 25fe6f9537..2787e2b187 100644 --- a/projects/clr/hipamd/api/hip/hip_internal.hpp +++ b/projects/clr/hipamd/api/hip/hip_internal.hpp @@ -49,7 +49,7 @@ typedef struct ihipIpcMemHandle_st { // This macro should be called at the beginning of every HIP API. #define HIP_INIT_API(...) \ - LogPrintfInfo("%s ( %s )", __func__, ToString( __VA_ARGS__ ).c_str()); \ + LogPrintfInfo("[%zx] %s ( %s )", std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \ amd::Thread* thread = amd::Thread::current(); \ if (!CL_CHECK_THREAD(thread)) { \ HIP_RETURN(hipErrorOutOfMemory); \ diff --git a/projects/clr/hipamd/api/hip/hip_module.cpp b/projects/clr/hipamd/api/hip/hip_module.cpp index 79c5733c74..ad13ab5ff2 100644 --- a/projects/clr/hipamd/api/hip/hip_module.cpp +++ b/projects/clr/hipamd/api/hip/hip_module.cpp @@ -219,15 +219,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, } if(startEvent != nullptr) { - eStart->stream_ = queue; - amd::Command* startCommand = new hip::TimerMarker(*eStart->stream_); + amd::Command* startCommand = new hip::TimerMarker(*queue); startCommand->enqueue(); - - if (eStart->event_ != nullptr) { - eStart->event_->release(); - } - - eStart->event_ = &startCommand->event(); + eStart->addMarker(queue, startCommand); } amd::NDRangeKernelCommand* command = new amd::NDRangeKernelCommand(*queue, waitList, *kernel, ndrange, sharedMemBytes); @@ -244,11 +238,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, command->enqueue(); if(stopEvent != nullptr) { - if (eStop->event_ != nullptr) { - eStop->event_->release(); - } - eStop->stream_ = queue; - eStop->event_ = &command->event(); + eStop->addMarker(queue, command); command->retain(); } diff --git a/projects/clr/hipamd/api/hip/hip_stream.cpp b/projects/clr/hipamd/api/hip/hip_stream.cpp index 1e4c820a44..72ede63ecb 100644 --- a/projects/clr/hipamd/api/hip/hip_stream.cpp +++ b/projects/clr/hipamd/api/hip/hip_stream.cpp @@ -192,24 +192,9 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int HIP_RETURN(hipErrorInvalidResourceHandle); } - amd::HostQueue* hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); hip::Event* e = reinterpret_cast(event); - cl_event clEvent = as_cl(e->event_); - amd::Command::EventWaitList eventWaitList; - cl_int err = amd::clSetEventWaitList(eventWaitList, *hostQueue, 1, &clEvent); - if (err != CL_SUCCESS) { - HIP_RETURN(hipErrorUnknown); - } - - amd::Command* command = new amd::Marker(*hostQueue, true, eventWaitList); - if (command == NULL) { - HIP_RETURN(hipErrorOutOfMemory); - } - command->enqueue(); - command->release(); - - HIP_RETURN(hipSuccess); + return HIP_RETURN(e->streamWait(stream, flags)); } hipError_t hipStreamQuery(hipStream_t stream) {