From 8cb232ddefedfc5fa0f746f6329ca803c6e1bc69 Mon Sep 17 00:00:00 2001
From: "Arya.Rafii" <Arya.Rafii@amd.com>
Date: Fri, 26 Feb 2021 09:57:36 -0500
Subject: [PATCH] SWDEV-269246 - Implementation of  hipDrvMemcpy2DUnaligned

Change-Id: I86fdd3b930cbd74c45cd31944f0ea52c0ff65a59
---
 hipamd/include/hip/amd_detail/hip_runtime_api.h | 8 ++++++++
 hipamd/rocclr/hip_hcc.def.in                    | 1 +
 hipamd/rocclr/hip_hcc.map.in                    | 1 +
 hipamd/rocclr/hip_memory.cpp                    | 9 +++++++++
 4 files changed, 19 insertions(+)

diff --git a/hipamd/include/hip/amd_detail/hip_runtime_api.h b/hipamd/include/hip/amd_detail/hip_runtime_api.h
index eb5890759c..8ddd7fd1e3 100644
--- a/hipamd/include/hip/amd_detail/hip_runtime_api.h
+++ b/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,
diff --git a/hipamd/rocclr/hip_hcc.def.in b/hipamd/rocclr/hip_hcc.def.in
index 65d15b3c6c..595b0ae226 100755
--- a/hipamd/rocclr/hip_hcc.def.in
+++ b/hipamd/rocclr/hip_hcc.def.in
@@ -265,6 +265,7 @@ hipExtStreamCreateWithCUMask
 hipStreamGetPriority
 hipMemcpy2DFromArray
 hipMemcpy2DFromArrayAsync
+hipDrvMemcpy2DUnaligned
 hipMemcpyAtoH
 hipMemcpyHtoA
 hipMemcpyParam2DAsync
diff --git a/hipamd/rocclr/hip_hcc.map.in b/hipamd/rocclr/hip_hcc.map.in
index e0b4893e67..2a513d5ebe 100755
--- a/hipamd/rocclr/hip_hcc.map.in
+++ b/hipamd/rocclr/hip_hcc.map.in
@@ -96,6 +96,7 @@ global:
     hipMemcpy2D;
     hipMemcpy2DAsync;
     hipMemcpy2DToArray;
+    hipDrvMemcpy2DUnaligned;
     hipMemcpy3D;
     hipMemcpy3DAsync;
     hipDrvMemcpy3D;
diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp
index 190d7c6d99..3404c55622 100755
--- a/hipamd/rocclr/hip_memory.cpp
+++ b/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));
+}
+