From e4eb8e6a1b1df929b4250accecf8b11bd124ce60 Mon Sep 17 00:00:00 2001 From: foreman Date: Fri, 13 Sep 2019 11:28:33 -0400 Subject: [PATCH] P4 to Git Change 1997981 by cpaquot@cpaquot-ocl-lc-lnx on 2019/09/13 11:17:32 SWDEV-203438 - [HIP] AllGather RCCL test issue The test tries to launch a kernel on two devices at once and they need to communicate with each other. For that, it uses a custom stream for each devices. Problem is in getNullStream we used to call syncStreams all the time and it was syncing all the streams even the ones on different devices. So that made the second kernel launch (on 2n dev) to wait for the first kernel to finish which would never occur since the first one was waiting for the second one. The fix is to not call syncStreams from getNullStream because we sync already anyway prior in general. Affected files ... ... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#21 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_event.cpp#16 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#40 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#70 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#41 edit ... //depot/stg/opencl/drivers/opencl/api/hip/hip_stream.cpp#24 edit [ROCm/clr commit: 0aa24f6d4d8b20cd14e678f3b80b2924d45ea898] --- projects/clr/hipamd/api/hip/hip_context.cpp | 14 +++- projects/clr/hipamd/api/hip/hip_event.cpp | 7 +- projects/clr/hipamd/api/hip/hip_internal.hpp | 1 + projects/clr/hipamd/api/hip/hip_memory.cpp | 70 +++----------------- projects/clr/hipamd/api/hip/hip_module.cpp | 13 ++-- projects/clr/hipamd/api/hip/hip_stream.cpp | 16 +---- 6 files changed, 30 insertions(+), 91 deletions(-) diff --git a/projects/clr/hipamd/api/hip/hip_context.cpp b/projects/clr/hipamd/api/hip/hip_context.cpp index d2c37c1b13..2ec1f9db50 100644 --- a/projects/clr/hipamd/api/hip/hip_context.cpp +++ b/projects/clr/hipamd/api/hip/hip_context.cpp @@ -75,6 +75,19 @@ void setCurrentContext(unsigned int index) { g_context = g_devices[index]; } +amd::HostQueue* getQueue(hipStream_t stream) { + if (stream == nullptr) { + syncStreams(); + return getNullStream(); + } else { + hip::Stream* s = reinterpret_cast(stream); + if ((s->flags & hipStreamNonBlocking) == 0) { + getNullStream()->finish(); + } + return s->asHostQueue(); + } +} + amd::HostQueue* getNullStream(amd::Context& context) { auto stream = g_nullStreams.find(&context); if (stream == g_nullStreams.end()) { @@ -86,7 +99,6 @@ amd::HostQueue* getNullStream(amd::Context& context) { g_nullStreams[&context] = queue; return queue; } - syncStreams(); return stream->second; } diff --git a/projects/clr/hipamd/api/hip/hip_event.cpp b/projects/clr/hipamd/api/hip/hip_event.cpp index 143d2eded9..81ae93e279 100644 --- a/projects/clr/hipamd/api/hip/hip_event.cpp +++ b/projects/clr/hipamd/api/hip/hip_event.cpp @@ -231,12 +231,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { hip::Event* e = reinterpret_cast(event); - amd::HostQueue* queue; - if (stream == nullptr) { - queue = hip::getNullStream(); - } else { - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); amd::Command* command = queue->getLastQueuedCommand(true); diff --git a/projects/clr/hipamd/api/hip/hip_internal.hpp b/projects/clr/hipamd/api/hip/hip_internal.hpp index 49b3757dda..7b0ff09ed9 100644 --- a/projects/clr/hipamd/api/hip/hip_internal.hpp +++ b/projects/clr/hipamd/api/hip/hip_internal.hpp @@ -72,6 +72,7 @@ namespace hip { extern amd::Context* getCurrentContext(); extern void setCurrentContext(unsigned int index); + extern amd::HostQueue* getQueue(hipStream_t s); extern amd::HostQueue* getNullStream(amd::Context&); extern amd::HostQueue* getNullStream(); extern void syncStreams(); diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp index 0d9b1bec06..b1f1f40aec 100644 --- a/projects/clr/hipamd/api/hip/hip_memory.cpp +++ b/projects/clr/hipamd/api/hip/hip_memory.cpp @@ -95,6 +95,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin if (((srcMemory == nullptr) && (dstMemory == nullptr)) || (kind == hipMemcpyHostToHost)) { + queue.finish(); memcpy(dst, src, sizeBytes); return hipSuccess; } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { @@ -169,17 +170,9 @@ hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes, } size_t offset = 0; - amd::HostQueue* queue = nullptr; + amd::HostQueue* queue = hip::getQueue(stream); amd::Memory* memory = getMemoryObject(dst, offset); - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } - if (memory != nullptr) { // Device memory amd::Command::EventWaitList waitList; @@ -764,15 +757,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, kind, stream); - amd::HostQueue* queue = nullptr; - - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue, true)); } @@ -782,15 +767,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, stream); - amd::HostQueue* queue = nullptr; - - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice, *queue, true)); @@ -800,15 +777,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, stream); - amd::HostQueue* queue = nullptr; - - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice, *queue, true)); @@ -818,15 +787,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, src, sizeBytes, stream); - amd::HostQueue* queue = nullptr; - - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN(ihipMemcpy(reinterpret_cast(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost, *queue, true)); @@ -921,15 +882,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp size_t height, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind, stream); - amd::HostQueue* queue; - - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue, true)); } @@ -1349,14 +1302,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height, hipStream_t stream) { HIP_INIT_API(dst, pitch, value, width, height, stream); - amd::HostQueue* queue = nullptr; - if (stream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(stream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(stream); HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue, true)); } diff --git a/projects/clr/hipamd/api/hip/hip_module.cpp b/projects/clr/hipamd/api/hip/hip_module.cpp index 542cd89551..0cd590886c 100644 --- a/projects/clr/hipamd/api/hip/hip_module.cpp +++ b/projects/clr/hipamd/api/hip/hip_module.cpp @@ -218,6 +218,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0, uint32_t params = 0) { + HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, + sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params); + hip::Function* function = hip::Function::asFunction(f); amd::Kernel* kernel = function->function_; amd::Device* device = hip::getCurrentContext()->devices()[0]; @@ -226,14 +229,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, hip::Event* eStart = reinterpret_cast(startEvent); hip::Event* eStop = reinterpret_cast(stopEvent); - amd::HostQueue* queue; - if (hStream == nullptr) { - hip::syncStreams(); - queue = hip::getNullStream(); - } else { - hip::getNullStream()->finish(); - queue = reinterpret_cast(hStream)->asHostQueue(); - } + amd::HostQueue* queue = hip::getQueue(hStream); + if ((params & amd::NDRangeKernelCommand::CooperativeGroups) && !device->info().cooperativeGroups_) { return hipErrorLaunchFailure; diff --git a/projects/clr/hipamd/api/hip/hip_stream.cpp b/projects/clr/hipamd/api/hip/hip_stream.cpp index bc681d8a7e..ab26b4fc2e 100644 --- a/projects/clr/hipamd/api/hip/hip_stream.cpp +++ b/projects/clr/hipamd/api/hip/hip_stream.cpp @@ -173,21 +173,9 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { hipError_t hipStreamSynchronize(hipStream_t stream) { HIP_INIT_API(stream); - amd::HostQueue* hostQueue; + amd::HostQueue* hostQueue = hip::getQueue(stream); + hostQueue->finish(); - if (stream == nullptr) { - hip::syncStreams(); - - hostQueue = hip::getNullStream(); - - hostQueue->finish(); - } else { - hip::getNullStream()->finish(); - - hip::Stream* hStream = reinterpret_cast(stream); - - hStream->finish(); - } HIP_RETURN(hipSuccess); }