From 57fa5938fe7181862642adffb388a1ba2cbbb478 Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Mon, 27 Feb 2023 17:23:01 +0000 Subject: [PATCH] SWDEV-383056 - Don't sync with dst device for hipMemcpyAsync. Change-Id: I28530e6bd870d617507592576295fc9e7eed1475 --- hipamd/src/hip_code_object.cpp | 2 -- hipamd/src/hip_internal.hpp | 3 ++- hipamd/src/hip_memory.cpp | 19 +++++++------------ hipamd/src/hip_peer.cpp | 9 +++++++-- hipamd/src/hip_platform.cpp | 2 -- hipamd/src/hip_texture.cpp | 3 --- 6 files changed, 16 insertions(+), 22 deletions(-) diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index 778783cfde..dd9de6374f 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -31,8 +31,6 @@ THE SOFTWARE. #include "platform/program.hpp" #include -hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hip::Stream& stream, bool isAsync = false); hipError_t ihipFree(void* ptr); // forward declaration of methods required for managed variables hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0); diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 49a8832650..ca924bd5f8 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -565,7 +565,8 @@ extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t de extern hipError_t ihipDeviceGet(hipDevice_t* device, int deviceId); extern hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr, uint64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes); - +hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, + hip::Stream& stream, bool isHostAsync = false, bool isGPUAsync = true); constexpr bool kOptionChangeable = true; constexpr bool kNewDevProg = false; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index d642354ab7..7929543756 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -451,7 +451,7 @@ void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, hip::Stream& s } // ================================================================================================ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hip::Stream& stream, bool isAsync = false) { + hip::Stream& stream, bool isHostAsync, bool isGPUAsync) { hipError_t status; if (sizeBytes == 0) { // Skip if nothing needs writing. @@ -464,7 +464,6 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin if (src == dst && kind == hipMemcpyDefault) { return hipSuccess; } - bool isP2P = false; size_t sOffset = 0; amd::Memory* srcMemory = getMemoryObject(src, sOffset); size_t dOffset = 0; @@ -473,24 +472,20 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin ihipHtoHMemcpy(dst, src, sizeBytes, stream); return hipSuccess; } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) { - isAsync = false; + isHostAsync = false; } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { - isAsync = false; - } else if ((srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) && - (srcMemory->getContext().devices().size() == 1) && - (dstMemory->getContext().devices().size() == 1)) { - isAsync = true; - isP2P = true; + isHostAsync = false; } + amd::Command* command = nullptr; - status = ihipMemcpyCommand(command, dst, src, sizeBytes, kind, stream, isAsync); + status = ihipMemcpyCommand(command, dst, src, sizeBytes, kind, stream, isHostAsync); if (status != hipSuccess) { return status; } command->enqueue(); - if (!isAsync) { + if (!isHostAsync) { command->awaitCompletion(); - } else if (isP2P) { + } else if (!isGPUAsync) { hip::Stream* pStream = hip::getNullStream(dstMemory->getContext()); amd::Command::EventWaitList waitList; waitList.push_back(command); diff --git a/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp index d5255103a6..17dc65da05 100644 --- a/hipamd/src/hip_peer.cpp +++ b/hipamd/src/hip_peer.cpp @@ -220,7 +220,8 @@ hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevic HIP_RETURN(hipErrorInvalidDevice); } - HIP_RETURN(hipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice)); + HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice, *hip::getNullStream(), + true, false)); } hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, @@ -235,7 +236,11 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } - HIP_RETURN(hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream)); + hip::Stream* hip_stream = hip::getStream(stream); + if (hip_stream == nullptr) { + return hipErrorInvalidValue; + } + HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice, *hip_stream, true, true)); } hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp index 84ba32ffb8..10b8d3f894 100644 --- a/hipamd/src/hip_platform.cpp +++ b/hipamd/src/hip_platform.cpp @@ -33,8 +33,6 @@ PlatformState* PlatformState::platform_; // Initiaized as nullptr by default // forward declaration of methods required for __hipRegisrterManagedVar hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0); -hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hip::Stream& stream, bool isAsync = false); struct __CudaFatBinaryWrapper { unsigned int magic; diff --git a/hipamd/src/hip_texture.cpp b/hipamd/src/hip_texture.cpp index 9fead8d84a..8c44373901 100644 --- a/hipamd/src/hip_texture.cpp +++ b/hipamd/src/hip_texture.cpp @@ -25,9 +25,6 @@ #include "hip_conversions.hpp" #include "platform/sampler.hpp" -hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hip::Stream& stream, bool isAsync = false); - hipError_t ihipFree(void* ptr); struct __hip_texture {