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(srcOrigin)[0] /= elementSize; + static_cast(copyRegion)[0] /= elementSize; amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(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(dstOrigin)[0] /= elementSize; + static_cast(copyRegion)[0] /= elementSize; amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) {