SWDEV-269246 - Implementation of hipDrvMemcpy2DUnaligned

Change-Id: I86fdd3b930cbd74c45cd31944f0ea52c0ff65a59

hipamd/include/hip/amd_detail/hip_runtime_api.h

@@ -3508,6 +3508,14 @@ hipError_t hipLaunchKernel(const void* function_address,
                            void** args,
                            size_t sharedMemBytes __dparm(0),
                            hipStream_t stream __dparm(0));
+/**
+ * Copies memory for 2D arrays.
+ *
+ * @param pCopy           - Parameters for the memory copy
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ */
+hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy);
 
 //TODO: Move this to hip_ext.h
 hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks,

hipamd/rocclr/hip_hcc.def.in

@@ -265,6 +265,7 @@ hipExtStreamCreateWithCUMask
 hipStreamGetPriority
 hipMemcpy2DFromArray
 hipMemcpy2DFromArrayAsync
+hipDrvMemcpy2DUnaligned
 hipMemcpyAtoH
 hipMemcpyHtoA
 hipMemcpyParam2DAsync

hipamd/rocclr/hip_hcc.map.in

@@ -96,6 +96,7 @@ global:
     hipMemcpy2D;
     hipMemcpy2DAsync;
     hipMemcpy2DToArray;
+    hipDrvMemcpy2DUnaligned;
     hipMemcpy3D;
     hipMemcpy3DAsync;
     hipDrvMemcpy3D;

hipamd/rocclr/hip_memory.cpp

@@ -2392,3 +2392,12 @@ hipError_t hipFreeHost(void *ptr) {
 
   HIP_RETURN(ihipFree(ptr));
 }
+
+hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy) {
+  HIP_INIT_API(hipDrvMemcpy2DUnaligned, pCopy);
+
+  HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*pCopy);
+
+  HIP_RETURN(ihipMemcpyParam3D(&desc, nullptr));
+}
+