From 9a62b2ffe71d1a7bead2e266b35e158ccd2c7989 Mon Sep 17 00:00:00 2001
From: Vladislav Sytchenko <vladislav.sytchenko@amd.com>
Date: Thu, 5 Mar 2020 14:00:43 -0500
Subject: [PATCH] Correct logic in ihipMemcpyAtoA()

HIP assumes that image width is in bytes, but VDI assumes that image width in pixels. Need to perform byte -> pixel coversion before doing anything.

Change-Id: Ia9fd1f46d05db3fbe8049add10b4d7e5118a2b9a


[ROCm/hip commit: 3fc50731ae2793cb8c93f2a5c2c0108eba7ebac1]
---
 projects/hip/vdi/hip_memory.cpp | 16 ++++++++++++----
 1 file changed, 12 insertions(+), 4 deletions(-)

diff --git a/projects/hip/vdi/hip_memory.cpp b/projects/hip/vdi/hip_memory.cpp
index 07bbb7f71c..cc3908ca64 100644
--- a/projects/hip/vdi/hip_memory.cpp
+++ b/projects/hip/vdi/hip_memory.cpp
@@ -1171,7 +1171,7 @@ hipError_t ihipMemcpyAtoA(hipArray* srcArray,
                           hipArray* dstArray,
                           amd::Coord3D srcOrigin,
                           amd::Coord3D dstOrigin,
-                          amd::Coord3D region,
+                          amd::Coord3D copyRegion,
                           hipStream_t stream,
                           bool isAsync = false) {
   cl_mem srcMemObj = reinterpret_cast<cl_mem>(srcArray->data);
@@ -1183,8 +1183,16 @@ hipError_t ihipMemcpyAtoA(hipArray* srcArray,
   amd::Image* srcImage = as_amd(srcMemObj)->asImage();
   amd::Image* dstImage = as_amd(dstMemObj)->asImage();
 
-  if (!srcImage->validateRegion(srcOrigin, region) ||
-      !dstImage->validateRegion(dstOrigin, region)) {
+  // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
+  // Note that src and dst should have the same element size.
+  assert(srcImage->getImageFormat().getElementSize() == dstImage->getImageFormat().getElementSize());
+  const size_t elementSize = srcImage->getImageFormat().getElementSize();
+  static_cast<size_t*>(srcOrigin)[0] /= elementSize;
+  static_cast<size_t*>(dstOrigin)[0] /= elementSize;
+  static_cast<size_t*>(copyRegion)[0] /= elementSize;
+
+  if (!srcImage->validateRegion(srcOrigin, copyRegion) ||
+      !dstImage->validateRegion(dstOrigin, copyRegion)) {
     return hipErrorInvalidValue;
   }
 
@@ -1195,7 +1203,7 @@ hipError_t ihipMemcpyAtoA(hipArray* srcArray,
                                                                *dstImage,
                                                                srcOrigin,
                                                                dstOrigin,
-                                                               region);
+                                                               copyRegion);
 
   if (command == nullptr) {
     return hipErrorOutOfMemory;