From aed26538575d4a501134a17ba59002be58a5b442 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 24 May 2018 23:30:27 +0530
Subject: [PATCH] Clean up and fix remaining bytes copy

---
 hipamd/src/hip_memory.cpp | 22 +++++++++-------------
 1 file changed, 9 insertions(+), 13 deletions(-)

diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 2f0c793033..fa68d7cfb1 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -1496,19 +1496,15 @@ __global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height,
     size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
     size_t idy = blockIdx.y * blockDim.y + threadIdx.y;
     size_t floorWidth = (width/sizeof(T));
-    if((idx < width) && (idy < height)) {
-        if((idx < floorWidth)){
-            T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch);
-            T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch);
-            dstPtr[idx] = srcPtr[idx];
-        } else {
-            size_t bytesToCopy = width - (floorWidth * sizeof(T));
-            uint8_t *dstPtr = (uint8_t *) ((uint8_t*) dst + idy * destPitch);
-            uint8_t *srcPtr = (uint8_t *) ((uint8_t*) src + idy * srcPitch);
-            for(int i =0 ; i < bytesToCopy ; i++) {
-                dstPtr[idx+i]= srcPtr[idx+i];
-            }
-       }
+    T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch);
+    T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch);
+    if((idx < floorWidth) && (idy < height)){
+        dstPtr[idx] = srcPtr[idx];
+    } else if((idx < width) && (idy < height)){
+        size_t bytesToCopy = width - (floorWidth * sizeof(T));
+        dstPtr += floorWidth;
+        srcPtr += floorWidth;
+        __builtin_memcpy(reinterpret_cast<uint8_t*>(dstPtr), reinterpret_cast<const uint8_t*>(srcPtr),bytesToCopy);
     }
 }
 }  // namespace