AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Fix texture tests, always pass row pitch to HSA API for linear layout images (#1785)

Просмотр исходного кода 
=> New ROCr calculates pitch as per HSA specification and addrlib is used to check whether HW can support that configuration. Hence few texture tests are failing with HSA_EXT_STATUS_ERROR_IMAGE_PITCH_UNSUPPORTED.

=> Determine pitch for linear images and always pass rowpitch to HSA API's.

[ROCm/hip commit: df20c17f12]
Этот коммит содержится в:
ansurya
2020-01-16 08:54:30 +05:30
коммит произвёл Maneesh Gupta
родитель dabbde21c6
Коммит ac41b1d0e0
3 изменённых файлов: 86 добавлений и 55 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip/src/hip_hcc_internal.h
+51
Просмотреть файл
@@ -91,6 +91,57 @@ extern int HIP_DUMP_CODE_OBJECT;
// TODO - remove when this is standard behavior.
extern int HCC_OPT_FLUSH;
#define IMAGE_PITCH_ALIGNMENT 256
template <typename T> inline T alignDown(T value, size_t alignment) {
return (T)(value & ~(alignment - 1));
}
template <typename T> inline T* alignDown(T* value, size_t alignment) {
return (T*)alignDown((intptr_t)value, alignment);
}
template <typename T> inline T alignUp(T value, size_t alignment) {
return alignDown((T)(value + alignment - 1), alignment);
}
template <typename T> inline T* alignUp(T* value, size_t alignment) {
return (T*)alignDown((intptr_t)(value + alignment - 1), alignment);
}
size_t getNumChannels(hsa_ext_image_channel_order_t channelOrder) {
switch (channelOrder) {
case HSA_EXT_IMAGE_CHANNEL_ORDER_RG:
return 2;
case HSA_EXT_IMAGE_CHANNEL_ORDER_RGB:
return 3;
case HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA:
return 4;
case HSA_EXT_IMAGE_CHANNEL_ORDER_R:
default:
return 1;
}
}
size_t getElementSize(hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType) {
size_t bytesPerPixel = getNumChannels(channelOrder);
switch (channelType) {
case HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8:
case HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8:
break;
case HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32:
case HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32:
case HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT:
bytesPerPixel *= 4;
break;
default:
bytesPerPixel *= 2;
break;
}
return bytesPerPixel;
}
// Class to assign a short TID to each new thread, for HIP debugging purposes.
class TidInfo {
public:
projects/hip/src/hip_memory.cpp
+22 -53
Просмотреть файл
@@ -37,6 +37,7 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
namespace hip_internal {
namespace {
inline
const char* hsa_to_string(hsa_status_t err) noexcept
{
@@ -791,21 +792,13 @@ hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width,
hsa_amd_memory_pool_t* allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
hsa_ext_image_descriptor_t imageDescriptor;
imageDescriptor.geometry = geometry;
imageDescriptor.width = width;
imageDescriptor.height = height;
imageDescriptor.depth = depth;
imageDescriptor.array_size = array_size;
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size) ;
else
imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth) ;
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(imageInfo.size == 0 || HSA_STATUS_SUCCESS != status){
return hipErrorRuntimeOther;
}
imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
const unsigned am_flags = 0;
*ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx,
@@ -898,55 +891,31 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height, int depth, hipChannelFormatDesc desc, hsa_ext_image_data_info_t &imageInfo,int array_size __dparm(0))
{
hsa_ext_image_descriptor_t imageDescriptor;
imageDescriptor.geometry = geometry;
imageDescriptor.width = width;
imageDescriptor.height = height;
imageDescriptor.depth = depth;
imageDescriptor.array_size = array_size;
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(desc, hipReadModeElementType, &channelOrder, &channelType);
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
// Get the current device agent.
hc::accelerator acc;
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
if (!agent)
return hipErrorInvalidHandle;
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;
}
size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size);
else
imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth);
imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
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;
imageDescriptor.geometry = geometry;
imageDescriptor.width = width;
imageDescriptor.height = height;
imageDescriptor.depth = depth;
imageDescriptor.array_size = array_size;
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.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;
}
size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size);
else
imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth);
imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
return hipSuccess;
}
projects/hip/src/hip_texture.cpp
+13 -2
Просмотреть файл
@@ -295,6 +295,10 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
pTexDesc->normalizedCoords);
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
if(hipResourceTypePitch2D != pResDesc->resType)
pitch = getElementSize(channelOrder, channelType) * alignUp(imageDescriptor.width, IMAGE_PITCH_ALIGNMENT);
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) ||
@@ -430,9 +434,11 @@ hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode r
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(size, IMAGE_PITCH_ALIGNMENT);
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS !=
hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return hipErrorRuntimeOther;
@@ -503,6 +509,9 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
if( 0 == pitch)
pitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, devPtr, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) ||
@@ -606,9 +615,11 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureRea
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(imageDescriptor.width, IMAGE_PITCH_ALIGNMENT);
if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
*agent, &imageDescriptor, array->data, permission,
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &(pTexture->image)) ||
HSA_STATUS_SUCCESS !=
hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
return hipErrorRuntimeOther;