diff --git a/src/common/algorithm.cuh b/src/common/algorithm.cuh
index 53acc65e1..b1c5a4271 100644
--- a/src/common/algorithm.cuh
+++ b/src/common/algorithm.cuh
@@ -148,8 +148,13 @@ void SegmentedArgSort(Context const *ctx, Span<U> values, Span<V> group_ptr,
       sorted_idx_out.data().get(), sorted_idx.size(), n_groups, group_ptr.data(),
       group_ptr.data() + 1, ctx->CUDACtx()->Stream());
 
+#if defined(XGBOOST_USE_HIP)
+  dh::safe_cuda(hipMemcpyAsync(sorted_idx.data(), sorted_idx_out.data().get(),
+                                sorted_idx.size_bytes(), hipMemcpyDeviceToDevice));
+#else
   dh::safe_cuda(cudaMemcpyAsync(sorted_idx.data(), sorted_idx_out.data().get(),
                                 sorted_idx.size_bytes(), cudaMemcpyDeviceToDevice));
+#endif
 }
 
 /**