From 68abfb4aff908c0dc62c2cbe3a1affdda747280c Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Thu, 18 Mar 2021 11:25:55 +0000
Subject: [PATCH] SWDEV-277152 - Add hipMemcpy2DToArrayAsync API in headers

Also, address few scenarios for hipMemcpy2DToArray/Async to return
proper error types

Change-Id: Idbadf666a5e7bebc0f3ea8b7048b79208e04cf6b
---
 .../include/hip/amd_detail/hip_runtime_api.h  | 22 +++++++++++++++++++
 hipamd/rocclr/hip_hcc.def.in                  |  1 +
 hipamd/rocclr/hip_hcc.map.in                  |  1 +
 hipamd/rocclr/hip_memory.cpp                  | 12 ++++++++++
 4 files changed, 36 insertions(+)

diff --git a/hipamd/include/hip/amd_detail/hip_runtime_api.h b/hipamd/include/hip/amd_detail/hip_runtime_api.h
index c6106d45b2..7c7ff1d981 100644
--- a/hipamd/include/hip/amd_detail/hip_runtime_api.h
+++ b/hipamd/include/hip/amd_detail/hip_runtime_api.h
@@ -2463,6 +2463,28 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
 hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
                               size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
 
+/**
+ *  @brief Copies data between host and device.
+ *
+ *  @param[in]   dst     Destination memory address
+ *  @param[in]   wOffset Destination starting X offset
+ *  @param[in]   hOffset Destination starting Y offset
+ *  @param[in]   src     Source memory address
+ *  @param[in]   spitch  Pitch of source memory
+ *  @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, hipMemcpyToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol,
+ * hipMemcpyAsync
+ */
+hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
+                                   size_t spitch, size_t width, size_t height, hipMemcpyKind kind,
+                                   hipStream_t stream __dparm(0));
+
 /**
  *  @brief Copies data between host and device.
  *
diff --git a/hipamd/rocclr/hip_hcc.def.in b/hipamd/rocclr/hip_hcc.def.in
index 595b0ae226..1a4e1a8eb9 100755
--- a/hipamd/rocclr/hip_hcc.def.in
+++ b/hipamd/rocclr/hip_hcc.def.in
@@ -95,6 +95,7 @@ hipMemcpyParam2D
 hipMemcpy2D
 hipMemcpy2DAsync
 hipMemcpy2DToArray
+hipMemcpy2DToArrayAsync
 hipMemcpy3D
 hipMemcpy3DAsync
 hipDrvMemcpy3D
diff --git a/hipamd/rocclr/hip_hcc.map.in b/hipamd/rocclr/hip_hcc.map.in
index 2a513d5ebe..b9a2d64301 100755
--- a/hipamd/rocclr/hip_hcc.map.in
+++ b/hipamd/rocclr/hip_hcc.map.in
@@ -96,6 +96,7 @@ global:
     hipMemcpy2D;
     hipMemcpy2DAsync;
     hipMemcpy2DToArray;
+    hipMemcpy2DToArrayAsync;
     hipDrvMemcpy2DUnaligned;
     hipMemcpy3D;
     hipMemcpy3DAsync;
diff --git a/hipamd/rocclr/hip_memory.cpp b/hipamd/rocclr/hip_memory.cpp
index 836bfbf626..0faf531dba 100755
--- a/hipamd/rocclr/hip_memory.cpp
+++ b/hipamd/rocclr/hip_memory.cpp
@@ -1643,6 +1643,10 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
 }
 
 hipError_t ihipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream, bool isAsync = false) {
+  if (dst == nullptr) {
+    HIP_RETURN(hipErrorInvalidResourceHandle);
+  }
+
   hip_Memcpy2D desc = {};
 
   desc.srcXInBytes = 0;
@@ -1670,6 +1674,10 @@ hipError_t ihipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, c
 hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) {
   HIP_INIT_API(hipMemcpy2DToArray, dst, wOffset, hOffset, src, spitch, width, height, kind);
 
+  if (spitch == 0) {
+    HIP_RETURN(hipErrorInvalidPitchValue);
+  }
+
   HIP_RETURN_DURATION(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, nullptr));
 }
 
@@ -2304,6 +2312,10 @@ hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffs
 hipError_t hipMemcpy2DToArrayAsync(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) {
   HIP_INIT_API(hipMemcpy2DToArrayAsync, dst, wOffset, hOffset, src, spitch, width, height, kind);
 
+  if (spitch == 0) {
+    HIP_RETURN(hipErrorInvalidPitchValue);
+  }
+
   HIP_RETURN_DURATION(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, spitch, width, height, kind, stream, true));
 }