diff --git a/projects/clr/hipamd/vdi/hip_conversions.hpp b/projects/clr/hipamd/vdi/hip_conversions.hpp index 0515c5c56b..d79a39c321 100644 --- a/projects/clr/hipamd/vdi/hip_conversions.hpp +++ b/projects/clr/hipamd/vdi/hip_conversions.hpp @@ -169,19 +169,19 @@ cl_mem_object_type getCLMemObjectType(const hipResourceType hipResType) { } inline -size_t getElementSize(const hipArray_Format arrayFormat) { - switch (arrayFormat) { +size_t getElementSize(const hipArray_const_t array) { + switch (array->Format) { case HIP_AD_FORMAT_UNSIGNED_INT8: case HIP_AD_FORMAT_SIGNED_INT8: - return 1; + return 1 * array->NumChannels; case HIP_AD_FORMAT_UNSIGNED_INT16: case HIP_AD_FORMAT_SIGNED_INT16: case HIP_AD_FORMAT_HALF: - return 2; + return 2 * array->NumChannels; case HIP_AD_FORMAT_UNSIGNED_INT32: case HIP_AD_FORMAT_SIGNED_INT32: case HIP_AD_FORMAT_FLOAT: - return 4; + return 4 * array->NumChannels; } ShouldNotReachHere(); @@ -617,7 +617,7 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.srcMemoryType = hipMemoryTypeArray; descDrv.srcArray = desc.srcArray; // When reffering to array memory, hipPos::x is in elements. - descDrv.srcXInBytes *= getElementSize(desc.srcArray->Format); + descDrv.srcXInBytes *= getElementSize(desc.srcArray); } if (desc.srcPtr.ptr != nullptr) { @@ -632,7 +632,7 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.dstMemoryType = hipMemoryTypeArray; descDrv.dstArray = desc.dstArray; // When reffering to array memory, hipPos::x is in elements. - descDrv.dstXInBytes *= getElementSize(desc.dstArray->Format); + descDrv.dstXInBytes *= getElementSize(desc.dstArray); } if (desc.dstPtr.ptr != nullptr) { @@ -645,11 +645,11 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { // If a HIP array is participating in the copy, the extent is defined in terms of that array's elements. if ((desc.srcArray != nullptr) && (desc.dstArray == nullptr)) { - descDrv.WidthInBytes *= getElementSize(desc.srcArray->Format); + descDrv.WidthInBytes *= getElementSize(desc.srcArray); } else if ((desc.srcArray == nullptr) && (desc.dstArray != nullptr)) { - descDrv.WidthInBytes *= getElementSize(desc.dstArray->Format); + descDrv.WidthInBytes *= getElementSize(desc.dstArray); } else if ((desc.srcArray != nullptr) && (desc.dstArray != nullptr)) { - descDrv.WidthInBytes *= getElementSize(desc.dstArray->Format); + descDrv.WidthInBytes *= getElementSize(desc.dstArray); } return descDrv; diff --git a/projects/clr/hipamd/vdi/hip_memory.cpp b/projects/clr/hipamd/vdi/hip_memory.cpp index 5438064e9b..e391bc28b3 100644 --- a/projects/clr/hipamd/vdi/hip_memory.cpp +++ b/projects/clr/hipamd/vdi/hip_memory.cpp @@ -1534,7 +1534,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; const size_t witdthInBytes = count / arrayHeight; - const size_t height = (count / dst->width) / (hip::getElementSize(dst->Format) * dst->NumChannels); + const size_t height = (count / dst->width) / hip::getElementSize(dst); HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, witdthInBytes, height, kind, nullptr)); } @@ -1574,7 +1574,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc const size_t arrayHeight = (src->height != 0) ? src->height : 1; const size_t witdthInBytes = count / arrayHeight; - const size_t height = (count / src->width) / (hip::getElementSize(src->Format) * src->NumChannels); + const size_t height = (count / src->width) / hip::getElementSize(src); HIP_RETURN(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, nullptr)); } @@ -1609,7 +1609,7 @@ hipError_t ihipMemcpy3D(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->Format) != hip::getElementSize(p->dstArray->Format))) { + (hip::getElementSize(p->dstArray) != hip::getElementSize(p->dstArray))) { return hipErrorInvalidValue; } @@ -2062,7 +2062,7 @@ hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffs const size_t arrayHeight = (src->height != 0) ? src->height : 1; const size_t widthInBytes = count / arrayHeight; - const size_t height = (count / src->width) / (hip::getElementSize(src->Format) * src->NumChannels); + const size_t height = (count / src->width) / hip::getElementSize(src); HIP_RETURN(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffsetSrc, widthInBytes, height, kind, stream, true)); } @@ -2083,7 +2083,7 @@ hipError_t hipMemcpyToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; const size_t widthInBytes = count / arrayHeight; - const size_t height = (count / dst->width) / (hip::getElementSize(dst->Format) * dst->NumChannels); + const size_t height = (count / dst->width) / hip::getElementSize(dst); HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, widthInBytes, height, kind, stream, true)); } diff --git a/projects/clr/hipamd/vdi/hip_texture.cpp b/projects/clr/hipamd/vdi/hip_texture.cpp index fcea0bbb47..0f8e818948 100644 --- a/projects/clr/hipamd/vdi/hip_texture.cpp +++ b/projects/clr/hipamd/vdi/hip_texture.cpp @@ -207,8 +207,8 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, hipTextureReadMode readMode = pTexDesc->readMode; // 32-bit integer format will not be promoted, regardless of whether or not // this hipTextureDesc::readMode is set hipReadModeNormalizedFloat is specified. - if ((hip::getElementSize(pResDesc->res.array.array->Format) == 4) && - (pResDesc->res.array.array->Format != HIP_AD_FORMAT_FLOAT)) { + if ((pResDesc->res.array.array->Format == HIP_AD_FORMAT_SIGNED_INT32) || + (pResDesc->res.array.array->Format == HIP_AD_FORMAT_UNSIGNED_INT32)) { readMode = hipReadModeElementType; }