From 250a44189978a39e01decdada7fd2b29c9e78dce Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 9 Aug 2019 20:41:13 -0400 Subject: [PATCH] P4 to Git Change 1981152 by cpaquot@cpaquot-ocl-lc-lnx on 2019/08/09 20:33:32 SWDEV-193430 - [HIP] Delay creating HostQueue till commands are enqueued. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#14 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.hpp#5 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#34 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#68 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#31 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#21 edit [ROCm/hip commit: 92840c805e52c76c57a305df38ad376db603b4a2] --- projects/hip/api/hip/hip_event.cpp | 6 +- projects/hip/api/hip/hip_event.hpp | 2 +- projects/hip/api/hip/hip_internal.hpp | 16 +++++ projects/hip/api/hip/hip_memory.cpp | 14 ++--- projects/hip/api/hip/hip_module.cpp | 4 +- projects/hip/api/hip/hip_stream.cpp | 87 ++++++++++++++++++--------- 6 files changed, 85 insertions(+), 44 deletions(-) diff --git a/projects/hip/api/hip/hip_event.cpp b/projects/hip/api/hip/hip_event.cpp index 746d93a8b2..8262a770d0 100644 --- a/projects/hip/api/hip/hip_event.cpp +++ b/projects/hip/api/hip/hip_event.cpp @@ -98,9 +98,7 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) { return hipSuccess; } -hipError_t Event::streamWait(hipStream_t stream, uint flags) { - amd::HostQueue* hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); - +hipError_t Event::streamWait(amd::HostQueue* hostQueue, uint flags) { if (stream_ == hostQueue) return hipSuccess; amd::ScopedLock lock(lock_); @@ -237,7 +235,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { if (stream == nullptr) { queue = hip::getNullStream(); } else { - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } amd::Command* command = queue->getLastQueuedCommand(true); diff --git a/projects/hip/api/hip/hip_event.hpp b/projects/hip/api/hip/hip_event.hpp index 4c6fb132e8..410c19c7b4 100644 --- a/projects/hip/api/hip/hip_event.hpp +++ b/projects/hip/api/hip/hip_event.hpp @@ -50,7 +50,7 @@ public: hipError_t query(); hipError_t synchronize(); hipError_t elapsedTime(Event& stop, float& ms); - hipError_t streamWait(hipStream_t stream, uint flags); + hipError_t streamWait(amd::HostQueue* queue, uint flags); void addMarker(amd::HostQueue* queue, amd::Command* command); diff --git a/projects/hip/api/hip/hip_internal.hpp b/projects/hip/api/hip/hip_internal.hpp index 86221b197e..6e7614f315 100644 --- a/projects/hip/api/hip/hip_internal.hpp +++ b/projects/hip/api/hip/hip_internal.hpp @@ -86,6 +86,22 @@ namespace hip { static Function* asFunction(hipFunction_t f) { return reinterpret_cast(f); } }; + + struct Stream { + amd::HostQueue* queue; + + amd::Device* device; + amd::Context* context; + amd::CommandQueue::Priority priority; + unsigned int flags; + + Stream(amd::Device* dev, amd::Context* ctx, amd::CommandQueue::Priority p, unsigned int f); + void create(); + amd::HostQueue* asHostQueue(); + void destroy(); + void finish(); + }; + }; struct ihipExec_t { diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp index 8ae5b36d3d..d00400c88b 100644 --- a/projects/hip/api/hip/hip_memory.cpp +++ b/projects/hip/api/hip/hip_memory.cpp @@ -177,7 +177,7 @@ hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes, queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } if (memory != nullptr) { @@ -771,7 +771,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue, true)); @@ -789,7 +789,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, @@ -807,7 +807,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, @@ -825,7 +825,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, @@ -928,7 +928,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue, true)); @@ -1355,7 +1355,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(stream))->asHostQueue(); + queue = reinterpret_cast(stream)->asHostQueue(); } HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue, true)); diff --git a/projects/hip/api/hip/hip_module.cpp b/projects/hip/api/hip/hip_module.cpp index f9c5a85137..315caffa31 100644 --- a/projects/hip/api/hip/hip_module.cpp +++ b/projects/hip/api/hip/hip_module.cpp @@ -231,7 +231,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, queue = hip::getNullStream(); } else { hip::getNullStream()->finish(); - queue = as_amd(reinterpret_cast(hStream))->asHostQueue(); + queue = reinterpret_cast(hStream)->asHostQueue(); } if ((params & amd::NDRangeKernelCommand::CooperativeGroups) && !device->info().cooperativeGroups_) { @@ -387,7 +387,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL for (int i = 0; i < numDevices; ++i) { hipSetDevice(i); const hipLaunchParams& launch = launchParamsList[i]; - amd::HostQueue* queue = as_amd(reinterpret_cast(launch.stream))->asHostQueue(); + amd::HostQueue* queue = reinterpret_cast(launch.stream)->asHostQueue(); hipFunction_t func = PlatformState::instance().getFunc(launch.func, i); if (func == nullptr) { HIP_RETURN(result); diff --git a/projects/hip/api/hip/hip_stream.cpp b/projects/hip/api/hip/hip_stream.cpp index cbba339efb..a87e0e0d79 100644 --- a/projects/hip/api/hip/hip_stream.cpp +++ b/projects/hip/api/hip/hip_stream.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "thread/monitor.hpp" static amd::Monitor streamSetLock("Guards global stream set"); -static std::unordered_set streamSet; +static std::unordered_set streamSet; // Internal structure for stream callback handler class StreamCallback { @@ -52,6 +52,37 @@ void syncStreams() { } } +Stream::Stream(amd::Device* dev, amd::Context* ctx, amd::CommandQueue::Priority p, unsigned int f) : + queue(nullptr), device(dev), context(ctx), priority(p), flags(f) {} + +void Stream::create() { + cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE; + queue = new amd::HostQueue(*context, *device, properties, + amd::CommandQueue::RealTimeDisabled, priority); + assert(queue != nullptr); + queue->create(); +} + +amd::HostQueue* Stream::asHostQueue() { + if (queue == nullptr) { + create(); + } + return queue; +} + +void Stream::destroy() { + if (queue != nullptr) { + queue->release(); + queue = nullptr; + } +} + +void Stream::finish() { + if (queue != nullptr) { + queue->finish(); + } +} + }; void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data) { @@ -66,12 +97,9 @@ void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, static hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags, amd::CommandQueue::Priority priority) { amd::Device* device = hip::getCurrentContext()->devices()[0]; - cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE; - amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, properties, - amd::CommandQueue::RealTimeDisabled, - priority); + hip::Stream* hStream = new hip::Stream(device, hip::getCurrentContext(), priority, flags); - if (queue == nullptr || !queue->create()) { + if (hStream == nullptr) { return hipErrorOutOfMemory; } @@ -80,11 +108,11 @@ static hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags, amd: { amd::ScopedLock lock(streamSetLock); - streamSet.insert(queue); + streamSet.insert(hStream); } } - *stream = reinterpret_cast(as_cl(queue)); + *stream = reinterpret_cast(hStream); return hipSuccess; } @@ -129,11 +157,10 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { HIP_INIT_API(stream, flags); - amd::HostQueue* hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); - auto it = streamSet.find(hostQueue); + hip::Stream* hStream = reinterpret_cast(stream); - if(flags != nullptr) { - *flags = (it == streamSet.end()) ? hipStreamNonBlocking : hipStreamDefault; + if(flags != nullptr && hStream != nullptr) { + *flags = hStream->flags; } else { HIP_RETURN(hipErrorInvalidValue); } @@ -150,18 +177,15 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { hip::syncStreams(); hostQueue = hip::getNullStream(); + + hostQueue->finish(); } else { hip::getNullStream()->finish(); - hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); + hip::Stream* hStream = reinterpret_cast(stream); + + hStream->finish(); } - - if (hostQueue == nullptr) { - HIP_RETURN(hipErrorUnknown); - } - - hostQueue->finish(); - HIP_RETURN(hipSuccess); } @@ -174,13 +198,12 @@ hipError_t hipStreamDestroy(hipStream_t stream) { amd::ScopedLock lock(streamSetLock); - amd::HostQueue* hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); + hip::Stream* hStream = reinterpret_cast(stream); - // Release last tracked command - hostQueue->setLastQueuedCommand(nullptr); + hStream->destroy(); + streamSet.erase(hStream); - hostQueue->release(); - streamSet.erase(hostQueue); + delete hStream; HIP_RETURN(hipSuccess); } @@ -188,8 +211,12 @@ hipError_t hipStreamDestroy(hipStream_t stream) { hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(stream, event, flags); + amd::HostQueue* queue; + if (stream == nullptr) { - stream = reinterpret_cast(as_cl(hip::getNullStream())); + queue = hip::getNullStream(); + } else { + queue = reinterpret_cast(stream)->asHostQueue(); } if (event == nullptr) { @@ -198,7 +225,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int hip::Event* e = reinterpret_cast(event); - return HIP_RETURN(e->streamWait(stream, flags)); + return HIP_RETURN(e->streamWait(queue, flags)); } hipError_t hipStreamQuery(hipStream_t stream) { @@ -208,7 +235,7 @@ hipError_t hipStreamQuery(hipStream_t stream) { if (stream == nullptr) { hostQueue = hip::getNullStream(); } else { - hostQueue = as_amd(reinterpret_cast(stream))->asHostQueue(); + hostQueue = reinterpret_cast(stream)->asHostQueue(); } amd::Command* command = hostQueue->getLastQueuedCommand(false); @@ -228,8 +255,8 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback unsigned int flags) { HIP_INIT_API(stream, callback, userData, flags); - amd::HostQueue* hostQueue = as_amd(reinterpret_cast - (stream))->asHostQueue(); + amd::HostQueue* hostQueue = reinterpret_cast + (stream)->asHostQueue(); amd::Command* command = hostQueue->getLastQueuedCommand(true); amd::Event& event = command->event(); StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);