From 73bb9a74bb4e426b65a5bd1935a7203b3a963800 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 28 Mar 2019 02:21:45 +0530
Subject: [PATCH] Handle D2D in memcpy2D

[ROCm/clr commit: 50d623981e25918caec6b78143af2ebd5616d9d9]
---
 projects/clr/hipamd/src/hip_memory.cpp | 22 ++++++++++++++--------
 1 file changed, 14 insertions(+), 8 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index e3823c504e..8196be91ff 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -1586,21 +1586,24 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
     if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue;
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
-    int isLocked = 0;
+    int isLockedOrD2D = 0;
     void *pinnedPtr=NULL;
     void *actualSrc = (void*)src;
     void *actualDest = dst;
     if(kind == hipMemcpyHostToDevice ) {
         if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){
-            isLocked = 1;
+            isLockedOrD2D = 1;
             actualSrc = pinnedPtr;
         }
     } else if(kind == hipMemcpyDeviceToHost) {
           if(getLockedPointer((void*)dst, dpitch, &pinnedPtr) == hipSuccess ){
-            isLocked = 1;
+            isLockedOrD2D = 1;
             actualDest = pinnedPtr;
           }
+    } else if(kind == hipMemcpyDeviceToDevice) {
+        isLockedOrD2D = 1;
     }
+
     hc::completion_future marker;
 
     hipError_t e = hipSuccess;
@@ -1608,7 +1611,7 @@ hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch
         stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false);
     } else {
         try {
-            if(!isLocked) {
+            if(!isLockedOrD2D) {
                 for (int i = 0; i < height; ++i)
                     stream->locked_copySync((unsigned char*)dst + i * dpitch,
                                     (unsigned char*)src + i * spitch, width, kind);
@@ -1639,27 +1642,30 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
     HIP_INIT_SPECIAL_API(hipMemcpy2DAsync, (TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream);
     if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return ihipLogStatus(hipErrorInvalidValue);
     hipError_t e = hipSuccess;
-    int isLocked = 0;
+    int isLockedOrD2D = 0;
     void *pinnedPtr=NULL;
     void *actualSrc = (void*)src;
     void *actualDest = dst;
     stream = ihipSyncAndResolveStream(stream);
     if(kind == hipMemcpyHostToDevice ) {
         if(getLockedPointer((void*)src, spitch, &pinnedPtr) == hipSuccess ){
-            isLocked = 1;
+            isLockedOrD2D = 1;
             actualSrc = pinnedPtr;
         }
     } else if(kind == hipMemcpyDeviceToHost) {
           if(getLockedPointer((void*)dst, dpitch, &pinnedPtr) == hipSuccess ){
-            isLocked = 1;
+            isLockedOrD2D = 1;
             actualDest = pinnedPtr;
           }
+    } else if(kind == hipMemcpyDeviceToDevice) {
+        isLockedOrD2D = 1;
     }
+
     if((width == dpitch) && (width == spitch)) {
             hip_internal::memcpyAsync(dst, src, width*height, kind, stream);
     } else {
         try {
-            if(!isLocked){
+            if(!isLockedOrD2D){
                 for (int i = 0; i < height; ++i)
                     e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
                                           (unsigned char*)src + i * spitch, width, kind, stream);