SWDEV-243576 - Fix hipMemcpy regression.

If the queue device doesn't match the device on any memory object, use the queue device from the memory object. Change-Id: I5fdcf00494f8391574f4443332c01788b8da44ef
2020-07-21 12:10:43 -04:00
parent 78fcbab8c2
commit e430506f70
1 changed files with 10 additions and 1 deletions
@@ -196,7 +196,16 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
        return hipErrorInvalidValue;
      }
    } else {
-      command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
+      amd::HostQueue* pQueue = &queue;
+      if (queueDevice != srcMemory->getContext().devices()[0]) {
+        pQueue = hip::getNullStream(srcMemory->getContext());
+        amd::Command* cmd = queue.getLastQueuedCommand(true);
+        if (cmd != nullptr) {
+          waitList.push_back(cmd);
+        }
+      }
+
+      command = new amd::CopyMemoryCommand(*pQueue, CL_COMMAND_COPY_BUFFER, waitList,
          *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
    }
  }