Skip to content

Commit 4098725

Browse files
przemektmalonaarongreig
przemektmalon
authored and
aarongreig
committed
Fix copy docs and implementation. (#19093)
This patch fixes the implementation of bindless image copies. Previously, source and destination pitch values were not being set correctly. This patch also updates the wording around the requirements for `ext_oneapi_copy`. A missing requirement was added to the specification. Namely that the `CopyExtent` parameter in the `ext_oneapi_copy` functions that take it, must not have `0` values in any of the three dimensions, they must be greater than or equal to `1`. The requirements for `ext_oneapi_copy` have also been re-written to prescribe what the functions expect, instead of providing a list of cases in which the function may fail. This should hopefully make it clearer and more prescriptive, rather than saying the copy function may fail if some condition is not met, we now say that the functions require that certain conditions be met. The coverage for sub-region copy testing has also been extended to prevent future regressions.
1 parent 2ee5e08 commit 4098725

3 files changed

Lines changed: 14 additions & 16 deletions

File tree

‎source/adapters/cuda/image.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -704,7 +704,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
704704
cpy_desc.dstY = pCopyRegion->dstOffset.y;
705705
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
706706
cpy_desc.Height = pCopyRegion->copyExtent.height;
707-
cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes;
707+
cpy_desc.srcPitch = pSrcImageDesc->rowPitch;
708708
if (pDstImageDesc->rowPitch == 0) {
709709
cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_ARRAY;
710710
cpy_desc.dstArray = (CUarray)pDst;
@@ -725,7 +725,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
725725
cpy_desc.dstZ = pCopyRegion->dstOffset.z;
726726
cpy_desc.srcMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST;
727727
cpy_desc.srcHost = pSrc;
728-
cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes;
728+
cpy_desc.srcPitch = pSrcImageDesc->rowPitch;
729729
cpy_desc.srcHeight = pSrcImageDesc->height;
730730
cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_ARRAY;
731731
cpy_desc.dstArray = (CUarray)pDst;
@@ -745,7 +745,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
745745
cpy_desc.dstZ = pCopyRegion->dstOffset.z;
746746
cpy_desc.srcMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST;
747747
cpy_desc.srcHost = pSrc;
748-
cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes;
748+
cpy_desc.srcPitch = pSrcImageDesc->rowPitch;
749749
cpy_desc.srcHeight = std::max(uint64_t{1}, pSrcImageDesc->height);
750750
cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_ARRAY;
751751
cpy_desc.dstArray = (CUarray)pDst;
@@ -793,7 +793,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
793793
cpy_desc.dstY = pCopyRegion->dstOffset.y;
794794
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
795795
cpy_desc.Height = pCopyRegion->copyExtent.height;
796-
cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes;
796+
cpy_desc.dstPitch = pDstImageDesc->rowPitch;
797797
cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST;
798798
cpy_desc.dstHost = pDst;
799799
if (pSrcImageDesc->rowPitch == 0) {
@@ -818,7 +818,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
818818
cpy_desc.srcArray = as_CUArray(pSrc);
819819
cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST;
820820
cpy_desc.dstHost = pDst;
821-
cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes;
821+
cpy_desc.dstPitch = pDstImageDesc->rowPitch;
822822
cpy_desc.dstHeight = pDstImageDesc->height;
823823
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
824824
cpy_desc.Height = pCopyRegion->copyExtent.height;
@@ -838,7 +838,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
838838
cpy_desc.srcArray = as_CUArray(pSrc);
839839
cpy_desc.dstMemoryType = CUmemorytype_enum::CU_MEMORYTYPE_HOST;
840840
cpy_desc.dstHost = pDst;
841-
cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes;
841+
cpy_desc.dstPitch = pDstImageDesc->rowPitch;
842842
cpy_desc.dstHeight = std::max(uint64_t{1}, pDstImageDesc->height);
843843
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
844844
cpy_desc.Height = std::max(uint64_t{1}, pCopyRegion->copyExtent.height);

‎source/adapters/hip/image.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -704,7 +704,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
704704
cpy_desc.srcY = pCopyRegion->srcOffset.y;
705705
cpy_desc.dstXInBytes = pCopyRegion->dstOffset.x * PixelSizeBytes;
706706
cpy_desc.dstY = pCopyRegion->dstOffset.y;
707-
cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes;
707+
cpy_desc.srcPitch = pSrcImageDesc->rowPitch;
708708
if (pDstImageDesc->rowPitch == 0) {
709709
cpy_desc.dstMemoryType = hipMemoryTypeArray;
710710
cpy_desc.dstArray = static_cast<hipArray_t>(pDst);
@@ -727,7 +727,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
727727
cpy_desc.dstZ = pCopyRegion->dstOffset.z;
728728
cpy_desc.srcMemoryType = hipMemoryTypeHost;
729729
cpy_desc.srcHost = pSrc;
730-
cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes;
730+
cpy_desc.srcPitch = pSrcImageDesc->rowPitch;
731731
cpy_desc.srcHeight = pSrcImageDesc->height;
732732
cpy_desc.dstMemoryType = hipMemoryTypeArray;
733733
cpy_desc.dstArray = static_cast<hipArray_t>(pDst);
@@ -749,7 +749,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
749749
cpy_desc.dstZ = pCopyRegion->dstOffset.z;
750750
cpy_desc.srcMemoryType = hipMemoryTypeHost;
751751
cpy_desc.srcHost = pSrc;
752-
cpy_desc.srcPitch = pSrcImageDesc->width * PixelSizeBytes;
752+
cpy_desc.srcPitch = pSrcImageDesc->rowPitch;
753753
cpy_desc.srcHeight = std::max(MinCopyHeight, pSrcImageDesc->height);
754754
cpy_desc.dstMemoryType = hipMemoryTypeArray;
755755
cpy_desc.dstArray = static_cast<hipArray_t>(pDst);
@@ -824,7 +824,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
824824
}
825825
cpy_desc.dstMemoryType = hipMemoryTypeHost;
826826
cpy_desc.dstHost = pDst;
827-
cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes;
827+
cpy_desc.dstPitch = pDstImageDesc->rowPitch;
828828
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
829829
cpy_desc.Height = pCopyRegion->copyExtent.height;
830830
UR_CHECK_ERROR(hipMemcpyParam2DAsync(&cpy_desc, Stream));
@@ -840,7 +840,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
840840
cpy_desc.srcArray = static_cast<hipArray_t>(const_cast<void *>(pSrc));
841841
cpy_desc.dstMemoryType = hipMemoryTypeHost;
842842
cpy_desc.dstHost = pDst;
843-
cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes;
843+
cpy_desc.dstPitch = pDstImageDesc->rowPitch;
844844
cpy_desc.dstHeight = pDstImageDesc->height;
845845
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
846846
cpy_desc.Height = pCopyRegion->copyExtent.height;
@@ -863,7 +863,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageCopyExp(
863863
cpy_desc.srcArray = static_cast<hipArray_t>(const_cast<void *>(pSrc));
864864
cpy_desc.dstMemoryType = hipMemoryTypeHost;
865865
cpy_desc.dstHost = pDst;
866-
cpy_desc.dstPitch = pDstImageDesc->width * PixelSizeBytes;
866+
cpy_desc.dstPitch = pDstImageDesc->rowPitch;
867867
cpy_desc.dstHeight = std::max(MinCopyHeight, pDstImageDesc->height);
868868
cpy_desc.WidthInBytes = PixelSizeBytes * pCopyRegion->copyExtent.width;
869869
cpy_desc.Height =

‎source/adapters/level_zero/image_common.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -779,8 +779,7 @@ ur_result_t bindlessImagesHandleCopyFlags(
779779

780780
switch (imageCopyFlags) {
781781
case UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE: {
782-
uint32_t SrcRowPitch =
783-
pSrcImageDesc->width * getPixelSizeBytes(pSrcImageFormat);
782+
uint32_t SrcRowPitch = pSrcImageDesc->rowPitch;
784783
uint32_t SrcSlicePitch = SrcRowPitch * pSrcImageDesc->height;
785784
if (pDstImageDesc->rowPitch == 0) {
786785
// Copy to Non-USM memory
@@ -824,8 +823,7 @@ ur_result_t bindlessImagesHandleCopyFlags(
824823
return UR_RESULT_SUCCESS;
825824
};
826825
case UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST: {
827-
uint32_t DstRowPitch =
828-
pDstImageDesc->width * getPixelSizeBytes(pDstImageFormat);
826+
uint32_t DstRowPitch = pDstImageDesc->rowPitch;
829827
uint32_t DstSlicePitch = DstRowPitch * pDstImageDesc->height;
830828
if (pSrcImageDesc->rowPitch == 0) {
831829
// Copy from Non-USM memory to host

0 commit comments

Comments
 (0)