diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 30258d7474..de2d08f4e1 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -835,7 +835,7 @@ hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width
                                HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT32,ptr,imageInfo);
 
     if(hip_status == hipSuccess)
-        *pitch = imageInfo.size/(height == 0 ? 1:height)/(depth == 0 ? 1:depth);
+        *pitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
 
     return hip_status;
 }
@@ -925,6 +925,24 @@ hipError_t GetImageInfo(hsa_ext_image_geometry_t geometry,int width, int height,
     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 = {.geometry = geometry, .width = width, .height= height, .depth = depth,
+        .array_size = array_size, .format.channel_order = channelOrder, .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;
+    }
+    return hipSuccess;
+}
+
 hipError_t ihipArrayToImageFormat(hipArray_Format format,hsa_ext_image_channel_type_t &channelType) {
    switch (format) {
        case HIP_AD_FORMAT_UNSIGNED_INT8:
@@ -1070,6 +1088,7 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAll
        case hipArrayDefault:
        case hipArrayCubemap:
        default:
+          array[0]->type = hipArrayCubemap;
           hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,pAllocateArray->Width,pAllocateArray->Height,
                                   pAllocateArray->Depth,channelOrder,channelType,ptr,imageInfo);
           array[0]->textureType = hipTextureType3D;
@@ -1113,6 +1132,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
        case hipArrayDefault:
        case hipArrayCubemap:
        default:
+          array[0]->type = hipArrayCubemap;
           hip_status = allocImage(tls,HSA_EXT_IMAGE_GEOMETRY_3D,extent.width,extent.height,extent.depth,channelOrder,channelType,ptr,imageInfo);
           array[0]->textureType = hipTextureType3D;
           break;
@@ -1546,25 +1566,46 @@ hipError_t ihipMemcpy3D(const struct hipMemcpy3DParms* p, hipStream_t stream, bo
                 ySize = p->srcPtr.ysize;
                 desc = p->dstArray->desc;
                 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 {
                 depth = p->Depth;
                 height = p->Height;
                 widthInBytes = p->WidthInBytes;
                 width =  p->dstArray->width;
-                desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindSigned);
+                hsa_ext_image_channel_order_t channelOrder;
+                switch(p->dstArray->NumChannels) {
+                    case 2:
+                       channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG;
+                       break;
+                    case 3:
+                       channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGB;
+                       break;
+                    case 4:
+                       channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA;
+                       break;
+                    case 1:
+                    default:
+                       channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R;
+                       break;
+                }
+                hsa_ext_image_channel_type_t channelType;
+                e = ihipArrayToImageFormat(p->dstArray->Format,channelType);
                 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, channelOrder, channelType, imageInfo, depth);
+                else
+                    GetImageInfo(HSA_EXT_IMAGE_GEOMETRY_3D, width, height, depth, channelOrder, channelType, imageInfo);
+                dstPitch = imageInfo.size/(height == 0 ? 1 : height)/(depth == 0 ? 1 : depth);
             }
-            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;