AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Take into an account the number of channels...

Bläddra i källkod 
when querying the element size of an array.

Change-Id: Id57d3374b14d80a59230ec8286704f2fbabb0fae


[ROCm/clr commit: c9084d0ad2]
This commit is contained in:
Vladislav Sytchenko
2020-03-27 13:47:44 -04:00
committad av Vladislav Sytchenko
förälder a54b720eba
incheckning 4ca75d8ee2
3 ändrade filer med 17 tillägg och 17 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/vdi/hip_conversions.hpp
+10 -10
Visa fil
@@ -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;
projects/clr/hipamd/vdi/hip_memory.cpp
+5 -5
Visa fil
@@ -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));
}
projects/clr/hipamd/vdi/hip_texture.cpp
+2 -2
Visa fil
@@ -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;
}