From 28524520e69a8650d827d92b318efec1d5f5dbe3 Mon Sep 17 00:00:00 2001 From: taosang2 Date: Thu, 28 Sep 2023 09:46:19 -0400 Subject: [PATCH] SWDEV-364236 - Fix layered Image issue Fix wrong logic to get layer index; Make layered image's layout match cuda spec; Fix wrong comparision of element size. Remove amd::BufferRect from ihipMemcpyAtoHCommand() and ihipMemcpyHtoACommand(). Change-Id: Icc6a4233fbce2e9b2dc6feb79e6bfbd761684c7d [ROCm/clr commit: 5a0085e5166640b1a93822454aa6652335740de4] --- projects/clr/hipamd/src/hip_conversions.hpp | 34 ++++++++++++ projects/clr/hipamd/src/hip_graph_helper.hpp | 4 +- .../clr/hipamd/src/hip_graph_internal.cpp | 9 ++-- projects/clr/hipamd/src/hip_memory.cpp | 52 ++++++++++++------- .../clr/rocclr/device/pal/palresource.cpp | 2 +- 5 files changed, 73 insertions(+), 28 deletions(-) diff --git a/projects/clr/hipamd/src/hip_conversions.hpp b/projects/clr/hipamd/src/hip_conversions.hpp index 2296a73516..17866d10b4 100644 --- a/projects/clr/hipamd/src/hip_conversions.hpp +++ b/projects/clr/hipamd/src/hip_conversions.hpp @@ -116,6 +116,16 @@ cl_mem_object_type getCLMemObjectType(const unsigned int hipWidth, return CL_MEM_OBJECT_ALLOCATION_FAILURE; } +inline cl_mem_object_type getCLMemObjectType(const hipArray* arr) { + const cl_mem dstMemObj = reinterpret_cast(arr->data); + const amd::Image* dstImage = as_amd(dstMemObj)->asImage(); + return dstImage ? dstImage->getType() : CL_MEM_OBJECT_ALLOCATION_FAILURE; +} + +inline bool isLayered1D(const hipArray* arr) { + return CL_MEM_OBJECT_IMAGE1D_ARRAY == getCLMemObjectType(arr); +} + inline cl_addressing_mode getCLAddressingMode(const hipTextureAddressMode hipAddressMode) { switch (hipAddressMode) { @@ -654,11 +664,15 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.dstZ = desc.dstPos.z; descDrv.dstLOD = 0; + bool is1DArraySrc = false; + bool is1DArrayDst = false; + if (desc.srcArray != nullptr) { descDrv.srcMemoryType = hipMemoryTypeArray; descDrv.srcArray = desc.srcArray; // When reffering to array memory, hipPos::x is in elements. descDrv.srcXInBytes *= getElementSize(desc.srcArray); + is1DArraySrc = isLayered1D(descDrv.srcArray); } if (desc.srcPtr.ptr != nullptr) { @@ -674,6 +688,7 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.dstArray = desc.dstArray; // When reffering to array memory, hipPos::x is in elements. descDrv.dstXInBytes *= getElementSize(desc.dstArray); + is1DArrayDst = isLayered1D(descDrv.dstArray); } if (desc.dstPtr.ptr != nullptr) { @@ -693,6 +708,25 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.WidthInBytes *= getElementSize(desc.dstArray); } + // The following will happen in + // buffer to/from layered 1D, layered 1D to/from layered 1D + size_t t = 0; + if (is1DArraySrc) { + t = descDrv.srcY; + descDrv.srcY = descDrv.srcZ; + descDrv.srcZ = t; + } + if (is1DArrayDst) { + t = descDrv.dstY; + descDrv.dstY = descDrv.dstZ; + descDrv.dstZ = t; + } + + if (is1DArraySrc || is1DArrayDst) { + t = descDrv.Height; + descDrv.Height = descDrv.Depth; + descDrv.Depth = t; + } return descDrv; } diff --git a/projects/clr/hipamd/src/hip_graph_helper.hpp b/projects/clr/hipamd/src/hip_graph_helper.hpp index b195267395..63f3ae139f 100644 --- a/projects/clr/hipamd/src/hip_graph_helper.hpp +++ b/projects/clr/hipamd/src/hip_graph_helper.hpp @@ -106,11 +106,11 @@ hipError_t ihipMemcpyAtoAValidate(hipArray_t srcArray, hipArray_t dstArray, amd: hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage, - amd::BufferRect& srcRect); + size_t &start); hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t dstRowPitch, size_t dstSlicePitch, amd::Image*& srcImage, - amd::BufferRect& dstRect); + size_t &start); hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams); diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index 5f5f40d4b6..9e7627974b 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -201,22 +201,21 @@ hipError_t GraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParams) } } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) { amd::Image* dstImage; - amd::BufferRect srcRect; - + size_t start = 0; status = ihipMemcpyHtoAValidate(pCopy.srcHost, pCopy.dstArray, srcOrigin, dstOrigin, copyRegion, - pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, dstImage, srcRect); + pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, dstImage, start); if (status != hipSuccess) { return status; } } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) { // Image to Host. amd::Image* srcImage; - amd::BufferRect dstRect; + size_t start = 0; status = ihipMemcpyAtoHValidate(pCopy.srcArray, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion, - pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcImage, dstRect); + pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcImage, start); if (status != hipSuccess) { return status; } diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 0d7af8267e..64195f2a48 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1543,13 +1543,15 @@ hipError_t ihipMemcpyAtoDValidate(hipArray_t srcArray, void* dstDevice, amd::Coo static_cast(srcOrigin)[0] /= elementSize; static_cast(copyRegion)[0] /= elementSize; + amd::Coord3D copyRegionReal = copyRegion; + if (hip::isLayered1D(srcArray)) copyRegionReal.c[1] = 1; - if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegionReal), srcImage->getRowPitch(), srcImage->getSlicePitch())) { return hipErrorInvalidValue; } - if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegionReal), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } @@ -1616,14 +1618,16 @@ hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray_t dstArray, amd::Coo static_cast(dstOrigin)[0] /= elementSize; static_cast(copyRegion)[0] /= elementSize; + amd::Coord3D copyRegionReal = copyRegion; + if (hip::isLayered1D(dstArray)) copyRegionReal.c[1] = 1; - if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegionReal), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } srcRect.start_ += srcOffset; - if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), + if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegionReal), dstImage->getRowPitch(), dstImage->getSlicePitch())) { return hipErrorInvalidValue; } @@ -1965,11 +1969,22 @@ hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray_t srcArray, hi return hipSuccess; } +size_t ihipGetbufferStart(const size_t* bufferOrigin, const size_t* region, size_t bufferRowPitch, + size_t bufferSlicePitch) { + size_t rowPitch_ = (bufferRowPitch != 0) ? bufferRowPitch : region[0]; + // Find the buffer's slice pitch + size_t slicePitch_ = (bufferSlicePitch != 0) ? bufferSlicePitch : rowPitch_ * region[1]; + // Find the region start offset + // For the layered 1D image, bufferOrigin[2] is always 0; for 1D image, bufferOrigin[1] is + // always 0. So the logic still holds true for all cases. + return bufferOrigin[2] * slicePitch_ + bufferOrigin[1] * rowPitch_ + bufferOrigin[0]; +} + hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray_t dstArray, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage, - amd::BufferRect& srcRect) { + size_t &start) { if ((srcHost == nullptr) || dstArray == nullptr) { return hipErrorInvalidValue; } @@ -1978,10 +1993,8 @@ hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray_t dstArray, return hipErrorInvalidValue; } - if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), - srcRowPitch, srcSlicePitch)) { - return hipErrorInvalidValue; - } + start = ihipGetbufferStart(static_cast(srcOrigin), static_cast(copyRegion), + srcRowPitch, srcSlicePitch); dstImage = as_amd(dstMemObj)->asImage(); // HIP assumes the width is in bytes, but OCL assumes it's in pixels. @@ -2000,10 +2013,10 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch, hip::Stream* stream, bool isAsync = false) { amd::Image* dstImage; - amd::BufferRect srcRect; + size_t start = 0; //!< Start offset for the copy region hipError_t status = ihipMemcpyHtoAValidate(srcHost, dstArray, srcOrigin, dstOrigin, copyRegion, - srcRowPitch, srcSlicePitch, dstImage, srcRect); + srcRowPitch, srcSlicePitch, dstImage, start); if (status != hipSuccess) { return status; } @@ -2011,7 +2024,7 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA); amd::WriteMemoryCommand* writeMemCmd = new amd::WriteMemoryCommand( *stream, CL_COMMAND_WRITE_IMAGE, amd::Command::EventWaitList{}, *dstImage, dstOrigin, - copyRegion, static_cast(srcHost) + srcRect.start_, srcRowPitch, srcSlicePitch, + copyRegion, static_cast(srcHost) + start, srcRowPitch, srcSlicePitch, copyMetadata); if (writeMemCmd == nullptr) { @@ -2029,7 +2042,7 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord3D& srcOrigin, amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion, size_t dstRowPitch, size_t dstSlicePitch, amd::Image*& srcImage, - amd::BufferRect& dstRect) { + size_t &start) { if (srcArray == nullptr || (dstHost == nullptr)) { return hipErrorInvalidValue; } @@ -2038,10 +2051,8 @@ hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord return hipErrorInvalidValue; } - if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), - dstRowPitch, dstSlicePitch)) { - return hipErrorInvalidValue; - } + start = ihipGetbufferStart(static_cast(dstOrigin), static_cast(copyRegion), + dstRowPitch, dstSlicePitch); srcImage = as_amd(srcMemObj)->asImage(); // HIP assumes the width is in bytes, but OCL assumes it's in pixels. @@ -2063,16 +2074,17 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray_t srcArray, vo amd::Image* srcImage; amd::BufferRect dstRect; amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA); + size_t start = 0; //!< Start offset for the copy region hipError_t status = ihipMemcpyAtoHValidate(srcArray, dstHost, srcOrigin, dstOrigin, copyRegion, - dstRowPitch, dstSlicePitch, srcImage, dstRect); + dstRowPitch, dstSlicePitch, srcImage, start); if (status != hipSuccess) { return status; } amd::ReadMemoryCommand* readMemCmd = new amd::ReadMemoryCommand( *stream, CL_COMMAND_READ_IMAGE, amd::Command::EventWaitList{}, *srcImage, srcOrigin, - copyRegion, static_cast(dstHost) + dstRect.start_, dstRowPitch, dstSlicePitch, + copyRegion, static_cast(dstHost) + start, dstRowPitch, dstSlicePitch, copyMetadata); if (readMemCmd == nullptr) { @@ -2732,7 +2744,7 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) { // If the source and destination are both arrays, hipMemcpy3D() will return an error if they do // not have the same element size. if (((p->srcArray != nullptr) && (p->dstArray != nullptr)) && - (hip::getElementSize(p->dstArray) != hip::getElementSize(p->dstArray))) { + (hip::getElementSize(p->srcArray) != hip::getElementSize(p->dstArray))) { return hipErrorInvalidValue; } diff --git a/projects/clr/rocclr/device/pal/palresource.cpp b/projects/clr/rocclr/device/pal/palresource.cpp index 42b29b34a7..370c585ede 100644 --- a/projects/clr/rocclr/device/pal/palresource.cpp +++ b/projects/clr/rocclr/device/pal/palresource.cpp @@ -1544,7 +1544,7 @@ bool Resource::partialMemCopyTo(VirtualGPU& gpu, const amd::Coord3D& srcOrigin, gpu.iCmd()->CmdCopyMemoryToImage(*iMem(), *dstResource.image_, imgLayout, 1, ©Region); } else if (!desc().buffer_ && dstResource.desc().buffer_) { Pal::MemoryImageCopyRegion copyRegion = {}; - uint32_t arraySliceIdx = img2Darray ? dstOrigin[2] : img1Darray ? dstOrigin[1] : 0; + uint32_t arraySliceIdx = img2Darray ? srcOrigin[2] : img1Darray ? srcOrigin[1] : 0; Pal::SubresId ImgSubresId = {0, desc().baseLevel_, arraySliceIdx}; copyRegion.imageSubres = ImgSubresId; copyRegion.imageOffset.x = srcOrigin[0];