From e07926ce0f485b3cc4c95ad6492206c749dc922d Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Thu, 7 Nov 2019 13:17:46 +0530 Subject: [PATCH] Fixed texture 2D mapping for pitched arrays & 3D Texture read (#1415) Texture 2D image mapping for pitched arrays: github issue: Texture Object's Buffer seems to be Misaligned #886 JIRA ticket: SWDEV-199313 SWDEV-151670 : Fixed issue with 3D texture with 4 components SWDEV-151671 : Issue with 2D layered texture with 4 components --- include/hip/hcc_detail/hip_runtime_api.h | 2 +- src/hip_memory.cpp | 595 ++++++------------ src/hip_texture.cpp | 20 +- tests/src/test_common.h | 30 + tests/src/texture/hipBindTex2DPitch.cpp | 87 +++ .../texture/hipNormalizedFloatValueTex.cpp | 33 +- tests/src/texture/hipTexObjPitch.cpp | 106 ++++ tests/src/texture/simpleTexture2DLayered.cpp | 106 ++++ tests/src/texture/simpleTexture3D.cpp | 133 ++++ 9 files changed, 698 insertions(+), 414 deletions(-) create mode 100644 tests/src/texture/hipBindTex2DPitch.cpp create mode 100644 tests/src/texture/hipTexObjPitch.cpp create mode 100644 tests/src/texture/simpleTexture2DLayered.cpp create mode 100644 tests/src/texture/simpleTexture3D.cpp diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index d955c078dd..ba7f2f4ade 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -3315,7 +3315,7 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, - size_t width, size_t height, textureReference* tex); + size_t width, size_t height, textureReference* tex, size_t pitch); template hipError_t hipBindTexture2D(size_t* offset, struct texture& tex, diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 009c2db505..3f82246cf9 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -429,58 +429,62 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { return hipHostMalloc(ptr, sizeBytes, flags); }; +hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width, int height, int depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType,void ** ptr, hsa_ext_image_data_info_t &imageInfo, int array_size __dparm(0)) { + auto ctx = ihipGetTlsDefaultCtx(); + if (ctx) { + hc::accelerator acc = ctx->getDevice()->_acc; + hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); + if (!agent) + return hipErrorInvalidResourceHandle; + 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; + 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; + + 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; + } + size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment; + const unsigned am_flags = 0; + *ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx, + false /*shareWithAll*/, am_flags, 0, alignment); + if (*ptr == NULL) { + return hipErrorMemoryAllocation; + } + return hipSuccess; + } + else { + return hipErrorMemoryAllocation; + } +} + // width in bytes hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) { hipError_t hip_status = hipSuccess; - if(ptr==NULL || pitch == NULL) - { - hip_status=hipErrorInvalidValue; - return hip_status; - } - // hardcoded 128 bytes - *pitch = ((((int)width - 1) / 128) + 1) * 128; - const size_t sizeBytes = (*pitch) * height * ((depth==0) ? 1 : depth); - - auto ctx = ihipGetTlsDefaultCtx(); - - if (ctx) { - 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); - - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = *pitch; - imageDescriptor.height = height; - imageDescriptor.depth = depth; - imageDescriptor.array_size = 0; - if (depth == 0) - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; - else - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.format.channel_order = HSA_EXT_IMAGE_CHANNEL_ORDER_R; - imageDescriptor.format.channel_type = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32; - - 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; - - const unsigned am_flags = 0; - *ptr = hip_internal::allocAndSharePtr("device_pitch", sizeBytes, ctx, - false /*shareWithAll*/, am_flags, 0, alignment); - - if (sizeBytes && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - } else { - hip_status = hipErrorMemoryAllocation; + if(ptr==NULL || pitch == NULL){ + return hipErrorInvalidValue; } + hsa_ext_image_data_info_t imageInfo; + if (depth == 0) + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,width,height,0,HSA_EXT_IMAGE_CHANNEL_ORDER_R, + HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo); + else + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,width,height,depth,HSA_EXT_IMAGE_CHANNEL_ORDER_R, + HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo); + + if(hip_status == hipSuccess) + *pitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth); return hip_status; } @@ -541,18 +545,74 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc, hsa_ext_image_channel_order_t* channelOrder, hsa_ext_image_channel_type_t* channelType); +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(acc.get_hsa_agent()); + if (!agent) + return hipErrorInvalidResourceHandle; + 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: + 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: + return hipErrorUnknown; + break; + } + return hipSuccess; +} + hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_SPECIAL_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; if (pAllocateArray->Width > 0) { - auto ctx = ihipGetTlsDefaultCtx(); *array = (hipArray*)malloc(sizeof(hipArray)); - HIP_ARRAY3D_DESCRIPTOR array3D; - array3D.Width = pAllocateArray->Width; - array3D.Height = pAllocateArray->Height; - array3D.Format = pAllocateArray->Format; - array3D.NumChannels = pAllocateArray->NumChannels; array[0]->width = pAllocateArray->Width; array[0]->height = pAllocateArray->Height; array[0]->Format = pAllocateArray->Format; @@ -560,100 +620,26 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat array[0]->isDrv = true; array[0]->textureType = hipTextureType2D; 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; - } - hsa_ext_image_channel_type_t channelType; - 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()); + hsa_ext_image_channel_type_t channelType; + hsa_ext_image_channel_order_t channelOrder; - 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); + hip_status = ihipArrayToImageFormat(pAllocateArray->Format,channelType); + if(hipSuccess != hip_status) + return ihipLogStatus(hip_status); - hsa_ext_image_descriptor_t imageDescriptor; - - imageDescriptor.width = pAllocateArray->Width; - imageDescriptor.height = pAllocateArray->Height; - imageDescriptor.depth = 0; - imageDescriptor.array_size = 0; - - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; - - hsa_ext_image_channel_order_t channelOrder; - - 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; - - *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; + 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; } + hsa_ext_image_data_info_t imageInfo; + return ihipLogStatus(allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,pAllocateArray->Width, + pAllocateArray->Height,0,channelOrder,channelType,ptr,imageInfo)); } else { - hip_status = hipErrorInvalidValue; + return ihipLogStatus(hipErrorInvalidValue); } - - return ihipLogStatus(hip_status); } hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, @@ -662,8 +648,6 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; if (width > 0) { - auto ctx = ihipGetTlsDefaultCtx(); - *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = flags; array[0]->width = width; @@ -674,67 +658,25 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si array[0]->textureType = hipTextureType2D; void** ptr = &array[0]->data; - if (ctx) { - const unsigned am_flags = 0; - size_t size = width; - if (height > 0) { - size = size * height; - } - - const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8); - - 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); - - hsa_ext_image_descriptor_t imageDescriptor; - - imageDescriptor.width = width; - imageDescriptor.height = height; - imageDescriptor.depth = 0; - imageDescriptor.array_size = 0; - switch (flags) { - case hipArrayLayered: - case hipArrayCubemap: - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - assert(0); - break; - case hipArrayDefault: - default: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; - break; - } - 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; - 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; + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); + hsa_ext_image_data_info_t imageInfo; + switch (flags) { + case hipArrayLayered: + case hipArrayCubemap: + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + default: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2D,width,height,0,channelOrder,channelType,ptr,imageInfo); + break; } } else { hip_status = hipErrorInvalidValue; } - return ihipLogStatus(hip_status); } @@ -742,8 +684,6 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll HIP_INIT_SPECIAL_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray); hipError_t hip_status = hipSuccess; - auto ctx = ihipGetTlsDefaultCtx(); - *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = pAllocateArray->Flags; array[0]->width = pAllocateArray->Width; @@ -752,111 +692,37 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll array[0]->Format = pAllocateArray->Format; array[0]->NumChannels = pAllocateArray->NumChannels; array[0]->isDrv = true; - array[0]->textureType = hipTextureType3D; void** ptr = &array[0]->data; - if (ctx) { - const unsigned am_flags = 0; - const size_t size = pAllocateArray->Width * pAllocateArray->Height * pAllocateArray->Depth; + hsa_ext_image_channel_order_t channelOrder; + 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; + } - size_t allocSize = 0; - hsa_ext_image_channel_type_t channelType; - 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); - - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = pAllocateArray->Width; - imageDescriptor.height = pAllocateArray->Height; - imageDescriptor.depth = 0; - imageDescriptor.array_size = 0; - switch (pAllocateArray->Flags) { - case hipArrayLayered: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA; - imageDescriptor.array_size = pAllocateArray->Depth; - break; - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - case hipArrayDefault: - assert(0); - break; - case hipArrayCubemap: - default: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.depth = pAllocateArray->Depth; - break; - } - 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; - - 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, am_flags, 0, - alignment); - - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - - } else { - hip_status = hipErrorMemoryAllocation; + hsa_ext_image_channel_type_t channelType; + hip_status = ihipArrayToImageFormat(pAllocateArray->Format,channelType); + hsa_ext_image_data_info_t imageInfo; + switch (pAllocateArray->Flags) { + case hipArrayLayered: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2DA,pAllocateArray->Width,pAllocateArray->Height,0, + channelOrder,channelType,ptr,imageInfo,pAllocateArray->Depth); + array[0]->textureType = hipTextureType2DLayered; + break; + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + case hipArrayCubemap: + default: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,pAllocateArray->Width,pAllocateArray->Height, + pAllocateArray->Depth,channelOrder,channelType,ptr,imageInfo); + array[0]->textureType = hipTextureType3D; + break; } return ihipLogStatus(hip_status); @@ -865,19 +731,13 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, struct hipExtent extent, unsigned int flags) { - - HIP_INIT_API(hipMalloc3DArray, array, desc, &extent, flags); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; - if(array==NULL ) - { - hip_status=hipErrorInvalidValue; - return ihipLogStatus(hip_status); + if(array==NULL ){ + return ihipLogStatus(hipErrorInvalidValue); } - auto ctx = ihipGetTlsDefaultCtx(); - *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->type = flags; array[0]->width = extent.width; @@ -885,68 +745,27 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* array[0]->depth = extent.depth; array[0]->desc = *desc; array[0]->isDrv = false; - array[0]->textureType = hipTextureType3D; void** ptr = &array[0]->data; - - if (ctx) { - const unsigned am_flags = 0; - const size_t size = extent.width * extent.height * extent.depth; - - const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8); - - 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); - - hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = extent.width; - imageDescriptor.height = extent.height; - imageDescriptor.depth = extent.depth; - imageDescriptor.array_size = 0; - switch (flags) { - case hipArrayLayered: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA; - imageDescriptor.array_size = extent.depth; - break; - case hipArraySurfaceLoadStore: - case hipArrayTextureGather: - case hipArrayDefault: - assert(0); - break; - case hipArrayCubemap: - default: - imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.depth = extent.depth; - break; - } - 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; - 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, am_flags, 0, - alignment); - - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } - - } else { - hip_status = hipErrorMemoryAllocation; + hsa_ext_image_channel_order_t channelOrder; + hsa_ext_image_channel_type_t channelType; + getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); + hsa_ext_image_data_info_t imageInfo; + switch (flags) { + case hipArrayLayered: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_2DA,extent.width,extent.height,0,channelOrder,channelType,ptr,imageInfo,extent.depth); + array[0]->textureType = hipTextureType2DLayered; + break; + case hipArraySurfaceLoadStore: + case hipArrayTextureGather: + assert(0); + break; + case hipArrayDefault: + case hipArrayCubemap: + default: + hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,extent.width,extent.height,extent.depth,channelOrder,channelType,ptr,imageInfo); + array[0]->textureType = hipTextureType3D; + break; } - return ihipLogStatus(hip_status); } @@ -1283,7 +1102,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); hc::completion_future marker; - try { stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false); } catch (ihipException& ex) { @@ -1452,15 +1270,9 @@ hipError_t hipMemcpyAtoH(void* dst, hipArray* srcArray, size_t srcOffset, size_t hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) { hipError_t e = hipSuccess; if(p) { - size_t byteSize; - size_t depth; - size_t height; - size_t widthInBytes; - size_t srcPitch; - size_t dstPitch; - void* srcPtr; - void* dstPtr; - size_t ySize; + size_t byteSize, width, height, depth, widthInBytes, srcPitch, dstPitch, ySize; + hipChannelFormatDesc desc; + void* srcPtr;void* dstPtr; if (p->dstArray != nullptr) { if (p->dstArray->isDrv == false) { switch (p->dstArray->desc.f) { @@ -1482,22 +1294,32 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo } depth = p->extent.depth; height = p->extent.height; + width = p->extent.width; widthInBytes = p->extent.width * byteSize; srcPitch = p->srcPtr.pitch; srcPtr = p->srcPtr.ptr; ySize = p->srcPtr.ysize; - dstPitch = p->dstArray->width * byteSize; + desc = p->dstArray->desc; dstPtr = p->dstArray->data; } else { depth = p->Depth; height = p->Height; widthInBytes = p->WidthInBytes; - dstPitch = p->dstArray->width * 4; + width = p->dstArray->width; + desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned); 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, 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; @@ -1509,6 +1331,7 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo ySize = p->srcPtr.ysize; dstPitch = p->dstPtr.pitch; } + stream = ihipSyncAndResolveStream(stream); hc::completion_future marker; try { diff --git a/src/hip_texture.cpp b/src/hip_texture.cpp index 8820878fc0..7ea8da8624 100644 --- a/src/hip_texture.cpp +++ b/src/hip_texture.cpp @@ -222,7 +222,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; void* devPtr = nullptr; - + size_t pitch = 0; switch (pResDesc->resType) { case hipResourceTypeArray: devPtr = pResDesc->res.array.array->data; @@ -279,6 +279,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou imageDescriptor.depth = 0; imageDescriptor.array_size = 0; imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2D; + pitch = pResDesc->res.pitch2D.pitchInBytes; getChannelOrderAndType(pResDesc->res.pitch2D.desc, pTexDesc->readMode, &channelOrder, &channelType); break; @@ -296,7 +297,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW; 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, pitch, 0, &(pTexture->image)) || HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { return ihipLogStatus(hipErrorRuntimeOther); @@ -456,7 +457,7 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, - size_t width, size_t height, textureReference* tex) { + size_t width, size_t height, textureReference* tex, size_t pitch) { hipError_t hip_status = hipSuccess; enum hipTextureAddressMode addressMode = tex->addressMode[0]; enum hipTextureFilterMode filterMode = tex->filterMode; @@ -504,7 +505,7 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode 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, pitch, 0, &(pTexture->image)) || HSA_STATUS_SUCCESS != hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) { return hipErrorRuntimeOther; @@ -522,8 +523,12 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d size_t pitch) { HIP_INIT_API(hipBindTexture2D, offset, tex, devPtr, desc, width, height, pitch); hipError_t hip_status = hipSuccess; + + //TODO: Fix when HSA accepts user defined pitch + if(pitch % 64) pitch =0; + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, offset, devPtr, - desc, width, height, tex); + desc, width, height, tex, pitch); return ihipLogStatus(hip_status); } @@ -797,7 +802,10 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT size_t offset; hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. + //TODO: Fix when HSA accepts user defined pitch + if(pitch % 64) pitch =0; + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, &offset, devPtr, - NULL, desc->Width, desc->Height, tex); + NULL, desc->Width, desc->Height, tex, pitch); return ihipLogStatus(hip_status); } diff --git a/tests/src/test_common.h b/tests/src/test_common.h index 67a8e5e60a..92528e38e3 100644 --- a/tests/src/test_common.h +++ b/tests/src/test_common.h @@ -153,6 +153,36 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg); unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N); +template // pointer type +void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) +{ + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + int offset = i*width*height + j*width + k; + if (hData[offset] != hOutputData[offset]) { + std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" << hOutputData[offset]<<" "; + failed("mistmatch at:%d %d %d",i,j,k); + } + } + } + } +} + +template +void checkArray(T input, T output, size_t height, size_t width) +{ + for(int i=0; i __global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { diff --git a/tests/src/texture/hipBindTex2DPitch.cpp b/tests/src/texture/hipBindTex2DPitch.cpp new file mode 100644 index 0000000000..f463942329 --- /dev/null +++ b/tests/src/texture/hipBindTex2DPitch.cpp @@ -0,0 +1,87 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/*HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "test_common.h" + +#define SIZE_H 8 +#define SIZE_W 12 +#define TYPE_t float +texture tex; + +// texture object is a kernel argument +__global__ void texture2dCopyKernel( TYPE_t* dst) { + + int x = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int y = hipThreadIdx_y + hipBlockIdx_y * hipBlockDim_y; + if ( (x< SIZE_W) && (y< SIZE_H) ){ + dst[SIZE_W*y+x] = tex2D(tex, x, y); + } +} + +int main (void) +{ + TYPE_t* B; + TYPE_t* A; + TYPE_t* devPtrB; + TYPE_t* devPtrA; + + B = new TYPE_t[SIZE_H*SIZE_W]; + A = new TYPE_t[SIZE_H*SIZE_W]; + for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){ + A[i-1] = i; + } + + size_t devPitchA, tex_ofs; + HIPCHECK(hipMallocPitch((void**)&devPtrA, &devPitchA ,SIZE_W*sizeof(TYPE_t), SIZE_H)) ; + HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); + + tex.normalized = false; + #if defined(__HIP_PLATFORM_NVCC__) + + cudaError_t status = cudaBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, + SIZE_W, SIZE_H, devPitchA); + if (status != cudaSuccess) { + printf("%serror: '%s'(%d) at %s:%d%s\n", KRED, cudaGetErrorString(status), + status,__FILE__, __LINE__, KNRM); + failed("API returned error code."); + } + #else + HIPCHECK(hipBindTexture2D(&tex_ofs, &tex, devPtrA, &tex.channelDesc, + SIZE_W, SIZE_H, devPitchA)); + #endif + HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(4,4,1), dim3(32,32,1), 0, 0, devPtrB); + hipDeviceSynchronize(); + HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost)); + + HipTest::checkArray(A, B, SIZE_H, SIZE_W); + delete []A; + delete []B; + hipFree(devPtrA); + hipFree(devPtrB); + passed(); +} diff --git a/tests/src/texture/hipNormalizedFloatValueTex.cpp b/tests/src/texture/hipNormalizedFloatValueTex.cpp index fd2e8ccf0f..0a9c879376 100644 --- a/tests/src/texture/hipNormalizedFloatValueTex.cpp +++ b/tests/src/texture/hipNormalizedFloatValueTex.cpp @@ -26,15 +26,6 @@ THE SOFTWARE. * HIT_END */ -#define CHECK(cmd) \ -{\ - hipError_t error = cmd;\ - if(error != hipSuccess) {\ - fprintf(stderr, "error: '%s' (%d) at %s:%d\n", hipGetErrorString(error), error, __FILE__, __LINE__);\ - exit(EXIT_FAILURE);\ - }\ -} -#include "hip/hip_runtime.h" #include "test_common.h" #define SIZE 10 @@ -54,31 +45,31 @@ bool textureTest(enum hipArray_Format texFormat) { T hData[] = {65, 66, 67, 68, 69, 70, 71, 72,73,74}; T *dData = NULL; - CHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE)); - CHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE)); + HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE)); + HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE)); textureReference* texRef = &textureNormalizedVal_1D; - CHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp)); - CHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp)); - CHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint)); - CHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES)); - CHECK(hipTexRefSetFormat(texRef, texFormat, 1)); + HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp)); + HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp)); + HIPCHECK(hipTexRefSetFilterMode(texRef, hipFilterModePoint)); + HIPCHECK(hipTexRefSetFlags(texRef, HIP_TRSF_NORMALIZED_COORDINATES)); + HIPCHECK(hipTexRefSetFormat(texRef, texFormat, 1)); HIP_ARRAY_DESCRIPTOR desc; desc.Width = SIZE; desc.Height = 1; desc.Format = texFormat; desc.NumChannels = 1; - CHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE)); + HIPCHECK(hipTexRefSetAddress2D(texRef, &desc, (hipDeviceptr_t)dData, sizeof(T)*SIZE)); bool testResult = true; float *dOutputData = NULL; - CHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE)); + HIPCHECK(hipMalloc((void **) &dOutputData, sizeof(float)*SIZE)); hipLaunchKernelGGL(HIP_KERNEL_NAME(normalizedValTextureTest), dim3(1,1,1), dim3(SIZE,1,1), 0, 0, SIZE, dOutputData); float *hOutputData = new float[SIZE]; - CHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE))); + HIPCHECK(hipMemcpyDtoH(hOutputData, (hipDeviceptr_t)dOutputData, (sizeof(float)*SIZE))); for(int i = 0; i < SIZE; i++) { @@ -100,9 +91,9 @@ int main(int argc, char** argv) { int device = 0; bool status = true; - CHECK(hipSetDevice(device)); + HIPCHECK(hipSetDevice(device)); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, device)); + HIPCHECK(hipGetDeviceProperties(&props, device)); std::cout << "Device :: " << props.name << std::endl; #ifdef __HIP_PLATFORM_HCC__ std::cout << "Arch - AMD GPU :: " << props.gcnArch << std::endl; diff --git a/tests/src/texture/hipTexObjPitch.cpp b/tests/src/texture/hipTexObjPitch.cpp new file mode 100644 index 0000000000..47648d5e73 --- /dev/null +++ b/tests/src/texture/hipTexObjPitch.cpp @@ -0,0 +1,106 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/*HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "test_common.h" + +#define SIZE_H 20 +#define SIZE_W 179 +// texture object is a kernel argument +template +__global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) { + + for(int i =0;i(texObj, j, i); + __syncthreads(); +} + +template +void texture2Dtest() +{ + TYPE_t* B; + TYPE_t* A; + TYPE_t* devPtrB; + TYPE_t* devPtrA; + + B = new TYPE_t[SIZE_H*SIZE_W]; + A = new TYPE_t[SIZE_H*SIZE_W]; + for(size_t i=1; i <= (SIZE_H*SIZE_W); i++){ + A[i-1] = i; + } + + size_t devPitchA; + HIPCHECK(hipMallocPitch((void**)&devPtrA, &devPitchA ,SIZE_W*sizeof(TYPE_t), SIZE_H)) ; + HIPCHECK(hipMemcpy2D(devPtrA, devPitchA, A, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyHostToDevice)); + + // Use the texture object + hipResourceDesc texRes; + hipMemset(&texRes, 0, sizeof(texRes)); + texRes.resType = hipResourceTypePitch2D; + texRes.res.pitch2D.devPtr = devPtrA; + texRes.res.pitch2D.height = SIZE_H; + texRes.res.pitch2D.width = SIZE_W; + texRes.res.pitch2D.pitchInBytes = devPitchA; + texRes.res.pitch2D.desc = hipCreateChannelDesc(); + + hipTextureDesc texDescr; + hipMemset(&texDescr, 0, sizeof(texDescr)); + texDescr.normalizedCoords = false; + texDescr.filterMode = hipFilterModePoint; + texDescr.mipmapFilterMode = hipFilterModePoint; + texDescr.addressMode[0] = hipAddressModeClamp; + texDescr.addressMode[1] = hipAddressModeClamp; + texDescr.addressMode[2] = hipAddressModeClamp; + texDescr.readMode = hipReadModeElementType; + + hipTextureObject_t texObj; + hipResourceViewDesc resDesc; + HIPCHECK( hipCreateTextureObject(&texObj, &texRes, &texDescr, &resDesc)); + + HIPCHECK(hipMalloc((void**)&devPtrB, SIZE_W*sizeof(TYPE_t)*SIZE_H)) ; + + hipLaunchKernelGGL(texture2dCopyKernel, dim3(1,1,1), dim3(1,1,1), 0, 0, + texObj, devPtrB, devPtrA); + + HIPCHECK(hipMemcpy2D(B, SIZE_W*sizeof(TYPE_t), devPtrB, SIZE_W*sizeof(TYPE_t), + SIZE_W*sizeof(TYPE_t), SIZE_H, hipMemcpyDeviceToHost)); + + HipTest::checkArray(A, B, SIZE_H, SIZE_W); + delete []A; + delete []B; + hipFree(devPtrA); + hipFree(devPtrB); +} + +int main() +{ + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + texture2Dtest(); + passed(); +} diff --git a/tests/src/texture/simpleTexture2DLayered.cpp b/tests/src/texture/simpleTexture2DLayered.cpp new file mode 100644 index 0000000000..e3ba6d9afe --- /dev/null +++ b/tests/src/texture/simpleTexture2DLayered.cpp @@ -0,0 +1,106 @@ +/* +Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ +#include "test_common.h" +typedef float T; +// Texture reference for 2D Layered texture +texture tex2DL; +__global__ void simpleKernelLayeredArray(T* outputData,int width,int height,int layer) +{ + unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; + unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; + outputData[layer*width*height + y*width + x] = tex2DLayered(tex2DL, x, y, layer); +} + +//////////////////////////////////////////////////////////////////////////////// +void runTest(int width,int height,int num_layers,texture *tex) +{ + unsigned int size = width * height * num_layers * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + + for (unsigned int layer = 0; layer < num_layers; layer++){ + for (int i = 0; i < (int)(width * height); i++){ + hData[layer*width*height + i] = i; + } + } + hipChannelFormatDesc channelDesc; + // Allocate array and copy image data + channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindFloat); + hipArray *arr; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, num_layers), hipArrayLayered)); + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width, height, num_layers); + //myparms.kind = hipMemcpyHostToDevice; + HIPCHECK(hipMemcpy3D(&myparms)); + + // set texture parameters + tex->addressMode[0] = hipAddressModeWrap; + tex->addressMode[1] = hipAddressModeWrap; + tex->filterMode = hipFilterModePoint; + tex->normalized = false; + + // Bind the array to the texture + HIPCHECK(hipBindTextureToArray(*tex, arr, channelDesc)); + + // Allocate device memory for result + T* dData = NULL; + hipMalloc((void **) &dData, size); + dim3 dimBlock(8, 8, 1); + dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); + for (unsigned int layer = 0; layer < num_layers; layer++) + hipLaunchKernelGGL(simpleKernelLayeredArray, dimGrid, dimBlock, 0, 0, dData, width, height, layer); + + HIPCHECK(hipDeviceSynchronize()); + // Allocate mem for the result on host side + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + + // copy result from device to host + HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + HipTest::checkArray(hData,hOutputData,width,height,num_layers); + + hipFree(dData); + hipFreeArray(arr); + free(hData); + free(hOutputData); +} + +//////////////////////////////////////////////////////////////////////////////// +// Program main +//////////////////////////////////////////////////////////////////////////////// +int main(int argc, char **argv) +{ + runTest(512,512,5,&tex2DL); + passed(); +} + diff --git a/tests/src/texture/simpleTexture3D.cpp b/tests/src/texture/simpleTexture3D.cpp new file mode 100644 index 0000000000..741b1671af --- /dev/null +++ b/tests/src/texture/simpleTexture3D.cpp @@ -0,0 +1,133 @@ +/* +Copyright (c) 2019-present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ +#include "test_common.h" + +//typedef char T; +const char *sampleName = "simpleTexture3D"; + +// Texture reference for 3D texture +texture texf; +texture texi; +texture texc; + +template +__global__ void simpleKernel3DArray(T* outputData, + int width, + int height,int depth) +{ + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + if(std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i); + else if(std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i); + else if(std::is_same::value) + outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i); + } + } + } +} + +//////////////////////////////////////////////////////////////////////////////// +//! Run a simple test for tex3D +//////////////////////////////////////////////////////////////////////////////// +template +void runTest(int width,int height,int depth,texture *tex) +{ + unsigned int size = width * height * depth * sizeof(T); + T* hData = (T*) malloc(size); + memset(hData, 0, size); + + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + hData[i*width*height + j*width +k] = i*width*height + j*width + k; + } + } + } + + // Allocate array and copy image data + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(T)*8, 0, 0, 0, hipChannelFormatKindSigned); + hipArray *arr; + + HIPCHECK(hipMalloc3DArray(&arr, &channelDesc, make_hipExtent(width, height, depth), hipArrayCubemap)); + hipMemcpy3DParms myparms = {0}; + myparms.srcPos = make_hipPos(0,0,0); + myparms.dstPos = make_hipPos(0,0,0); + myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height); + myparms.dstArray = arr; + myparms.extent = make_hipExtent(width, height, depth); + myparms.kind = hipMemcpyHostToDevice; + HIPCHECK(hipMemcpy3D(&myparms)); + + // set texture parameters + tex->addressMode[0] = hipAddressModeWrap; + tex->addressMode[1] = hipAddressModeWrap; + tex->filterMode = hipFilterModePoint; + tex->normalized = false; + + // Bind the array to the texture + HIPCHECK(hipBindTextureToArray(*tex, arr, channelDesc)); + + // Allocate device memory for result + T* dData = NULL; + hipMalloc((void **) &dData, size); + + hipLaunchKernelGGL(simpleKernel3DArray, dim3(1,1,1), dim3(1,1,1), 0, 0, dData, width, height, depth); + HIPCHECK(hipDeviceSynchronize()); + + // Allocate mem for the result on host side + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + + // copy result from device to host + HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + HipTest::checkArray(hData,hOutputData,width,height,depth); + + hipFree(dData); + hipFreeArray(arr); + free(hData); + free(hOutputData); +} + +//////////////////////////////////////////////////////////////////////////////// +// Program main +//////////////////////////////////////////////////////////////////////////////// +int main(int argc, char **argv) +{ + printf("%s starting...\n", sampleName); + for(int i=1;i<25;i++) + { + runTest(i,i,i,&texf); + runTest(i+1,i,i,&texi); + runTest(i,i+1,i,&texc); + } + passed(); +} +