General sync memcpy improvements. Add hipMemcpyWithStream (#1673)

* General sync memcpy improvements. Add `hipMemcpyWithStream`

* Update hip_memory.cpp

include/hip/hcc_detail/hip_runtime_api.h

@@ -1277,7 +1277,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
  *  @param[in]  height Requested pitched allocation height
  *
  *  If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
- *  The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array. 
+ *  The intended usage of pitch is as a separate parameter of the allocation, used to compute addresses within the 2D array.
  *  Given the row and column of an array element of type T, the address is computed as:
  *  T* pElement = (T*)((char*)BaseAddress + Row * Pitch) + Column;
  *
@@ -1361,6 +1361,10 @@ hipError_t hipHostFree(void* ptr);
  */
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
 
+// TODO: Add description
+hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
+                               hipMemcpyKind kind, hipStream_t stream);
+
 /**
  *  @brief Copy data from Host to Device
  *
@@ -1498,7 +1502,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
 #else
 hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
 
-#ifdef __cplusplus //Start : Not supported in gcc 
+#ifdef __cplusplus //Start : Not supported in gcc
 namespace hip_impl {
 inline
 __attribute__((visibility("hidden")))

include/hip/nvcc_detail/hip_runtime_api.h

@@ -610,6 +610,16 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
         cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind)));
 }
 
+inline hipError_t hipMemcpyWithStream(void* dst, const void* src,
+                                      size_t sizeBytes, hipMemcpyKind copyKind,
+                                      hipStream_t stream) {
+    cudaError_t error = cudaMemcpyAsync(dst, src, sizeBytes,
+                                        hipMemcpyKindToCudaMemcpyKind(copyKind),
+                                        stream);
+    if (error != cudaSuccess) return hipCUDAErrorTohipError(error);
+
+    return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
+}
 
 inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
                                         hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) {