diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 3a81305ba3..624f3615e1 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -2063,6 +2063,45 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const
 hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
                               size_t count, hipMemcpyKind kind);
 
+/**
+ *  @brief Copies data between host and device.
+ *
+ *  @param[in]   dst       Destination memory address
+ *  @param[in]   dpitch    Pitch of destination memory
+ *  @param[in]   src       Source memory address
+ *  @param[in]   wOffset   Source starting X offset
+ *  @param[in]   hOffset   Source starting Y offset
+ *  @param[in]   width     Width of matrix transfer (columns in bytes)
+ *  @param[in]   height    Height of matrix transfer (rows)
+ *  @param[in]   kind      Type of transfer
+ *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
+ * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
+ *
+ *  @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol,
+ * hipMemcpyAsync
+ */
+hipError_t hipMemcpy2DFromArray( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind);
+
+/**
+ *  @brief Copies data between host and device asynchronously.
+ *
+ *  @param[in]   dst       Destination memory address
+ *  @param[in]   dpitch    Pitch of destination memory
+ *  @param[in]   src       Source memory address
+ *  @param[in]   wOffset   Source starting X offset
+ *  @param[in]   hOffset   Source starting Y offset
+ *  @param[in]   width     Width of matrix transfer (columns in bytes)
+ *  @param[in]   height    Height of matrix transfer (rows)
+ *  @param[in]   kind      Type of transfer
+ *  @param[in]   stream    Accelerator view which the copy is being enqueued
+ *  @return      #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue,
+ * #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
+ *
+ *  @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol,
+ * hipMemcpyAsync
+ */
+hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream __dparm(0));
+
 /**
  *  @brief Copies data between host and device.
  *
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index c8369685ec..140e5bb319 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1827,6 +1827,24 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
     return ihipLogStatus(e);
 }
 
+hipError_t ihip2dOffsetMemcpy(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
+                            size_t height, size_t srcXOffsetInBytes, size_t srcYOffset,
+                            size_t dstXOffsetInBytes, size_t dstYOffset,hipMemcpyKind kind,
+                            hipStream_t stream, bool isAsync) {
+    if((spitch < width + srcXOffsetInBytes) || (srcYOffset >= height)){
+        return hipErrorInvalidValue;
+    } else if((dpitch < width + dstXOffsetInBytes) || (dstYOffset >= height)){
+        return hipErrorInvalidValue;
+    }
+    src = (void*)((char*)src+ srcYOffset*spitch + srcXOffsetInBytes);
+    dst = (void*)((char*)dst+ dstYOffset*dpitch + dstXOffsetInBytes);
+    if(isAsync){
+        return ihipMemcpy2DAsync(dst, dpitch, src, spitch, width, height, hipMemcpyDefault, stream);
+    } else{
+        return ihipMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyDefault);
+    }
+}
+
 hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool isAsync) {
     if (pCopy == nullptr) {
         return hipErrorInvalidValue;
@@ -1864,18 +1882,10 @@ hipError_t ihipMemcpyParam2D(const hip_Memcpy2D* pCopy, hipStream_t stream, bool
         default:
             return hipErrorInvalidValue;
     }
-    if(pCopy->srcPitch < pCopy->WidthInBytes + pCopy->srcXInBytes || pCopy->srcY >= pCopy->Height){
-        return hipErrorInvalidValue;
-    } else if(pCopy->dstPitch < pCopy->WidthInBytes + pCopy->dstXInBytes || pCopy->dstY >= pCopy->Height){
-        return hipErrorInvalidValue;
-    }
-    src = (void*)((char*)src+pCopy->srcY*pCopy->srcPitch + pCopy->srcXInBytes);
-    dst = (void*)((char*)dst+pCopy->dstY*pCopy->dstPitch + pCopy->dstXInBytes);
-    if(isAsync){
-        return ihipMemcpy2DAsync(dst, dpitch, src, spitch, pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault, stream);
-    } else{
-        return ihipMemcpy2D(dst, dpitch, src, spitch, pCopy->WidthInBytes, pCopy->Height, hipMemcpyDefault);
-    }
+    return ihip2dOffsetMemcpy(dst, dpitch, src, spitch, pCopy->WidthInBytes,
+                            pCopy->Height, pCopy->srcXInBytes, pCopy->srcY,
+                            pCopy->dstXInBytes, pCopy->dstY, hipMemcpyDefault,
+                            stream, isAsync);
 }
 
 hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
@@ -1888,6 +1898,60 @@ hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream)
     return ihipLogStatus(ihipMemcpyParam2D(pCopy, stream, true));
 }
 
+hipError_t hipMemcpy2DFromArray( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind ){
+    HIP_INIT_SPECIAL_API(hipMemcpyParam2D, (TRACE_MCMD), dst, dpitch, src, wOffset, hOffset, width, height, kind);
+    size_t byteSize;
+    if(src) {
+        switch (src->desc.f) {
+            case hipChannelFormatKindSigned:
+                byteSize = sizeof(int);
+                break;
+            case hipChannelFormatKindUnsigned:
+                byteSize = sizeof(unsigned int);
+                break;
+            case hipChannelFormatKindFloat:
+                byteSize = sizeof(float);
+                break;
+            case hipChannelFormatKindNone:
+                byteSize = sizeof(size_t);
+                break;
+            default:
+                byteSize = 0;
+                break;
+        }
+    } else {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+    return ihipLogStatus(ihip2dOffsetMemcpy(dst, dpitch, src->data, src->width*byteSize, width, height, wOffset, hOffset, 0, 0, kind, hipStreamNull, false));
+}
+
+hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream ){
+    HIP_INIT_SPECIAL_API(hipMemcpyParam2D, (TRACE_MCMD), dst, dpitch, src, wOffset, hOffset, width, height, kind);
+    size_t byteSize;
+    if(src) {
+        switch (src->desc.f) {
+            case hipChannelFormatKindSigned:
+                byteSize = sizeof(int);
+                break;
+            case hipChannelFormatKindUnsigned:
+                byteSize = sizeof(unsigned int);
+                break;
+            case hipChannelFormatKindFloat:
+                byteSize = sizeof(float);
+                break;
+            case hipChannelFormatKindNone:
+                byteSize = sizeof(size_t);
+                break;
+            default:
+                byteSize = 0;
+                break;
+        }
+    } else {
+        return ihipLogStatus(hipErrorInvalidValue);
+    }
+    return ihipLogStatus(ihip2dOffsetMemcpy(dst, dpitch, src->data, src->width*byteSize, width, height, wOffset, hOffset, 0, 0, kind, stream, true));
+}
+
 // TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
 hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) {
     HIP_INIT_SPECIAL_API(hipMemsetAsync, (TRACE_MCMD), dst, value, sizeBytes, stream);