diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h
index 6787bc7e5e..0510015c42 100644
--- a/projects/hip/src/hip_hcc_internal.h
+++ b/projects/hip/src/hip_hcc_internal.h
@@ -91,6 +91,57 @@ extern int HIP_DUMP_CODE_OBJECT;
 // TODO - remove when this is standard behavior.
 extern int HCC_OPT_FLUSH;
 
+#define IMAGE_PITCH_ALIGNMENT 256
+template <typename T> inline T alignDown(T value, size_t alignment) {
+    return (T)(value & ~(alignment - 1));
+}
+
+template <typename T> inline T* alignDown(T* value, size_t alignment) {
+    return (T*)alignDown((intptr_t)value, alignment);
+}
+
+template <typename T> inline T alignUp(T value, size_t alignment) {
+    return alignDown((T)(value + alignment - 1), alignment);
+}
+
+template <typename T> inline T* alignUp(T* value, size_t alignment) {
+    return (T*)alignDown((intptr_t)(value + alignment - 1), alignment);
+}
+
+size_t getNumChannels(hsa_ext_image_channel_order_t channelOrder) {
+    switch (channelOrder) {
+      case HSA_EXT_IMAGE_CHANNEL_ORDER_RG:
+        return 2;
+      case HSA_EXT_IMAGE_CHANNEL_ORDER_RGB:
+        return 3;
+      case HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA:
+        return 4;
+      case HSA_EXT_IMAGE_CHANNEL_ORDER_R:
+      default:
+        return 1;
+  }
+}
+
+size_t getElementSize(hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType) {
+    size_t bytesPerPixel = getNumChannels(channelOrder);
+    switch (channelType) {
+      case HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8:
+      case HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT8:
+        break;
+
+      case HSA_EXT_IMAGE_CHANNEL_TYPE_SIGNED_INT32:
+      case HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32:
+      case HSA_EXT_IMAGE_CHANNEL_TYPE_FLOAT:
+        bytesPerPixel *= 4;
+        break;
+
+      default:
+        bytesPerPixel *= 2;
+        break;
+    }
+    return bytesPerPixel;
+}
+
 // Class to assign a short TID to each new thread, for HIP debugging purposes.
 class TidInfo {
    public:
diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp
index 43399bfdb5..444e41107a 100644
--- a/projects/hip/src/hip_memory.cpp
+++ b/projects/hip/src/hip_memory.cpp
@@ -37,6 +37,7 @@ __device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
 namespace hip_internal {
 
 namespace {
+
     inline
     const char* hsa_to_string(hsa_status_t err) noexcept
     {
@@ -791,21 +792,13 @@ hipError_t allocImage(TlsData* tls,hsa_ext_image_geometry_t geometry, int width,
       hsa_amd_memory_pool_t* allocRegion = static_cast<hsa_amd_memory_pool_t*>(acc.get_hsa_am_region());
       hsa_amd_memory_pool_get_info(*allocRegion, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &allocGranularity);
 
-      hsa_ext_image_descriptor_t imageDescriptor;
-      imageDescriptor.geometry = geometry;
-      imageDescriptor.width = width;
-      imageDescriptor.height = height;
-      imageDescriptor.depth = depth;
-      imageDescriptor.array_size = array_size;
-      imageDescriptor.format.channel_order = channelOrder;
-      imageDescriptor.format.channel_type = channelType;
+      size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
+      if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
+          imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size) ;
+      else
+          imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth) ;
 
-      hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
-      hsa_status_t status =
-         hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
-      if(imageInfo.size == 0 || HSA_STATUS_SUCCESS != status){
-         return hipErrorRuntimeOther;
-      }
+      imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
       size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
       const unsigned am_flags = 0;
       *ptr = hip_internal::allocAndSharePtr("device_array", imageInfo.size, ctx,
@@ -898,55 +891,31 @@ extern void getChannelOrderAndType(const hipChannelFormatDesc& desc,
 
 hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height, int depth, hipChannelFormatDesc desc, hsa_ext_image_data_info_t &imageInfo,int array_size __dparm(0))
 {
-    hsa_ext_image_descriptor_t imageDescriptor;
-    imageDescriptor.geometry = geometry;
-    imageDescriptor.width = width;
-    imageDescriptor.height = height;
-    imageDescriptor.depth = depth;
-    imageDescriptor.array_size = array_size;
     hsa_ext_image_channel_order_t channelOrder;
     hsa_ext_image_channel_type_t channelType;
     getChannelOrderAndType(desc, hipReadModeElementType, &channelOrder, &channelType);
-    imageDescriptor.format.channel_order = channelOrder;
-    imageDescriptor.format.channel_type = channelType;
-
-    hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
-    // Get the current device agent.
-    hc::accelerator acc;
-    hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
-    if (!agent)
-        return hipErrorInvalidHandle;
-    hsa_status_t status =
-        hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
-    if(HSA_STATUS_SUCCESS != status){
-        return hipErrorRuntimeOther;
-    }
 
+    size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
+    if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
+       imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size);
+    else
+       imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth);
+    imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
     return hipSuccess;
 }
 
 hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,size_t width, size_t height, size_t depth, hsa_ext_image_channel_order_t channelOrder, hsa_ext_image_channel_type_t channelType, hsa_ext_image_data_info_t &imageInfo,size_t array_size __dparm(0))
 {
-    hsa_ext_image_descriptor_t imageDescriptor;
-    imageDescriptor.geometry = geometry;
-    imageDescriptor.width = width;
-    imageDescriptor.height = height;
-    imageDescriptor.depth = depth;
-    imageDescriptor.array_size = array_size;
-    imageDescriptor.format.channel_order = channelOrder;
-    imageDescriptor.format.channel_type = channelType;
 
-    // Get the current device agent.
-    hc::accelerator acc;
-    hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
-    if (!agent)
-        return hipErrorInvalidResourceHandle;
-    hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
-    hsa_status_t status =
-        hsa_ext_image_data_get_info_with_layout(*agent, &imageDescriptor, permission, HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &imageInfo);
-    if(HSA_STATUS_SUCCESS != status){
-        return hipErrorRuntimeOther;
-    }
+    size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
+
+    if(HSA_EXT_IMAGE_GEOMETRY_2DA == geometry)
+       imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (array_size == 0 ? 1 : array_size);
+    else
+       imageInfo.size = rowPitch * (height == 0 ? 1 : height) * (depth == 0 ? 1 : depth);
+
+    imageInfo.alignment = IMAGE_PITCH_ALIGNMENT;
+
     return hipSuccess;
 }
 
diff --git a/projects/hip/src/hip_texture.cpp b/projects/hip/src/hip_texture.cpp
index 7ea8da8624..81eb5ad272 100644
--- a/projects/hip/src/hip_texture.cpp
+++ b/projects/hip/src/hip_texture.cpp
@@ -295,6 +295,10 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou
                               pTexDesc->normalizedCoords);
 
         hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
+
+        if(hipResourceTypePitch2D != pResDesc->resType)
+            pitch = getElementSize(channelOrder, channelType) * alignUp(imageDescriptor.width, IMAGE_PITCH_ALIGNMENT);
+
         if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
                                       *agent, &imageDescriptor, devPtr, permission,
                                       HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) ||
@@ -430,9 +434,11 @@ hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode r
 
         hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
 
+        size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(size, IMAGE_PITCH_ALIGNMENT);
+
         if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
                                       *agent, &imageDescriptor, devPtr, permission,
-                                      HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
+                                      HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &(pTexture->image)) ||
             HSA_STATUS_SUCCESS !=
                 hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
             return hipErrorRuntimeOther;
@@ -503,6 +509,9 @@ hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode
 
         hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
 
+        if( 0 == pitch)
+            pitch = getElementSize(channelOrder, channelType) * alignUp(width, IMAGE_PITCH_ALIGNMENT);
+
         if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
                                       *agent, &imageDescriptor, devPtr, permission,
                                       HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, pitch, 0, &(pTexture->image)) ||
@@ -606,9 +615,11 @@ hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureRea
 
         hsa_access_permission_t permission = HSA_ACCESS_PERMISSION_RW;
 
+        size_t rowPitch = getElementSize(channelOrder, channelType) * alignUp(imageDescriptor.width, IMAGE_PITCH_ALIGNMENT);
+
         if (HSA_STATUS_SUCCESS != hsa_ext_image_create_with_layout(
                                       *agent, &imageDescriptor, array->data, permission,
-                                      HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, 0, 0, &(pTexture->image)) ||
+                                      HSA_EXT_IMAGE_DATA_LAYOUT_LINEAR, rowPitch, 0, &(pTexture->image)) ||
             HSA_STATUS_SUCCESS !=
                 hsa_ext_sampler_create(*agent, &samplerDescriptor, &(pTexture->sampler))) {
             return hipErrorRuntimeOther;