From 1bd640b6590687c820486f3220b614e7a222a0bc Mon Sep 17 00:00:00 2001 From: Vladislav Sytchenko Date: Mon, 6 Apr 2020 13:03:32 -0400 Subject: [PATCH] (SWDEV-228488) These fixes address regressions caused by http://gerrit-git.amd.com/c/compute/ec/hip/+/337601 Currently we're converting a 1D offset into a 3D offset, which doesn't make much sense once you consider the fact that this offset is relative to a different origin than our current 3D offset. I traced through our blit kernels in VDI - the copy buffer rect path is able to handle immediate offsets in the 3D buffer via the amd::BufferRect::start_ parameter. Instead of adjusting the offset, simply adjust the start of the region. Change-Id: Ic8797a2c8ac0ad106f246f61ff06ca1ca03d3058 --- vdi/hip_memory.cpp | 31 ++++++++++++------------------- 1 file changed, 12 insertions(+), 19 deletions(-) diff --git a/vdi/hip_memory.cpp b/vdi/hip_memory.cpp index afe67a867d..1246f53fe8 100644 --- a/vdi/hip_memory.cpp +++ b/vdi/hip_memory.cpp @@ -884,19 +884,6 @@ hipError_t hipMemcpyDtoHAsync(void* dstHost, HIP_RETURN(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *hip::getQueue(stream), true)); } -inline void adjustOrigin(amd::Coord3D &origin, - size_t offset, - size_t rowPitch, - size_t slicePitch) { - size_t zOffset = offset / (slicePitch ? slicePitch : 1); - size_t yOffset = (offset - slicePitch * zOffset) / (rowPitch ? rowPitch : 1); - size_t xOffset = (offset - slicePitch * zOffset - rowPitch * yOffset); - - static_cast(origin)[0] += xOffset; - static_cast(origin)[1] += yOffset; - static_cast(origin)[2] += zOffset; -} - hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin, @@ -914,7 +901,6 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray, amd::Image* srcImage = as_amd(srcMemObj)->asImage(); size_t dstOffset = 0; amd::Memory* dstMemory = getMemoryObject(dstDevice, dstOffset); - adjustOrigin(dstOrigin, dstOffset, dstRowPitch, dstSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcImage->getRowPitch(), srcImage->getSlicePitch())) { @@ -925,6 +911,8 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray, if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } + dstRect.start_ += dstOffset; + dstRect.end_ += dstOffset; const size_t copySizeInBytes = copyRegion[0] * copyRegion[1] * copyRegion[2] * srcImage->getImageFormat().getElementSize(); if (!srcImage->validateRegion(srcOrigin, copyRegion) || @@ -972,13 +960,14 @@ hipError_t ihipMemcpyDtoA(void* srcDevice, size_t srcOffset = 0; amd::Memory* srcMemory = getMemoryObject(srcDevice, srcOffset); - adjustOrigin(srcOrigin, srcOffset, srcRowPitch, srcSlicePitch); amd::Image* dstImage = as_amd(dstMemObj)->asImage(); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } + srcRect.start_ += srcOffset; + srcRect.end_ += srcOffset; amd::BufferRect dstRect; if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstImage->getRowPitch(), dstImage->getSlicePitch())) { @@ -1028,15 +1017,15 @@ hipError_t ihipMemcpyDtoD(void* srcDevice, bool isAsync = false) { size_t srcOffset = 0; amd::Memory *srcMemory = getMemoryObject(srcDevice, srcOffset); - adjustOrigin(srcOrigin, srcOffset, srcRowPitch, srcSlicePitch); size_t dstOffset = 0; amd::Memory *dstMemory = getMemoryObject(dstDevice, dstOffset); - adjustOrigin(dstOrigin, dstOffset, dstRowPitch, dstSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } + srcRect.start_ += srcOffset; + srcRect.end_ += srcOffset; amd::Coord3D srcStart(srcRect.start_, 0, 0); amd::Coord3D srcEnd(srcRect.end_, 1, 1); @@ -1048,6 +1037,8 @@ hipError_t ihipMemcpyDtoD(void* srcDevice, if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } + dstRect.start_ += dstOffset; + dstRect.end_ += dstOffset; amd::Coord3D dstStart(dstRect.start_, 0, 0); amd::Coord3D dstEnd(dstRect.end_, 1, 1); @@ -1092,12 +1083,13 @@ hipError_t ihipMemcpyDtoH(void* srcDevice, bool isAsync = false) { size_t srcOffset = 0; amd::Memory *srcMemory = getMemoryObject(srcDevice, srcOffset); - adjustOrigin(srcOrigin, srcOffset, srcRowPitch, srcSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } + srcRect.start_ += srcOffset; + srcRect.end_ += srcOffset; amd::Coord3D srcStart(srcRect.start_, 0, 0); amd::Coord3D srcEnd(srcRect.end_, 1, 1); @@ -1146,7 +1138,6 @@ hipError_t ihipMemcpyHtoD(const void* srcHost, bool isAsync = false) { size_t dstOffset = 0; amd::Memory *dstMemory = getMemoryObject(dstDevice, dstOffset); - adjustOrigin(dstOrigin, dstOffset, dstRowPitch, dstSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { @@ -1157,6 +1148,8 @@ hipError_t ihipMemcpyHtoD(const void* srcHost, if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } + dstRect.start_ += dstOffset; + dstRect.end_ += dstOffset; amd::Coord3D dstStart(dstRect.start_, 0, 0); amd::Coord3D dstEnd(dstRect.end_, 1, 1);