From 67731a083ac77b4c368329668277ef406a6fc20e Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Wed, 25 Apr 2018 12:32:11 -0400
Subject: [PATCH] P4 to Git Change 1546224 by
 skudchad@skudchad_test2_win_opencl on 2018/04/25 12:24:40

	SWDEV-145570 - [HIP] - Add hipMemcpyTo/FromArray APIs.

	ReviewBoardURL = http://ocltc.amd.com/reviews/r/14696/diff/

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#12 edit


[ROCm/hip commit: 18c7c164c6a1182fc70a75d6b590b45bb8b80d48]
---
 projects/hip/api/hip/hip_memory.cpp | 84 ++++++++++++++++++++++++++---
 1 file changed, 78 insertions(+), 6 deletions(-)

diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp
index 68eb04ac70..23a631881f 100644
--- a/projects/hip/api/hip/hip_memory.cpp
+++ b/projects/hip/api/hip/hip_memory.cpp
@@ -555,22 +555,94 @@ hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, con
   return hipErrorUnknown;
 }
 
-hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
+hipError_t hipMemcpyToArray(hipArray* dstArray, size_t wOffset, size_t hOffset, const void* src,
                             size_t count, hipMemcpyKind kind) {
-  HIP_INIT_API(dst, wOffset, hOffset, src, count, kind);
+  HIP_INIT_API(dstArray, wOffset, hOffset, src, count, kind);
 
-  assert(0 && "Unimplemented");
+  amd::Device* device = g_context->devices()[0];
 
-  return hipErrorUnknown;
+  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
+                                             amd::CommandQueue::RealTimeDisabled,
+                                             amd::CommandQueue::Priority::Normal);
+  if (!queue) {
+    return hipErrorOutOfMemory;
+  }
+
+  amd::Command* command;
+  amd::Command::EventWaitList waitList;
+  amd::Memory* memory;
+
+  amd::Coord3D dstOffset(wOffset, hOffset, 0);
+
+  switch (kind) {
+  case hipMemcpyDeviceToHost:
+    assert(!"Invalid case");
+    /* fall thru */
+  case hipMemcpyHostToDevice:
+    memory = amd::SvmManager::FindSvmBuffer(dstArray->data);
+    command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER, waitList,
+      *memory->asBuffer(), dstOffset, count, src);
+    break;
+  default:
+    assert(!"Shouldn't reach here");
+    break;
+  }
+  if (!command) {
+    return hipErrorOutOfMemory;
+  }
+
+  command->enqueue();
+  command->awaitCompletion();
+  command->release();
+
+  queue->release();
+
+  return hipSuccess;
 }
 
 hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
                               size_t count, hipMemcpyKind kind) {
   HIP_INIT_API(dst, srcArray, wOffset, hOffset, count, kind);
 
-  assert(0 && "Unimplemented");
+  amd::Device* device = g_context->devices()[0];
 
-  return hipErrorUnknown;
+  amd::HostQueue* queue = new amd::HostQueue(*g_context, *device, 0,
+                                             amd::CommandQueue::RealTimeDisabled,
+                                             amd::CommandQueue::Priority::Normal);
+  if (!queue) {
+    return hipErrorOutOfMemory;
+  }
+
+  amd::Command* command;
+  amd::Command::EventWaitList waitList;
+  amd::Memory* memory;
+
+  amd::Coord3D srcOffset(wOffset, hOffset, 0);
+
+  switch (kind) {
+  case hipMemcpyHostToDevice:
+    assert(!"Invalid case");
+    /* fall thru */
+  case hipMemcpyDeviceToHost:
+    memory = amd::SvmManager::FindSvmBuffer(srcArray->data);
+    command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER, waitList,
+      *memory->asBuffer(), srcOffset, count, dst);
+    break;
+  default:
+    assert(!"Shouldn't reach here");
+    break;
+  }
+  if (!command) {
+    return hipErrorOutOfMemory;
+  }
+
+  command->enqueue();
+  command->awaitCompletion();
+  command->release();
+
+  queue->release();
+
+  return hipSuccess;
 }
 
 hipError_t hipMemcpyHtoA(hipArray* dstArray, size_t dstOffset, const void* srcHost, size_t count) {