UCP/RNDV/ROCM: Address CR comments #3

ROCm · Dec 4, 2020 · 00de81e · 00de81e
1 parent 30ff487
commit 00de81e
Showing 1 changed file with 8 additions and 7 deletions.
diff --git 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);