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 [ROCm/clr commit: dc8f556460]
2019-11-07 13:17:46 +05:30
parent cf648f7dbe
commit f5a92e5cb6
9 changed files with 698 additions and 414 deletions
@@ -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 <class T, int dim, enum hipTextureReadMode readMode>
 hipError_t hipBindTexture2D(size_t* offset, struct texture<T, dim, readMode>& tex,
@@ -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<hsa_agent_t*>(acc.get_hsa_agent());
+      if (!agent)
+         return hipErrorInvalidResourceHandle;
+      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;
+      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<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);
-
-        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<hsa_agent_t*>(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<hsa_agent_t*>(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<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);
+        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<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);
-
-            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<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);
-
-        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<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);
-
-        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 {
@@ -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);
 }
@@ -153,6 +153,36 @@ int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg);

 unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N);

+template<typename T> // 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<typename T> 
+void checkArray(T input, T output, size_t height, size_t width)
+{
+    for(int i=0; i<height; i++ ){
+        for(int j=0; j<width; j++ ){
+            int offset = i*width + j;
+            if( input[offset] !=  output[offset] ){
+                 std::cerr << '[' << i << ',' << j << ',' << "]:" << input[offset] << "----" << output[offset]<<"  ";
+                 failed("mistmatch at:%d %d",i,j);
+            }
+        }
+    }
+}
+

 template <typename T>
 __global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) {
@@ -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<TYPE_t, 2, hipReadModeElementType> 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();
+}
@@ -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;
@@ -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 <typename TYPE_t>
+__global__ void texture2dCopyKernel( hipTextureObject_t texObj, TYPE_t* dst,TYPE_t* A) {
+
+    for(int i =0;i<SIZE_H;i++)
+        for(int j = 0;j<SIZE_W;j++)
+            dst[SIZE_W*i+j] = tex2D<TYPE_t>(texObj, j, i);
+    __syncthreads();
+}
+
+template <typename TYPE_t>
+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<TYPE_t>();
+
+    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<float>();
+    texture2Dtest<int>();
+    texture2Dtest<unsigned char>();
+    texture2Dtest<short>();
+    texture2Dtest<char>();
+    texture2Dtest<unsigned int>();
+    passed();
+}
@@ -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<float, hipTextureType2DLayered> 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<T, hipTextureType2DLayered> *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();
+}
+
@@ -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<float, hipTextureType3D, hipReadModeElementType> texf;
+texture<int, hipTextureType3D, hipReadModeElementType> texi;
+texture<char, hipTextureType3D, hipReadModeElementType> texc;
+
+template <typename T>
+__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<T, float>::value)
+                   outputData[i*width*height + j*width + k] = tex3D(texf, texf.textureObject, k, j, i);
+               else if(std::is_same<T, int>::value)
+                   outputData[i*width*height + j*width + k] = tex3D(texi, texi.textureObject, k, j, i);
+               else if(std::is_same<T, char>::value)
+                   outputData[i*width*height + j*width + k] = tex3D(texc, texc.textureObject, k, j, i);
+           }
+       }
+    }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+//! Run a simple test for tex3D
+////////////////////////////////////////////////////////////////////////////////
+template <typename T>
+void runTest(int width,int height,int depth,texture<T, hipTextureType3D, hipReadModeElementType> *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<float>(i,i,i,&texf);
+        runTest<int>(i+1,i,i,&texi);
+        runTest<char>(i,i+1,i,&texc);
+    }
+    passed();
+}
+