From bde082af6eb66f0dd226168f0e240478bf95604e Mon Sep 17 00:00:00 2001
From: Vladislav Sytchenko <Vladislav.Sytchenko@amd.com>
Date: Mon, 15 Mar 2021 21:09:58 -0400
Subject: [PATCH] SWDEV-275317 - Correctly pass HIP image width to ROCclr APIs
 in AtoD/DtoA copies

HIP assumes that image width is in bytes, but OCL/ROCclr assumes that
it's in pixels. AtoD/DtoA need to account for this.

Change-Id: I275bd41d8b03e141caaf951bc6b714e51ca72dfc


[ROCm/hip commit: e967adbdaf4de8bc11a529cc31d5c127b1087faa]
---
 projects/hip/rocclr/hip_memory.cpp | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/projects/hip/rocclr/hip_memory.cpp b/projects/hip/rocclr/hip_memory.cpp
index 694c2c515b..2500bc227a 100755
--- a/projects/hip/rocclr/hip_memory.cpp
+++ b/projects/hip/rocclr/hip_memory.cpp
@@ -971,6 +971,10 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray,
   }
 
   amd::Image* srcImage = as_amd(srcMemObj)->asImage();
+  // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
+  const size_t elementSize = srcImage->getImageFormat().getElementSize();
+  static_cast<size_t*>(srcOrigin)[0] /= elementSize;
+  static_cast<size_t*>(copyRegion)[0] /= elementSize;
 
   amd::BufferRect srcRect;
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion), srcImage->getRowPitch(), srcImage->getSlicePitch())) {
@@ -1031,6 +1035,10 @@ hipError_t ihipMemcpyDtoA(void* srcDevice,
   }
 
   amd::Image* dstImage = as_amd(dstMemObj)->asImage();
+  // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
+  const size_t elementSize = dstImage->getImageFormat().getElementSize();
+  static_cast<size_t*>(dstOrigin)[0] /= elementSize;
+  static_cast<size_t*>(copyRegion)[0] /= elementSize;
 
   amd::BufferRect srcRect;
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion), srcRowPitch, srcSlicePitch)) {