From 2d5140cb27f6beed714f3cac7d88faf6ff2ded59 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 9 Sep 2016 10:21:52 -0500
Subject: [PATCH] Added async memcpy driver api

Change-Id: I90e8a078d668a408e79f9e1142e7534771467a4f
---
 include/hcc_detail/hip_runtime_api.h |  6 ++-
 src/hip_memory.cpp                   | 79 ++++++++++++++++++++++++++++
 tests/src/hipDrvMemcpy.cpp           | 28 ++++++++--
 3 files changed, 107 insertions(+), 6 deletions(-)

diff --git a/include/hcc_detail/hip_runtime_api.h b/include/hcc_detail/hip_runtime_api.h
index 71e46612de..c782696753 100644
--- a/include/hcc_detail/hip_runtime_api.h
+++ b/include/hcc_detail/hip_runtime_api.h
@@ -883,7 +883,11 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes);
 
 hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes);
 
-hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes);
+hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream);
+
+hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream);
+
+hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream);
 
 
 /**
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 8516d7520d..19b4745b59 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -576,6 +576,85 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
     return ihipLogStatus(e);
 }
 
+hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream)
+{
+    HIP_INIT_API(dst, src, sizeBytes, stream);
+
+    hipError_t e = hipSuccess;
+
+    stream = ihipSyncAndResolveStream(stream);
+
+    hipMemcpyKind kind = hipMemcpyHostToDevice;
+
+    if ((dst == NULL) || (src == NULL)) {
+        e= hipErrorInvalidValue;
+    } else if (stream) {
+        try {
+            stream->copyAsync((void*)dst, src, sizeBytes, kind);
+        }
+        catch (ihipException ex) {
+            e = ex._code;
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+}
+
+hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
+{
+    HIP_INIT_API(dst, src, sizeBytes, stream);
+
+    hipError_t e = hipSuccess;
+
+    hipMemcpyKind kind = hipMemcpyDeviceToDevice;
+
+    stream = ihipSyncAndResolveStream(stream);
+
+
+    if ((dst == NULL) || (src == NULL)) {
+        e= hipErrorInvalidValue;
+    } else if (stream) {
+        try {
+            stream->copyAsync((void*)dst, (void*)src, sizeBytes, kind);
+        }
+        catch (ihipException ex) {
+            e = ex._code;
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+}
+
+hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
+{
+    HIP_INIT_API(dst, src, sizeBytes, stream);
+
+    hipError_t e = hipSuccess;
+
+    stream = ihipSyncAndResolveStream(stream);
+
+    hipMemcpyKind kind = hipMemcpyDeviceToHost;
+
+    if ((dst == NULL) || (src == NULL)) {
+        e= hipErrorInvalidValue;
+    } else if (stream) {
+        try {
+            stream->copyAsync(dst, (void*)src, sizeBytes, kind);
+        }
+        catch (ihipException ex) {
+            e = ex._code;
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+}
+
 // dpitch, spitch, and width in bytes
 hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,
         size_t width, size_t height, hipMemcpyKind kind) {
diff --git a/tests/src/hipDrvMemcpy.cpp b/tests/src/hipDrvMemcpy.cpp
index 55c9b18818..8322464964 100644
--- a/tests/src/hipDrvMemcpy.cpp
+++ b/tests/src/hipDrvMemcpy.cpp
@@ -6,21 +6,39 @@
 #define SIZE LEN<<2
 
 int main(){
-  int *A, *B, *C;
-  hipDeviceptr Ad, Bd;
+  int *A, *B;
+  hipDeviceptr_t Ad, Bd;
   A = new int[LEN];
   B = new int[LEN];
-  C = new int[LEN];
+
   for(int i=0;i<LEN;i++){
     A[i] = i;
   }
+
   hipMalloc((void**)&Ad, SIZE);
   hipMalloc((void**)&Bd, SIZE);
+
   hipMemcpyHtoD(Ad, A, SIZE);
   hipMemcpyDtoD(Bd, Ad, SIZE);
   hipMemcpyDtoH(B, Bd, SIZE);
-  hipMemcpyHtoH(C, B,SIZE);
+
   for(int i=0;i<16;i++){
-    std::cout<<A[i]<<" "<<C[i]<<std::endl;
+    std::cout<<A[i]<<" "<<B[i]<<std::endl;
   }
+
+  int *Ah, *Bh;
+  hipHostMalloc(&Ah, SIZE, 0);
+  hipHostMalloc(&Bh, SIZE, 0);
+  memcpy(Ah, A, SIZE);
+  hipStream_t stream;
+  hipStreamCreate(&stream);
+
+  hipMemcpyHtoDAsync(Ad, Ah, SIZE, stream);
+  hipStreamSynchronize(stream);
+  hipMemcpyDtoDAsync(Bd, Ad, SIZE, stream);
+  hipStreamSynchronize(stream);
+  hipMemcpyDtoHAsync(Bh, Bd, SIZE, stream);
+  hipStreamSynchronize(stream);
+
+  std::cout<<Ah[10]<<" "<<Bh[10]<<std::endl;
 }