SWDEV-477324 - Capture Memcpy1D pinned H2D D2H

Change-Id: I1f4744f20a9caeed005ec68da44e5fde737e09f7


[ROCm/clr commit: 742b0210d3]

projects/clr/hipamd/src/hip_graph_helper.hpp

@@ -121,4 +121,6 @@ hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord
                                   size_t &start);
 
 hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams);
+
+hip::MemcpyType ihipGetMemcpyType(const void* src, void* dst, hipMemcpyKind kind);
 }  // namespace hip

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -1576,6 +1576,13 @@ class GraphMemcpyNode1D : public GraphMemcpyNode {
       WorkerThreadLock_.lock();
     }
     status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *stream);
+    hip::MemcpyType type = ihipGetMemcpyType(src_, dst_, kind_);
+    if (type == hipCopyBuffer) {
+      amd::CopyMemoryCommand* cpycmd = reinterpret_cast<amd::CopyMemoryCommand*>(command);
+      amd::CopyMetadata copyMetadata = cpycmd->copyMetadata();
+      copyMetadata.copyEnginePreference_ = amd::CopyMetadata::CopyEnginePreference::BLIT;
+      cpycmd->SetCopyMetadata(copyMetadata);
+    }
     if (!AMD_DIRECT_DISPATCH) {
       WorkerThreadLock_.unlock();
     }
@@ -1720,8 +1727,9 @@ class GraphMemcpyNode1D : public GraphMemcpyNode {
   virtual bool GraphCaptureEnabled() override {
     bool isGraphCapture = false;
     if (DEBUG_CLR_GRAPH_PACKET_CAPTURE) {
-      switch (kind_) {
-        case hipMemcpyDeviceToDevice:
+      hip::MemcpyType type = ihipGetMemcpyType(src_, dst_, kind_);
+      switch (type) {
+        case hipCopyBuffer:
           isGraphCapture = true;
           break;
         default:

projects/clr/hipamd/src/hip_internal.hpp

@@ -69,6 +69,21 @@ typedef struct hipArray {
 }hipArray;
 
 namespace hip {
+enum MemcpyType {
+  /// Memcpy from host to host
+  hipHostToHost,
+  /// Memcpy from host to device
+  hipWriteBuffer,
+  /// Memcpy from device to host
+  hipReadBuffer,
+  /// Memcpy from device A to device A
+  /// Memcpy from pinned host buffer to device/device to pinned host buffer
+  hipCopyBuffer,
+  /// Memcpy from device A to device A, user forced to SDMA
+  hipCopyBufferSDMA,
+  /// Memcpy from device A to device B
+  hipCopyBufferP2P,
+};
 struct Graph;
 struct GraphNode;
 struct GraphExec;

projects/clr/hipamd/src/hip_memory.cpp

@@ -463,6 +463,35 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
   return hipSuccess;
 }
 
+hip::MemcpyType ihipGetMemcpyType(const void* src, void* dst, hipMemcpyKind kind) {
+  size_t sOffset = 0;
+  amd::Memory* srcMemory = getMemoryObject(src, sOffset);
+  size_t dOffset = 0;
+  amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
+  hip::MemcpyType type;
+  if(srcMemory == nullptr && dstMemory == nullptr) {
+    type = hipHostToHost;
+  } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
+    type = hipWriteBuffer;
+  } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
+    type = hipReadBuffer;
+  } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
+    // Check if the queue device doesn't match the device on any memory object.
+    // And any of them are not host allocation.
+    // Hence it's a P2P transfer, because the app has requested access to another GPU
+    if ((srcMemory->GetDeviceById() != dstMemory->GetDeviceById()) &&
+        ((srcMemory->getContext().devices().size() == 1) &&
+         (dstMemory->getContext().devices().size() == 1))) {
+      type = hipCopyBufferP2P;
+    } else if( kind == hipMemcpyDeviceToDeviceNoCU) {
+      type = hipCopyBufferSDMA;
+    } else {
+      type = hipCopyBuffer;
+    }
+  }
+  return type;
+}
+
 // ================================================================================================
 hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes,
                              hipMemcpyKind kind, hip::Stream& stream, bool isAsync) {
@@ -473,37 +502,37 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
   amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
   amd::Device* queueDevice = &stream.device();
   amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::NONE);
-  if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
-    hip::Stream* pStream = &stream;
-    if (queueDevice != dstMemory->GetDeviceById()) {
-      pStream = hip::getNullStream(dstMemory->GetDeviceById()->context());
-      amd::Command* cmd = stream.getLastQueuedCommand(true);
-      if (cmd != nullptr) {
-        waitList.push_back(cmd);
+  hip::MemcpyType type = ihipGetMemcpyType(src, dst, kind);
+  hip::Stream* pStream = &stream;
+  switch (type) {
+    case hipWriteBuffer:
+      if (queueDevice != dstMemory->GetDeviceById()) {
+        pStream = hip::getNullStream(dstMemory->GetDeviceById()->context());
+        amd::Command* cmd = stream.getLastQueuedCommand(true);
+        if (cmd != nullptr) {
+          waitList.push_back(cmd);
+        }
       }
-    }
-    command = new amd::WriteMemoryCommand(*pStream, CL_COMMAND_WRITE_BUFFER, waitList,
-              *dstMemory->asBuffer(), dOffset, sizeBytes, src, 0, 0, copyMetadata);
-  } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
-    hip::Stream* pStream = &stream;
-    if (queueDevice != srcMemory->GetDeviceById()) {
-      pStream = hip::getNullStream(srcMemory->GetDeviceById()->context());
-      amd::Command* cmd = stream.getLastQueuedCommand(true);
-      if (cmd != nullptr) {
-        waitList.push_back(cmd);
+      command = new amd::WriteMemoryCommand(*pStream, CL_COMMAND_WRITE_BUFFER, waitList,
+                                            *dstMemory->asBuffer(), dOffset, sizeBytes, src, 0, 0,
+                                            copyMetadata);
+      break;
+    case hipReadBuffer:
+      if (queueDevice != srcMemory->GetDeviceById()) {
+        pStream = hip::getNullStream(srcMemory->GetDeviceById()->context());
+        amd::Command* cmd = stream.getLastQueuedCommand(true);
+        if (cmd != nullptr) {
+          waitList.push_back(cmd);
+        }
       }
-    }
-    command = new amd::ReadMemoryCommand(*pStream, CL_COMMAND_READ_BUFFER, waitList,
-              *srcMemory->asBuffer(), sOffset, sizeBytes, dst, 0, 0, copyMetadata);
-  } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
-    // Check if the queue device doesn't match the device on any memory object.
-    // And any of them are not host allocation.
-    // Hence it's a P2P transfer, because the app has requested access to another GPU
-    if ((srcMemory->GetDeviceById() != dstMemory->GetDeviceById()) &&
-        ((srcMemory->getContext().devices().size() == 1) &&
-         (dstMemory->getContext().devices().size() == 1))) {
+      command = new amd::ReadMemoryCommand(*pStream, CL_COMMAND_READ_BUFFER, waitList,
+                                           *srcMemory->asBuffer(), sOffset, sizeBytes, dst, 0, 0,
+                                           copyMetadata);
+      break;
+    case hipCopyBufferP2P:
       command = new amd::CopyMemoryP2PCommand(stream, CL_COMMAND_COPY_BUFFER, waitList,
-          *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
+                                              *srcMemory->asBuffer(), *dstMemory->asBuffer(),
+                                              sOffset, dOffset, sizeBytes);
       if (command == nullptr) {
         return hipErrorOutOfMemory;
       }
@@ -513,10 +542,12 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
         delete command;
         return hipErrorInvalidValue;
       }
-    } else {
-      hip::Stream* pStream = &stream;
+      break;
+    case hipCopyBufferSDMA:
+      copyMetadata.copyEnginePreference_ = amd::CopyMetadata::CopyEnginePreference::SDMA;
+    case hipCopyBuffer:
       if ((srcMemory->GetDeviceById() == dstMemory->GetDeviceById()) &&
-          (queueDevice != srcMemory->GetDeviceById())) {
+          queueDevice != srcMemory->GetDeviceById()) {
         pStream = hip::getNullStream(srcMemory->GetDeviceById()->context());
         amd::Command* cmd = stream.getLastQueuedCommand(true);
         if (cmd != nullptr) {
@@ -531,7 +562,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
           if (cmd != nullptr) {
             waitList.push_back(cmd);
           }
-        // Scenarios such as HtoD where src is pinned memory
+          // Scenarios such as HtoD where src is pinned memory
         } else if ((queueDevice != dstMemory->GetDeviceById()) &&
                    (srcMemory->getContext().devices().size() != 1)) {
           pStream = hip::getNullStream(dstMemory->GetDeviceById()->context());
@@ -541,14 +572,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);
-    }
+                                           *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset,
+                                           dOffset, sizeBytes, copyMetadata);
+      break;
   }
   if (command == nullptr) {
     return hipErrorOutOfMemory;
@@ -558,6 +585,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
   }
   return hipSuccess;
 }
+
 bool IsHtoHMemcpy(void* dst, const void* src) {
   size_t sOffset = 0;
   amd::Memory* srcMemory = getMemoryObject(src, sOffset);

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

@@ -2263,12 +2263,12 @@ bool KernelBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& ds
 #endif
 
   bool useShaderCopyPath = setup_.disableHwlCopyBuffer_ ||
-                           (sizeIn[0] <= dev().settings().sdmaCopyThreshold_) ||
-                           (!srcMemory.isHostMemDirectAccess() &&
-                            !dstMemory.isHostMemDirectAccess() &&
-                            !(p2p || asan) && !ipcShared &&
-                            !(copyMetadata.copyEnginePreference_ ==
-                              amd::CopyMetadata::CopyEnginePreference::SDMA));
+      (sizeIn[0] <= dev().settings().sdmaCopyThreshold_) ||
+      (!(p2p || asan || ipcShared) &&
+           (!srcMemory.isHostMemDirectAccess() && !dstMemory.isHostMemDirectAccess() &&
+            !(copyMetadata.copyEnginePreference_ ==
+              amd::CopyMetadata::CopyEnginePreference::SDMA)) ||
+       (copyMetadata.copyEnginePreference_ == amd::CopyMetadata::CopyEnginePreference::BLIT));
 
   if (!useShaderCopyPath) {
     if (amd::IS_HIP) {

projects/clr/rocclr/platform/command.hpp

@@ -952,6 +952,8 @@ class CopyMemoryCommand : public TwoMemoryArgsCommand {
   const BufferRect& dstRect() const { return dstRect_; }
   //! Return the copy MetaData
   amd::CopyMetadata copyMetadata() const { return copyMetadata_; }
+  //! Updates copy MetaData
+  void SetCopyMetadata(amd::CopyMetadata copyMetadata) { copyMetadata_ = copyMetadata; }
   //! Updates the host memory to read from
   void setSource(Memory& srcMemory) { memory1_ = &srcMemory; }
   //! Updates the memory object to write to.