replace hipMemcpy with hipMemcpyWithStream (#734)

include/ck/tensor_operation/gpu/device/device_grouped_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

@@ -611,10 +611,11 @@ struct DeviceGroupedGemmSoftmaxGemmPermute_Xdl_CShuffle
                 some_has_main_k_block_loop |= y;
             }
 
-            hipGetErrorString(hipMemcpy(arg.p_workspace_,
+            hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_,
                                                   arg.group_kernel_args_.data(),
                                                   arg.group_kernel_args_.size() * sizeof(GroupKernelArg),
-                                        hipMemcpyHostToDevice));
+                                                  hipMemcpyHostToDevice,
+                                                  stream_config.stream_id_));
 
             float ave_time = 0;
 

include/ck/tensor_operation/gpu/device/impl/device_grouped_contraction_multiple_d_xdl_cshuffle.hpp

@@ -652,11 +652,12 @@ struct DeviceGroupedContractionMultipleD_Xdl_CShuffle
                 }
             }
 
-            hipGetErrorString(hipMemcpy(arg.p_workspace_,
+            hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_,
                                                   arg.contraction_multi_d_kernel_args_.data(),
                                                   arg.contraction_multi_d_kernel_args_.size() *
                                                   sizeof(ContractionMultiDKernelArg),
-                                        hipMemcpyHostToDevice));
+                                                  hipMemcpyHostToDevice,
+                                                  stream_config.stream_id_));
 
             float ave_time = 0;
 

include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_multiple_d_dl.hpp

@@ -597,10 +597,11 @@ struct DeviceGroupedGemmMultipleD_Dl : public DeviceGroupedGemm<ALayout,
                 }
             }
 
-            hipGetErrorString(hipMemcpy(arg.p_workspace_,
+            hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_,
                                                   arg.gemm_desc_kernel_arg_.data(),
                                                   arg.gemm_desc_kernel_arg_.size() * sizeof(GemmKernelArg),
-                                        hipMemcpyHostToDevice));
+                                                  hipMemcpyHostToDevice,
+                                                  stream_config.stream_id_));
 
             auto launch_kernel = [&](auto has_main_k_block_loop,
                                      auto has_double_tail_k_block_loop) {

include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_xdl.hpp

@@ -549,10 +549,11 @@ struct DeviceGroupedGemm_Xdl : public DeviceGroupedGemm<ALayout,
             }
 
             hipGetErrorString(
-                hipMemcpy(arg.p_workspace_,
+                hipMemcpyWithStream(arg.p_workspace_,
                                     arg.gemm_desc_kernel_arg_.data(),
                                     arg.gemm_desc_kernel_arg_.size() * sizeof(GemmBiasTransKernelArg),
-                          hipMemcpyHostToDevice));
+                                    hipMemcpyHostToDevice,
+                                    stream_config.stream_id_));
 
             float ave_time = 0;
 

include/ck/tensor_operation/gpu/device/impl/device_grouped_gemm_xdl_splitk_cshuffle.hpp

@@ -406,10 +406,11 @@ struct DeviceGroupedGemmXdlSplitKCShuffle : public DeviceGroupedGemmSplitK<ALayo
                 }
             }
 
-            hip_check_error(hipMemcpy(arg.p_workspace_,
+            hip_check_error(hipMemcpyWithStream(arg.p_workspace_,
                                                 arg.gemm_kernel_args_.data(),
                                                 arg.gemm_kernel_args_.size() * sizeof(GemmTransKernelArg),
-                                      hipMemcpyHostToDevice));
+                                                hipMemcpyHostToDevice,
+                                                stream_config.stream_id_));
 
             float ave_time = 0;