From 56862b1c3544aed5cb6a639340bf8a8904af9d2c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 21 Nov 2017 21:19:06 +0530 Subject: [PATCH] Fixed review comments --- include/hip/hip_runtime_api.h | 6 +- src/hip_memory.cpp | 168 +++++++++++++++++----------------- src/hip_module.cpp | 23 ++--- src/hip_texture.cpp | 128 +++++++++++++------------- 4 files changed, 162 insertions(+), 163 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index a605406753..f764599164 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -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 }; diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 2b291b061b..75c054f5e1 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -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(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(acc.get_hsa_agent()); - size_t allocGranularity = 0; - hsa_amd_memory_pool_t *allocRegion = static_cast(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(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; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index af73cde3cf..f4e4b01ac2 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -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(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(it->second); + ret = hipSuccess; + } } + return ihipLogStatus(ret); } diff --git a/src/hip_texture.cpp b/src/hip_texture.cpp index b91d1dd494..7bf540ecc1 100644 --- a/src/hip_texture.cpp +++ b/src/hip_texture.cpp @@ -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;