Fix memory copy in device_grouped_mha_bwd_xdl_cshuffle_kloop_v1.hpp

6fca1efd · Xinya Zhang · 6c2d60d3 · 6fca1efd
Commit 6fca1efd authored Sep 25, 2023 by Xinya Zhang
Hide whitespace changes
Inline Side-by-side

Showing with 5 additions and 5 deletions

include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_bwd_xdl_cshuffle_kloop_v1.hpp ...ice/impl/device_grouped_mha_bwd_xdl_cshuffle_kloop_v1.hpp +5 -5

No files found.
--- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_bwd_xdl_cshuffle_kloop_v1.hpp
+++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_mha_bwd_xdl_cshuffle_kloop_v1.hpp
@@ -934,11 +934,11 @@ struct DeviceGroupedMultiheadAttentionBackward_Kloop_Xdl_CShuffle_V1
            //     some_has_main_k_block_loop |= y;
            // }

-            hipGetErrorString(hipMemcpyAsync(arg.p_workspace_,
-                                             arg.group_kernel_args_.data(),
-                                             arg.group_kernel_args_.size() * sizeof(GroupKernelArg),
-                                             hipMemcpyHostToDevice,
-                                             stream_config.stream_id_));
+            hipGetErrorString(hipMemcpyWithStream(arg.p_workspace_,
+                                                  arg.group_kernel_args_.data(),
+                                                  arg.group_kernel_args_.size() * sizeof(GroupKernelArg),
+                                                  hipMemcpyHostToDevice,
+                                                  stream_config.stream_id_));

            float ave_time = 0;