AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

Fix texture 3D & 2D layered with N components (#1746)

Przeglądaj źródła 
SWDEV-151670: Issue with 3D texture with 4 components
SWDEV-151671: Issue with 2D layered texture with 4 components

Fixed memcpy when memory is allocated with driver API's.

Github issues: #1755
Fixed 3D default case when array type is not set during memory allocation.
This commit is contained in:
ansurya
2019-12-30 12:41:42 +05:30
zatwierdzone przez Maneesh Gupta
rodzic 006e6844c1
commit 651c7a8e27
1 zmienionych plików z 51 dodań i 10 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
hipamd/src/hip_memory.cpp
+51 -10
Wyświetl plik
@@ -835,7 +835,7 @@ hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width
HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo);
if(hip_status == hipSuccess)
*pitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
*pitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
return hip_status;
}
@@ -925,6 +925,24 @@ hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height,
return hipSuccess;
}
hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,size_t width, size_t height, size_t depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType, hsa_ext_image_data_info_t &imageInfo,size_t array_size __dparm(0))
{
hsa_ext_image_descriptor_t imageDescriptor = {.geometry = geometry, .width = width, .height= height, .depth = depth,
.array_size = array_size, .format.channel_order = channelOrder, .format.channel_type = channelType};
// Get the current device agent.
hc::accelerator acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if (!agent)
return hipErrorInvalidResourceHandle;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
hsa_status_t status =
hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
if(HSA_STATUS_SUCCESS != status){
return hipErrorRuntimeOther;
}
return hipSuccess;
}
hipError_t ihipArrayToImageFormat(hipArray_Format format,hsa_ext_image_channel_type_t &channelType) {
switch (format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
@@ -1070,6 +1088,7 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
case hipArrayDefault:
case hipArrayCubemap:
default:
array[0]->type = hipArrayCubemap;
hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,pAllocateArray->Width,pAllocateArray->Height,
pAllocateArray->Depth,channelOrder,channelType,ptr,imageInfo);
array[0]->textureType = hipTextureType3D;
@@ -1113,6 +1132,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
case hipArrayDefault:
case hipArrayCubemap:
default:
array[0]->type = hipArrayCubemap;
hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,extent.width,extent.height,extent.depth,channelOrder,channelType,ptr,imageInfo);
array[0]->textureType = hipTextureType3D;
break;
@@ -1546,25 +1566,46 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo
ySize = p->srcPtr.ysize;
desc = p->dstArray->desc;
dstPtr = p->dstArray->data;
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
} else {
depth = p->Depth;
height = p->Height;
widthInBytes = p->WidthInBytes;
width = p->dstArray->width;
desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
hsa_ext_image_channel_order_t channelOrder;
switch(p->dstArray->NumChannels) {
case 2:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
break;
case 3:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
break;
case 4:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
break;
case 1:
default:
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
break;
}
hsa_ext_image_channel_type_t channelType;
e = ihipArrayToImageFormat(p->dstArray->Format,channelType);
srcPitch = p->srcPitch;
srcPtr = (void*)p->srcHost;
ySize = p->srcHeight;
dstPtr = p->dstArray->data;
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, channelOrder, channelType, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, channelOrder, channelType, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
}
hsa_ext_image_data_info_t imageInfo;
if(hipTextureType2DLayered == p->dstArray->textureType)
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_2DA, width, height, 0, desc, imageInfo, depth);
else
GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, desc, imageInfo);
dstPitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
} else {
// Non array destination
depth = p->extent.depth;