SWDEV-383056 - Don't sync with dst device for hipMemcpyAsync.

Change-Id: I28530e6bd870d617507592576295fc9e7eed1475

hipamd/src/hip_code_object.cpp

@@ -31,8 +31,6 @@ THE SOFTWARE.
 #include "platform/program.hpp"
 #include <elf/elf.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);
 // forward declaration of methods required for managed variables
 hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0);

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;
 

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);

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) {

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;

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 {