AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fixed review comments

Browse Source 
[ROCm/clr commit: 38029f2849]
This commit is contained in:
Rahul Garg
2017-11-21 21:19:06 +05:30
parent dc2c56c43d
commit 212afdf56b
4 changed files with 162 additions and 163 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hip_runtime_api.h
+3 -3
View File
@@ -116,10 +116,10 @@ typedef struct hipDeviceProp_t {
* Memory type (for pointer attributes)
*/
enum hipMemoryType {
hipMemoryTypeHost, ///< Memory is physically located on host
hipMemoryTypeHost, ///< Memory is physically located on host
hipMemoryTypeDevice, ///< Memory is physically located on device. (see deviceId for specific device)
hipMemoryTypeArray,
hipMemoryTypeUnified
hipMemoryTypeArray, ///< Array memory, physically located on device. (see deviceId for specific device)
hipMemoryTypeUnified ///< Not used currently
};
projects/clr/hipamd/src/hip_memory.cpp
+83 -85
View File
@@ -407,8 +407,8 @@ hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannel
extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
enum hipTextureReadMode readMode,
hsa_ext_image_channel_order_t& channelOrder,
hsa_ext_image_channel_type_t& channelType);
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType);
hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray )
{
@@ -417,106 +417,104 @@ hipError_t hipArrayCreate ( hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAlloc
hipError_t hip_status = hipSuccess;
if(pAllocateArray->width >0) {
auto ctx = ihipGetTlsDefaultCtx();
*array = (hipArray*)malloc(sizeof(hipArray));
*array = (hipArray*)malloc(sizeof(hipArray));
array[0]->drvDesc = *pAllocateArray;
array[0]->width = pAllocateArray->width;
array[0]->height = pAllocateArray->height;
array[0]->isDrv = true;
void ** ptr = &array[0]->data;
if (ctx) {
const unsigned am_flags = 0;
size_t size = pAllocateArray->width;
if(pAllocateArray->height > 0) {
size = size * pAllocateArray->height;
}
const unsigned am_flags = 0;
size_t size = pAllocateArray->width;
if(pAllocateArray->height > 0) {
size = size * pAllocateArray->height;
}
hsa_ext_image_channel_type_t channelType;
size_t allocSize = 0;
switch(pAllocateArray->format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
allocSize = size * sizeof(unsigned char);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
allocSize = size * sizeof(unsigned short);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
allocSize = size * sizeof(unsigned int);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case HIP_AD_FORMAT_SIGNED_INT8:
allocSize = size * sizeof(char);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
allocSize = size * sizeof(short);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
allocSize = size * sizeof(int);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case HIP_AD_FORMAT_HALF:
allocSize = size * sizeof(short);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case HIP_AD_FORMAT_FLOAT:
allocSize = size * sizeof(float);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
default:
hip_status = hipErrorUnknown;
break;
}
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
size_t allocSize = 0;
switch(pAllocateArray->format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
allocSize = size * sizeof(uint8_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
allocSize = size * sizeof(uint16_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
allocSize = size * sizeof(uint32_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case HIP_AD_FORMAT_SIGNED_INT8:
allocSize = size * sizeof(int8_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
allocSize = size * sizeof(int16_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
allocSize = size * sizeof(int32_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case HIP_AD_FORMAT_HALF:
allocSize = size * sizeof(int16_t);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case HIP_AD_FORMAT_FLOAT:
allocSize = size * sizeof(float);
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
default:
hip_status = hipErrorUnknown;
break;
}
hc::accelerator acc = ctx->getDevice()->_acc;
hsa_agent_t* agent =static_cast<hsa_agent_t*>(acc.get_hsa_agent());
size_t allocGranularity = 0;
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);
size_t allocGranularity = 0;
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;
hsa_ext_image_descriptor_t imageDescriptor;
imageDescriptor.width = pAllocateArray->width;
imageDescriptor.height = pAllocateArray->height;
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.width = pAllocateArray->width;
imageDescriptor.height = pAllocateArray->height;
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_order_t channelOrder;
//getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
if (pAllocateArray->numChannels == 4) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (pAllocateArray->numChannels == 2) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (pAllocateArray->numChannels == 1) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
if (pAllocateArray->numChannels == 4) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (pAllocateArray->numChannels == 2) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (pAllocateArray->numChannels == 1) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
hsa_ext_image_data_info_t imageInfo;
hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
hsa_ext_image_data_info_t imageInfo;
hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
*ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment);
if (size && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorInvalidValue;
*ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false/*shareWithAll*/, am_flags, 0, alignment);
if (size && (*ptr == NULL)) {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorMemoryAllocation;
}
} else {
hip_status = hipErrorInvalidValue;
}
return ihipLogStatus(hip_status);
}
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
size_t width, size_t height, unsigned int flags)
{
@@ -588,7 +586,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
}
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -685,7 +683,7 @@ hipError_t hipMalloc3DArray(hipArray_t *array,
}
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
getChannelOrderAndType(*desc, hipReadModeElementType, channelOrder, channelType);
getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType);
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
projects/clr/hipamd/src/hip_module.cpp
+12 -11
View File
@@ -841,17 +841,18 @@ hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image, unsigned
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name)
{
HIP_INIT_API(texRef, hmod, name);
hipError_t ret = hipSuccess;
hipError_t ret = hipErrorNotFound;
if(texRef == NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
if(name == NULL || hmod == NULL){
return ihipLogStatus(hipErrorNotInitialized);
}
else{
const auto it = hmod->coGlobals.find(name);
if (it == hmod->coGlobals.end()) return ihipLogStatus(hipErrorInvalidValue);
*texRef = reinterpret_cast<textureReference*>(it->second);
return ihipLogStatus(ret);
ret = hipErrorInvalidValue;
} else {
if(name == NULL || hmod == NULL){
ret = hipErrorNotInitialized;
} else{
const auto it = hmod->coGlobals.find(name);
if (it == hmod->coGlobals.end()) return ihipLogStatus(hipErrorInvalidValue);
*texRef = reinterpret_cast<textureReference*>(it->second);
ret = hipSuccess;
}
}
return ihipLogStatus(ret);
}
projects/clr/hipamd/src/hip_texture.cpp
+64 -64
View File
@@ -34,59 +34,59 @@ void saveTextureInfo(const hipTexture* pTexture,
void getDrvChannelOrderAndType(const enum hipArray_Format Format,
unsigned int NumChannels,
hsa_ext_image_channel_order_t& channelOrder,
hsa_ext_image_channel_type_t& channelType)
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType)
{
switch(Format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case HIP_AD_FORMAT_SIGNED_INT8:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case HIP_AD_FORMAT_HALF:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case HIP_AD_FORMAT_FLOAT:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
default:
break;
}
switch(Format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case HIP_AD_FORMAT_SIGNED_INT8:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
case HIP_AD_FORMAT_SIGNED_INT16:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case HIP_AD_FORMAT_SIGNED_INT32:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case HIP_AD_FORMAT_HALF:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case HIP_AD_FORMAT_FLOAT:
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
default:
break;
}
if (NumChannels == 4) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (NumChannels == 2) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (NumChannels == 1) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
if (NumChannels == 4) {
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (NumChannels == 2) {
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (NumChannels == 1) {
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
}
}
void getChannelOrderAndType(const hipChannelFormatDesc& desc,
enum hipTextureReadMode readMode,
hsa_ext_image_channel_order_t& channelOrder,
hsa_ext_image_channel_type_t& channelType)
hsa_ext_image_channel_order_t* channelOrder,
hsa_ext_image_channel_type_t* channelType)
{
if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w != 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
} else if (desc.x != 0 && desc.y != 0 && desc.z != 0 && desc.w == 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
} else if (desc.x != 0 && desc.y != 0 && desc.z == 0 && desc.w == 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
} else if (desc.x != 0 && desc.y == 0 && desc.z == 0 && desc.w == 0) {
channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
*channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
} else {
}
@@ -94,49 +94,49 @@ void getChannelOrderAndType(const hipChannelFormatDesc& desc,
case hipChannelFormatKindUnsigned:
switch(desc.x) {
case 32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
break;
case 16:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT16 :
HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16;
break;
case 8:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_UNORM_INT8 :
HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8;
break;
default:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32;
}
break;
case hipChannelFormatKindSigned:
switch(desc.x) {
case 32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
break;
case 16:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT16 :
HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT16;
break;
case 8:
channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 :
*channelType = readMode == hipReadModeNormalizedFloat ? HSA_EXT_IMAGE_CHANNEL_TYPE_SNORM_INT8 :
HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8;
break;
default:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32;
}
break;
case hipChannelFormatKindFloat:
switch(desc.x) {
case 32:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
break;
case 16:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_HALF_FLOAT;
break;
case 8:
break;
default:
channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
*channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT;
}
break;
case hipChannelFormatKindNone:
@@ -255,7 +255,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.array_size = 0;
break;
}
getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.array.array->desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
case hipResourceTypeMipmappedArray:
devPtr = pResDesc->res.mipmap.mipmap->data;
@@ -264,7 +264,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.depth = pResDesc->res.mipmap.mipmap->depth;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
getChannelOrderAndType(pResDesc->res.mipmap.mipmap->desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.mipmap.mipmap->desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
case hipResourceTypeLinear:
devPtr = pResDesc->res.linear.devPtr;
@@ -273,7 +273,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_1D; // ? HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR
getChannelOrderAndType(pResDesc->res.linear.desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.linear.desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
case hipResourceTypePitch2D:
devPtr = pResDesc->res.pitch2D.devPtr;
@@ -282,7 +282,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
imageDescriptor.depth = 0;
imageDescriptor.array_size = 0;
imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D;
getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, channelOrder, channelType);
getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, &channelOrder, &channelType);
break;
default:
break;
@@ -416,9 +416,9 @@ hipError_t ihipBindTextureImpl(int dim,
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
if(NULL == desc) {
getDrvChannelOrderAndType(tex->format, tex->numChannels, channelOrder, channelType);
getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(*desc, readMode, channelOrder, channelType);
getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType);
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -493,9 +493,9 @@ hipError_t ihipBindTexture2DImpl(int dim,
hsa_ext_image_channel_type_t channelType;
if(NULL == desc) {
getDrvChannelOrderAndType(tex->format, tex->numChannels, channelOrder, channelType);
getDrvChannelOrderAndType(tex->format, tex->numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(*desc, readMode, channelOrder, channelType);
getChannelOrderAndType(*desc, readMode, &channelOrder, &channelType);
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;
@@ -593,9 +593,9 @@ hipError_t ihipBindTextureToArrayImpl(int dim,
hsa_ext_image_channel_order_t channelOrder;
hsa_ext_image_channel_type_t channelType;
if(array->isDrv) {
getDrvChannelOrderAndType(array->drvDesc.format, array->drvDesc.numChannels, channelOrder, channelType);
getDrvChannelOrderAndType(array->drvDesc.format, array->drvDesc.numChannels, &channelOrder, &channelType);
} else {
getChannelOrderAndType(desc, readMode, channelOrder, channelType);
getChannelOrderAndType(desc, readMode, &channelOrder, &channelType);
}
imageDescriptor.format.channel_order = channelOrder;
imageDescriptor.format.channel_type = channelType;