Fix late-coming issues. (#1724)

Implementation for hipMemcpyWithStream.


[ROCm/hip commit: 75a11330aa]

projects/hip/include/hip/hcc_detail/grid_launch.h

@@ -48,8 +48,6 @@ typedef struct grid_launch_parm
   //! Value of packet fences to apply to launch.
   //! The correspond to the value of bits 9:14 in the AQL packet,
   //! see HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE and hsa_fence_scope_t.
-  //! Set to -1 for conservative defaults.
-  //! Placeholder, is not used to control packet dispatch yet
   unsigned int  launch_fence;
 
   //! Pointer to the accelerator_view where the kernel should execute.

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -1361,6 +1361,9 @@ 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
  *

projects/hip/include/hip/nvcc_detail/hip_runtime_api.h

@@ -611,6 +611,18 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes,
 }
 
 
+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)) {
     return hipCUDAErrorTohipError(