From 00de81e80e77b817c76f28bbc5f11b554a9d82ef Mon Sep 17 00:00:00 2001 From: Sourav Chakraborty Date: Fri, 4 Dec 2020 11:22:10 -0600 Subject: [PATCH] UCP/RNDV/ROCM: Address CR comments #3 --- src/uct/rocm/copy/rocm_copy_ep.c | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/src/uct/rocm/copy/rocm_copy_ep.c b/src/uct/rocm/copy/rocm_copy_ep.c index f80d28097bee..e1752ccdb6f0 100644 --- a/src/uct/rocm/copy/rocm_copy_ep.c +++ b/src/uct/rocm/copy/rocm_copy_ep.c @@ -51,17 +51,17 @@ ucs_status_t uct_rocm_copy_ep_zcopy(uct_ep_h tl_ep, uct_rkey_t rkey, int is_put) { + size_t size = uct_iov_get_length(iov); + uct_rocm_copy_iface_t *iface = ucs_derived_of(tl_ep->iface, uct_rocm_copy_iface_t); + hsa_signal_t signal = iface->hsa_signal; + uct_rocm_copy_key_t *rocm_copy_key = (uct_rocm_copy_key_t *) rkey; + hsa_status_t status; hsa_agent_t agent; void *src_addr, *dst_addr; void *host_ptr, *dev_ptr, *mapped_ptr; size_t offset; - size_t size = uct_iov_get_length(iov); - uct_rocm_copy_iface_t *iface = ucs_derived_of(tl_ep->iface, uct_rocm_copy_iface_t); - hsa_signal_t signal = iface->hsa_signal; - uct_rocm_copy_key_t *rocm_copy_key = (uct_rocm_copy_key_t *) rkey; - ucs_trace("remote addr %p rkey %p size %zu", (void*)remote_addr, (void*)rkey, size); @@ -93,6 +93,7 @@ ucs_status_t uct_rocm_copy_ep_zcopy(uct_ep_h tl_ep, src_addr = dev_ptr; dst_addr = mapped_ptr; } + hsa_signal_store_screlease(signal, 1); ucs_trace("hsa async copy from src %p to dst %p, len %ld", src_addr, dst_addr, size); @@ -109,9 +110,9 @@ ucs_status_t uct_rocm_copy_ep_get_zcopy(uct_ep_h tl_ep, const uct_iov_t *iov, si uint64_t remote_addr, uct_rkey_t rkey, uct_completion_t *comp) { - ucs_status_t status; size_t size = uct_iov_get_length(iov); uct_rocm_copy_iface_t *iface = ucs_derived_of(tl_ep->iface, uct_rocm_copy_iface_t); + ucs_status_t status; if (size < iface->config.d2h_thresh) { uct_rocm_memcpy_d2h(iov->buffer, (void *)remote_addr, size); @@ -131,9 +132,9 @@ ucs_status_t uct_rocm_copy_ep_put_zcopy(uct_ep_h tl_ep, const uct_iov_t *iov, si uint64_t remote_addr, uct_rkey_t rkey, uct_completion_t *comp) { - ucs_status_t status; size_t size = uct_iov_get_length(iov); uct_rocm_copy_iface_t *iface = ucs_derived_of(tl_ep->iface, uct_rocm_copy_iface_t); + ucs_status_t status; if (size < iface->config.h2d_thresh) { uct_rocm_memcpy_h2d((void *)remote_addr, iov->buffer, size);