AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-364236 - Fix layered Image issue

Browse Source 
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: 5a0085e516]
This commit is contained in:
taosang2
2023-09-28 09:46:19 -04:00
committed by Tao Sang
parent 648bfdc5cd
commit 28524520e6
5 changed files with 73 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_conversions.hpp
+34
View File
@@ -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<const cl_mem>(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;
}
projects/clr/hipamd/src/hip_graph_helper.hpp
+2 -2
View File
@@ -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);
projects/clr/hipamd/src/hip_graph_internal.cpp
+4 -5
View File
@@ -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;
}
projects/clr/hipamd/src/hip_memory.cpp
+32 -20
View File
@@ -1543,13 +1543,15 @@ hipError_t ihipMemcpyAtoDValidate(hipArray_t srcArray, void* dstDevice, amd::Coo
static_cast<size_t*>(srcOrigin)[0] /= elementSize;
static_cast<size_t*>(copyRegion)[0] /= elementSize;
amd::Coord3D copyRegionReal = copyRegion;
if (hip::isLayered1D(srcArray)) copyRegionReal.c[1] = 1;
if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegionReal),
srcImage->getRowPitch(), srcImage->getSlicePitch())) {
return hipErrorInvalidValue;
}
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegionReal),
dstRowPitch, dstSlicePitch)) {
return hipErrorInvalidValue;
}
@@ -1616,14 +1618,16 @@ hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray_t dstArray, amd::Coo
static_cast<size_t*>(dstOrigin)[0] /= elementSize;
static_cast<size_t*>(copyRegion)[0] /= elementSize;
amd::Coord3D copyRegionReal = copyRegion;
if (hip::isLayered1D(dstArray)) copyRegionReal.c[1] = 1;
if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegionReal),
srcRowPitch, srcSlicePitch)) {
return hipErrorInvalidValue;
}
srcRect.start_ += srcOffset;
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(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<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
srcRowPitch, srcSlicePitch)) {
return hipErrorInvalidValue;
}
start = ihipGetbufferStart(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(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<const char*>(srcHost) + srcRect.start_, srcRowPitch, srcSlicePitch,
copyRegion, static_cast<const char*>(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<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
dstRowPitch, dstSlicePitch)) {
return hipErrorInvalidValue;
}
start = ihipGetbufferStart(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(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<char*>(dstHost) + dstRect.start_, dstRowPitch, dstSlicePitch,
copyRegion, static_cast<char*>(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;
}
projects/clr/rocclr/device/pal/palresource.cpp
+1 -1
View File
@@ -1544,7 +1544,7 @@ bool Resource::partialMemCopyTo(VirtualGPU& gpu, const amd::Coord3D& srcOrigin,
gpu.iCmd()->CmdCopyMemoryToImage(*iMem(), *dstResource.image_, imgLayout, 1, &copyRegion);
} 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];