SWDEV-408180 - Add a new hipMemcpyKind

Add hipMemcpyDeviceToDeviceNoCU to force a non blit copy path. This
helps in cases where an app may determine that CU may be busy and copies
with SDMA may be quicker.

Change-Id: I59b415dd8f6022c244e8d75f265464d5c635df1e


[ROCm/clr commit: f316a30e5d]

projects/clr/hipamd/src/hip_memory.cpp

@@ -379,7 +379,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
   size_t dOffset = 0;
   amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
   amd::Device* queueDevice = &stream.device();
-  amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA);
+  amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::NONE);
   if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
     hip::Stream* pStream = &stream;
     if (queueDevice != dstMemory->getContext().devices()[0]) {
@@ -424,7 +424,6 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
       hip::Stream* pStream = &stream;
       if ((srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) &&
           (queueDevice != srcMemory->getContext().devices()[0])) {
-        copyMetadata.copyEnginePreference_ = amd::CopyMetadata::CopyEnginePreference::NONE;
         pStream = hip::getNullStream(srcMemory->getContext());
         amd::Command* cmd = stream.getLastQueuedCommand(true);
         if (cmd != nullptr) {
@@ -449,6 +448,10 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
           }
         }
       }
+
+      copyMetadata.copyEnginePreference_ = (kind == hipMemcpyDeviceToDeviceNoCU) ?
+                                            amd::CopyMetadata::CopyEnginePreference::SDMA :
+                                            amd::CopyMetadata::CopyEnginePreference::NONE;
       command = new amd::CopyMemoryCommand(*pStream, CL_COMMAND_COPY_BUFFER, waitList,
           *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes,
           copyMetadata);
@@ -2807,7 +2810,7 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
       }
     }
   }
-  
+
   if (p->kind < hipMemcpyHostToHost || p->kind > hipMemcpyDefault) {
     return hipErrorInvalidMemcpyDirection;
   }

projects/clr/rocclr/device/rocm/rocblit.cpp

@@ -398,7 +398,7 @@ bool DmaBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& dstMe
     return HostBlitManager::copyBuffer(srcMemory, dstMemory, srcOrigin, dstOrigin, size, false,
                                        copyMetadata);
   } else {
-    return hsaCopy(gpuMem(srcMemory), gpuMem(dstMemory), srcOrigin, dstOrigin, size);
+    return hsaCopy(gpuMem(srcMemory), gpuMem(dstMemory), srcOrigin, dstOrigin, size, copyMetadata);
   }
 
   return true;
@@ -638,7 +638,7 @@ bool DmaBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMem
 // ================================================================================================
 bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
                              const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin,
-                             const amd::Coord3D& size) const {
+                             const amd::Coord3D& size, amd::CopyMetadata copyMetadata) const {
   address src = reinterpret_cast<address>(srcMemory.getDeviceMemory());
   address dst = reinterpret_cast<address>(dstMemory.getDeviceMemory());
 
@@ -684,10 +684,14 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
   uint32_t copyMask = 0;
   uint32_t freeEngineMask = 0;
   bool kUseRegularCopyApi = 0;
-
+  bool forceSDMA = (copyMetadata.copyEnginePreference_ ==
+                      amd::CopyMetadata::CopyEnginePreference::SDMA);
   HwQueueEngine engine = HwQueueEngine::Unknown;
+
+  // Determine engine and assign a copy mask for the new versatile ROCr API
+  // If engine preferred is SDMA, assign the SdmaWrite path
   if ((srcAgent.handle == dev().getCpuAgent().handle) &&
-      (dstAgent.handle != dev().getCpuAgent().handle)) {
+      (dstAgent.handle != dev().getCpuAgent().handle) || forceSDMA) {
     engine = HwQueueEngine::SdmaWrite;
     copyMask = kUseRegularCopyApi ? 0 : dev().fetchSDMAMask(this, false);
   } else if ((srcAgent.handle != dev().getCpuAgent().handle) &&
@@ -711,8 +715,8 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
       if (copyMask == 0) {
         // Check SDMA engine status
         status = hsa_amd_memory_copy_engine_status(dstAgent, srcAgent, &freeEngineMask);
-        ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Query copy engine status %x, free_engine mask 0x%x",
-                status, freeEngineMask);
+        ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Query copy engine status %x, "
+                "free_engine mask 0x%x", status, freeEngineMask);
         // Return a mask with the rightmost bit set
         copyMask = freeEngineMask - (freeEngineMask & (freeEngineMask - 1));
         gpu().setLastUsedSdmaEngine(copyMask);
@@ -725,13 +729,14 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
 
       ClPrint(amd::LOG_DEBUG, amd::LOG_COPY,
               "HSA Async Copy on copy_engine=0x%x, dst=0x%zx, src=0x%zx, "
-              "size=%ld, wait_event=0x%zx, completion_signal=0x%zx", copyEngine,
-              dst, src, size[0], (wait_events.size() != 0) ? wait_events[0].handle : 0,
+              "size=%ld, forceSDMA=%d, wait_event=0x%zx, completion_signal=0x%zx", copyEngine,
+              dst, src, size[0], forceSDMA, (wait_events.size() != 0) ? wait_events[0].handle : 0,
               active.handle);
 
       status = hsa_amd_memory_async_copy_on_engine(dst, dstAgent, src, srcAgent,
                                                   size[0], wait_events.size(),
-                                                  wait_events.data(), active, copyEngine, false);
+                                                  wait_events.data(), active, copyEngine,
+                                                  forceSDMA);
     } else {
       kUseRegularCopyApi = true;
     }
@@ -2258,7 +2263,9 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
   bool useShaderCopyPath = setup_.disableHwlCopyBuffer_ ||
                            (!srcMemory.isHostMemDirectAccess() &&
                             !dstMemory.isHostMemDirectAccess() &&
-                            !(p2p || asan) && !ipcShared);
+                            !(p2p || asan) && !ipcShared &&
+                            !(copyMetadata.copyEnginePreference_ ==
+                              amd::CopyMetadata::CopyEnginePreference::SDMA));
 
   if (!useShaderCopyPath) {
     if (amd::IS_HIP) {

projects/clr/rocclr/device/rocm/rocblit.hpp

@@ -230,7 +230,8 @@ class DmaBlitManager : public device::HostBlitManager {
   //! Assits in transferring data from Host to Local or vice versa
   //! taking into account the Hsail profile supported by Hsa Agent
   bool hsaCopy(const Memory& srcMemory, const Memory& dstMemory, const amd::Coord3D& srcOrigin,
-               const amd::Coord3D& dstOrigin, const amd::Coord3D& size) const;
+               const amd::Coord3D& dstOrigin, const amd::Coord3D& size,
+               amd::CopyMetadata copyMetadata) const;
 
   const size_t MinSizeForPinnedTransfer;
   bool completeOperation_;                    //!< DMA blit manager must complete operation