From a896cb9af08a9d8cb3f1e2195e06e40692970228 Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 15 May 2018 16:26:16 -0400 Subject: [PATCH] P4 to Git Change 1555193 by cpaquot@cpaquot-ocl-lc-lnx on 2018/05/15 16:19:50 SWDEV-145570 - [HIP] Implemented events Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#4 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.hpp#2 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#25 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#9 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#9 edit --- hipamd/api/hip/hip_event.cpp | 139 +++++++++++++++++++++++++++------- hipamd/api/hip/hip_event.hpp | 21 ++++- hipamd/api/hip/hip_memory.cpp | 38 +++++++--- hipamd/api/hip/hip_module.cpp | 3 +- hipamd/api/hip/hip_stream.cpp | 29 ++++++- 5 files changed, 187 insertions(+), 43 deletions(-) diff --git a/hipamd/api/hip/hip_event.cpp b/hipamd/api/hip/hip_event.cpp index b9930636bb..127148e63c 100644 --- a/hipamd/api/hip/hip_event.cpp +++ b/hipamd/api/hip/hip_event.cpp @@ -24,37 +24,68 @@ THE SOFTWARE. #include "hip_event.hpp" +hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { + if (event == nullptr) { + return hipErrorInvalidValue; + } + + unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming | + hipEventReleaseToDevice | hipEventReleaseToSystem; + const unsigned releaseFlags = (hipEventReleaseToDevice | hipEventReleaseToSystem); + + const bool illegalFlags = + (flags & ~supportedFlags) || // can't set any unsupported flags. + (flags & releaseFlags) == releaseFlags; // can't set both release flags + + if (!illegalFlags) { + hip::Event* e = new hip::Event(flags); + + if (e == nullptr) { + return hipErrorOutOfMemory; + } + + *event = reinterpret_cast(e); + } else { + return hipErrorInvalidValue; + } + return hipSuccess; +} + +hipError_t ihipEventQuery(hipEvent_t event) { + if (event == nullptr) { + return hipErrorInvalidResourceHandle; + } + + hip::Event* e = reinterpret_cast(event); + + if (e->event_ == nullptr) { + return hipErrorInvalidResourceHandle; + } + + e->event_->notifyCmdQueue(); + + return (e->event_->status() == CL_COMPLETE)? hipSuccess : hipErrorNotReady; +} + hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { HIP_INIT_API(event, flags); - hip::Event* e = new hip::Event(flags); - - if (e == nullptr) { - return hipErrorOutOfMemory; - } - - *event = reinterpret_cast(e); - - return hipSuccess; -} + return ihipEventCreateWithFlags(event, flags); +} hipError_t hipEventCreate(hipEvent_t* event) { HIP_INIT_API(event); - hip::Event* e = new hip::Event(0); - - if (e == nullptr) { - return hipErrorOutOfMemory; - } - - *event = reinterpret_cast(e); - - return hipSuccess; + return ihipEventCreateWithFlags(event, 0); } hipError_t hipEventDestroy(hipEvent_t event) { HIP_INIT_API(event); + if (event == nullptr) { + return hipErrorInvalidResourceHandle; + } + delete reinterpret_cast(event); return hipSuccess; @@ -63,29 +94,83 @@ hipError_t hipEventDestroy(hipEvent_t event) { hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { HIP_INIT_API(ms, start, stop); - return hipErrorUnknown; + if (start == nullptr || stop == nullptr) { + return hipErrorInvalidResourceHandle; + } + + hip::Event* eStart = reinterpret_cast(start); + hip::Event* eStop = reinterpret_cast(stop); + + if (eStart->event_ == nullptr || + eStop->event_ == nullptr) { + return hipErrorInvalidResourceHandle; + } + + if ((eStart->flags | eStop->flags) & hipEventDisableTiming) { + return hipErrorInvalidResourceHandle; + } + + if (ihipEventQuery(start) == hipErrorNotReady || + ihipEventQuery(stop) == hipErrorNotReady) { + return hipErrorNotReady; + } + + if (ms == nullptr) { + return hipErrorInvalidValue; + } + + *ms = static_cast(eStop->event_->profilingInfo().submitted_ - eStart->event_->profilingInfo().submitted_)/1000000.f; + + return hipSuccess; } hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { HIP_INIT_API(event, stream); - assert(0 && "Unimplemented"); + if (event == nullptr) { + return hipErrorInvalidResourceHandle; + } - return hipErrorUnknown; + hip::Event* e = reinterpret_cast(event); + + if (stream == nullptr) { + e->stream_ = hip::getNullStream(); + } else { + e->stream_ = as_amd(reinterpret_cast(stream))->asHostQueue(); + } + amd::Command* command = (e->flags & hipEventDisableTiming)? new amd::Marker(*e->stream_, true) : + new hip::TimerMarker(*e->stream_); + command->enqueue(); + + if (e->event_ != nullptr) { + e->event_->release(); + } + + e->event_ = &command->event(); + + return hipSuccess; } hipError_t hipEventSynchronize(hipEvent_t event) { HIP_INIT_API(event); - assert(0 && "Unimplemented"); + if (event == nullptr) { + return hipErrorInvalidResourceHandle; + } - return hipErrorUnknown; + hip::Event* e = reinterpret_cast(event); + + if (e->event_ == nullptr) { + return hipErrorInvalidResourceHandle; + } + + e->event_->awaitCompletion(); + + return hipSuccess; } hipError_t hipEventQuery(hipEvent_t event) { HIP_INIT_API(event); - assert(0 && "Unimplemented"); - - return hipErrorUnknown; + return ihipEventQuery(event); } diff --git a/hipamd/api/hip/hip_event.hpp b/hipamd/api/hip/hip_event.hpp index 3ac1ea8bfe..19f93a5c27 100644 --- a/hipamd/api/hip/hip_event.hpp +++ b/hipamd/api/hip/hip_event.hpp @@ -27,12 +27,27 @@ THE SOFTWARE. namespace hip { +class TimerMarker: public amd::Marker { +public: + TimerMarker(amd::HostQueue& queue) : amd::Marker(queue, true) { + profilingInfo_.enabled_ = true; + profilingInfo_.callback_ = nullptr; + profilingInfo_.start_ = profilingInfo_.end_ = 0; + } +}; + class Event { public: - Event(unsigned int flags) : flags(flags) {} - ~Event() {} + Event(unsigned int flags) : flags(flags), stream_(nullptr), event_(nullptr) {} + ~Event() { + if (event_ != nullptr) { + event_->release(); + } + } unsigned int flags; -private: + + amd::HostQueue* stream_; + amd::Event* event_; }; }; diff --git a/hipamd/api/hip/hip_memory.cpp b/hipamd/api/hip/hip_memory.cpp index 568ea557ec..a754745ebb 100644 --- a/hipamd/api/hip/hip_memory.cpp +++ b/hipamd/api/hip/hip_memory.cpp @@ -68,7 +68,6 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, amd::HostQueue& queue, bool isAsync = false) { - amd::Command* command = nullptr; amd::Command::EventWaitList waitList; size_t sOffset = 0; @@ -116,15 +115,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin return hipErrorOutOfMemory; } -// FIXME: virtualize MemoryCommand::validateMemory() -#if 0 - // Make sure we have memory for the command execution - if (CL_SUCCESS != command->validateMemory()) { - delete command; - return hipErrorMemoryAllocation; - } -#endif - command->enqueue(); if (!isAsync) { command->awaitCompletion(); @@ -193,6 +183,7 @@ hipError_t hipFree(void* ptr) { hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { HIP_INIT_API(dst, src, sizeBytes, kind); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemcpy(dst, src, sizeBytes, kind, *queue); } @@ -203,8 +194,10 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st amd::HostQueue* queue; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -214,6 +207,7 @@ hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t st hipError_t hipMemset(void* dst, int value, size_t sizeBytes) { HIP_INIT_API(dst, value, sizeBytes); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemset(dst, value, sizeBytes, *queue); @@ -533,6 +527,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { HIP_INIT_API(dst, src, sizeBytes); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, *queue); @@ -541,6 +536,7 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) { hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_API(dst, src, sizeBytes); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, *queue); @@ -549,6 +545,7 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) { hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) { HIP_INIT_API(dst, src, sizeBytes); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, *queue); @@ -557,6 +554,7 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { HIP_INIT_API(dst, src, sizeBytes); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToHost, *queue); @@ -569,8 +567,10 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, amd::HostQueue* queue = nullptr; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -585,8 +585,10 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, amd::HostQueue* queue = nullptr; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -601,8 +603,10 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz amd::HostQueue* queue = nullptr; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -617,8 +621,10 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, amd::HostQueue* queue = nullptr; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -723,6 +729,7 @@ hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t height, hipMemcpyKind kind) { HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue); @@ -736,8 +743,10 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp amd::HostQueue* queue; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -757,6 +766,7 @@ hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, size_t count, hipMemcpyKind kind) { HIP_INIT_API(dstArray, wOffset, hOffset, src, count, kind); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); amd::Command* command = nullptr; @@ -795,6 +805,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs size_t count, hipMemcpyKind kind) { HIP_INIT_API(dst, srcArray, wOffset, hOffset, count, kind); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); amd::Command* command = nullptr; @@ -833,6 +844,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffs hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) { HIP_INIT_API(dstArray, dstOffset, srcHost, count); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); amd::Command::EventWaitList waitList; @@ -856,6 +868,7 @@ hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHo hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t count) { HIP_INIT_API(dst, srcArray, srcOffset, count); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); amd::Command::EventWaitList waitList; @@ -879,6 +892,7 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p) { HIP_INIT_API(p); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); size_t byteSize; @@ -1035,6 +1049,7 @@ hipError_t ihipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height) { HIP_INIT_API(dst, pitch, value, width, height); + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); return ihipMemset2D(dst, pitch, value, width, height, *queue); } @@ -1045,8 +1060,10 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, amd::HostQueue* queue = nullptr; if (stream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -1060,6 +1077,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dst, unsigned char value, size_t sizeBytes return hipErrorInvalidValue; } + hip::syncStreams(); amd::HostQueue* queue = hip::getNullStream(); size_t offset = 0; amd::Command::EventWaitList waitList; diff --git a/hipamd/api/hip/hip_module.cpp b/hipamd/api/hip/hip_module.cpp index 0a5675114c..0feefad7a3 100644 --- a/hipamd/api/hip/hip_module.cpp +++ b/hipamd/api/hip/hip_module.cpp @@ -146,8 +146,10 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, amd::HostQueue* queue; if (hStream == nullptr) { + hip::syncStreams(); queue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); queue = as_amd(reinterpret_cast(hStream))->asHostQueue(); } if (!queue) { @@ -195,7 +197,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, } command->enqueue(); - command->awaitCompletion(); command->release(); return hipSuccess; diff --git a/hipamd/api/hip/hip_stream.cpp b/hipamd/api/hip/hip_stream.cpp index f4995bcd81..6fc8140f6c 100644 --- a/hipamd/api/hip/hip_stream.cpp +++ b/hipamd/api/hip/hip_stream.cpp @@ -22,6 +22,7 @@ THE SOFTWARE. #include #include "hip_internal.hpp" +#include "hip_event.hpp" #include "thread/monitor.hpp" static amd::Monitor streamSetLock("Guards global stream set"); @@ -97,8 +98,12 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { amd::HostQueue* hostQueue; if (stream == nullptr) { + hip::syncStreams(); + hostQueue = hip::getNullStream(); } else { + hip::getNullStream()->finish(); + hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); } @@ -130,9 +135,29 @@ hipError_t hipStreamDestroy(hipStream_t stream) { hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(stream, event, flags); - assert(0 && "Unimplemented"); + if (stream == nullptr || event == nullptr) { + return hipErrorInvalidResourceHandle; + } - return hipErrorUnknown; + 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->context(), 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; } hipError_t hipStreamQuery(hipStream_t stream) {