diff --git a/src/ucp/api/device/ucp_device_impl.h b/src/ucp/api/device/ucp_device_impl.h
index 86388082b84..ce004d7677a 100644
--- a/src/ucp/api/device/ucp_device_impl.h
+++ b/src/ucp/api/device/ucp_device_impl.h
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -105,14 +105,25 @@ UCS_F_DEVICE void ucp_device_request_init(uct_device_ep_t *device_ep,
     })
 
 
+#define UCP_DEVICE_GET_LANE(_handle, _channel_id, _lane, _uct_channel_id) \
+    _lane           = _channel_id % _handle->num_lanes; \
+    _uct_channel_id = _channel_id / _handle->num_lanes;
+
+
+#define UCP_DEVICE_GET_ELEM(_handle, _index) \
+    static_cast<ucs_typeof(_handle->mem_elements[0])*>(UCS_PTR_BYTE_OFFSET( \
+            _handle->mem_elements, \
+            (sizeof(_handle->mem_elements[0]) + \
+             (sizeof(_handle->mem_elements[0].tl[0]) * _handle->num_lanes)) * \
+                    _index))
+
+
 UCS_F_DEVICE ucs_status_t ucp_device_prepare_send_remote(
         const ucp_device_remote_mem_list_h dst_mem_list_h,
-        unsigned dst_mem_list_index, uint64_t &remote_address,
+        unsigned dst_mem_list_index, uint64_t &remote_address, unsigned lane,
         ucp_device_request_t *req, uct_device_ep_t *&device_ep,
-        const uct_device_mem_element_t *&uct_elem,
-        uct_device_completion_t *&comp)
+        const uct_device_mem_elem_t *&uct_elem, uct_device_completion_t *&comp)
 {
-    const size_t elem_size = sizeof(uct_device_remote_mem_list_elem_t);
     ucs_status_t status;
 
     status = ucp_device_check_params(dst_mem_list_h, dst_mem_list_index);
@@ -120,30 +131,26 @@ UCS_F_DEVICE ucs_status_t ucp_device_prepare_send_remote(
         return status;
     }
 
-    const auto dst_mem_element = static_cast<uct_device_remote_mem_list_elem_t*>(
-            UCS_PTR_BYTE_OFFSET(dst_mem_list_h->mem_elements,
-                                dst_mem_list_index * elem_size));
+    const auto dst_mem_element = UCP_DEVICE_GET_ELEM(dst_mem_list_h,
+                                                     dst_mem_list_index);
     remote_address             = dst_mem_element->addr;
-    device_ep                  = dst_mem_element->device_ep;
-    uct_elem                   = &dst_mem_element->uct_mem_element;
+    device_ep                  = dst_mem_element->tl[lane].ep;
+    uct_elem                   = &dst_mem_element->tl[lane].uct;
     ucp_device_request_init(device_ep, req, comp);
 
     return UCS_OK;
 }
 
 
-UCS_F_DEVICE ucs_status_t
-ucp_device_prepare_send(const ucp_device_local_mem_list_h src_mem_list_h,
-                        unsigned src_mem_list_index,
-                        const ucp_device_remote_mem_list_h dst_mem_list_h,
-                        unsigned dst_mem_list_index, const void *&address,
-                        uint64_t &remote_address, ucp_device_request_t *req,
-                        uct_device_ep_t *&device_ep,
-                        const uct_device_local_mem_list_elem_t *&src_uct_elem,
-                        const uct_device_mem_element_t *&uct_elem,
-                        uct_device_completion_t *&comp)
+UCS_F_DEVICE ucs_status_t ucp_device_prepare_send(
+        const ucp_device_local_mem_list_h src_mem_list_h,
+        unsigned src_mem_list_index,
+        const ucp_device_remote_mem_list_h dst_mem_list_h,
+        unsigned dst_mem_list_index, const void *&address,
+        uint64_t &remote_address, unsigned lane, ucp_device_request_t *req,
+        uct_device_ep_t *&device_ep, const uct_device_mem_elem_t *&src_uct_elem,
+        const uct_device_mem_elem_t *&uct_elem, uct_device_completion_t *&comp)
 {
-    const size_t elem_size = sizeof(uct_device_local_mem_list_elem_t);
     ucs_status_t status;
 
     status = ucp_device_check_params(src_mem_list_h, src_mem_list_index);
@@ -152,16 +159,16 @@ ucp_device_prepare_send(const ucp_device_local_mem_list_h src_mem_list_h,
     }
 
     status = ucp_device_prepare_send_remote(dst_mem_list_h, dst_mem_list_index,
-                                            remote_address, req, device_ep,
-                                            uct_elem, comp);
+                                            remote_address, lane, req,
+                                            device_ep, uct_elem, comp);
     if (status != UCS_OK) {
         return status;
     }
 
-    src_uct_elem = static_cast<uct_device_local_mem_list_elem_t*>(
-            UCS_PTR_BYTE_OFFSET(src_mem_list_h->mem_elements,
-                                src_mem_list_index * elem_size));
-    address      = src_uct_elem->addr;
+    const auto src_mem_elem = UCP_DEVICE_GET_ELEM(src_mem_list_h,
+                                                  src_mem_list_index);
+    src_uct_elem            = src_mem_elem->tl + lane;
+    address                 = src_mem_elem->addr;
 
     return UCS_OK;
 }
@@ -212,18 +219,22 @@ ucp_device_put(const ucp_device_local_mem_list_h src_mem_list_h,
                unsigned dst_mem_list_index, size_t dst_offset, size_t length,
                unsigned channel_id, uint64_t flags, ucp_device_request_t *req)
 {
+    unsigned uct_channel_id;
+    unsigned lane;
     const void *address;
-    const uct_device_mem_element_t *uct_elem;
-    const uct_device_local_mem_list_elem_t *src_uct_elem;
+    const uct_device_mem_elem_t *uct_elem;
+    const uct_device_mem_elem_t *src_uct_elem;
     uint64_t remote_address;
     uct_device_completion_t *comp;
     uct_device_ep_t *device_ep;
     ucs_status_t status;
 
+    UCP_DEVICE_GET_LANE(dst_mem_list_h, channel_id, lane, uct_channel_id);
+
     status = ucp_device_prepare_send(src_mem_list_h, src_mem_list_index,
                                      dst_mem_list_h, dst_mem_list_index,
-                                     address, remote_address, req, device_ep,
-                                     src_uct_elem, uct_elem, comp);
+                                     address, remote_address, lane, req,
+                                     device_ep, src_uct_elem, uct_elem, comp);
     if (status != UCS_OK) {
         return status;
     }
@@ -232,7 +243,7 @@ ucp_device_put(const ucp_device_local_mem_list_h src_mem_list_h,
                                     src_uct_elem, uct_elem,
                                     UCS_PTR_BYTE_OFFSET(address, src_offset),
                                     remote_address + dst_offset, length,
-                                    channel_id, flags, comp);
+                                    uct_channel_id, flags, comp);
 }
 
 
@@ -274,23 +285,27 @@ UCS_F_DEVICE ucs_status_t ucp_device_counter_inc(
         unsigned mem_list_index, size_t offset, unsigned channel_id,
         uint64_t flags, ucp_device_request_t *req)
 {
+    unsigned uct_channel_id;
+    unsigned lane;
     uint64_t remote_address;
-    const uct_device_mem_element_t *uct_elem;
+    const uct_device_mem_elem_t *uct_elem;
     uct_device_completion_t *comp;
     uct_device_ep_t *device_ep;
     ucs_status_t status;
 
+    UCP_DEVICE_GET_LANE(mem_list_h, channel_id, lane, uct_channel_id);
+
     status = ucp_device_prepare_send_remote(mem_list_h, mem_list_index,
-                                            remote_address, req, device_ep,
-                                            uct_elem, comp);
+                                            remote_address, lane, req,
+                                            device_ep, uct_elem, comp);
     if (status != UCS_OK) {
         return status;
     }
 
     return UCP_DEVICE_SEND_BLOCKING(level, uct_device_ep_atomic_add, device_ep,
                                     req, uct_elem, inc_value,
-                                    remote_address + offset, channel_id, flags,
-                                    comp);
+                                    remote_address + offset, uct_channel_id,
+                                    flags, comp);
 }
 
 
@@ -311,7 +326,6 @@ UCS_F_DEVICE ucs_status_t
 ucp_device_get_ptr(const ucp_device_remote_mem_list_h mem_list_h,
                    unsigned mem_list_index, void **addr_p)
 {
-    const size_t elem_size = sizeof(uct_device_remote_mem_list_elem_t);
     ucs_status_t status;
 
     status = ucp_device_check_params(mem_list_h, mem_list_index);
@@ -319,12 +333,9 @@ ucp_device_get_ptr(const ucp_device_remote_mem_list_h mem_list_h,
         return status;
     }
 
-    const auto mem_element = static_cast<uct_device_remote_mem_list_elem_t*>(
-            UCS_PTR_BYTE_OFFSET(mem_list_h->mem_elements,
-                                mem_list_index * elem_size));
+    const auto mem_element = UCP_DEVICE_GET_ELEM(mem_list_h, mem_list_index);
 
-    return uct_device_ep_get_ptr(mem_element->device_ep,
-                                 &mem_element->uct_mem_element,
+    return uct_device_ep_get_ptr(mem_element->tl[0].ep, &mem_element->tl[0].uct,
                                  mem_element->addr, addr_p);
 }
 
diff --git a/src/ucp/api/device/ucp_device_types.h b/src/ucp/api/device/ucp_device_types.h
index 7847f314fdd..121607345cb 100644
--- a/src/ucp/api/device/ucp_device_types.h
+++ b/src/ucp/api/device/ucp_device_types.h
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -30,17 +30,22 @@ typedef struct ucp_device_remote_mem_list {
      * Structure version. Allow runtime ABI compatibility checks between host
      * and device code.
      */
-    uint16_t                          version;
+    uint16_t                     version;
+
+    /**
+     * Number of lanes in each memory descriptor.
+     */
+    uint16_t                     num_lanes;
 
     /**
      * Number of entries in the memory descriptors array @a elems.
      */
-    uint32_t                          length;
+    uint32_t                     length;
 
     /**
      * UCT memory element objects are allocated contiguously.
      */
-    uct_device_remote_mem_list_elem_t mem_elements[0];
+    uct_device_remote_mem_elem_t mem_elements[0];
 } ucp_device_remote_mem_list_t;
 
 
@@ -59,14 +64,19 @@ typedef struct ucp_device_local_mem_list {
      * Structure version. Allow runtime ABI compatibility checks between host
      * and device code.
      */
-    uint16_t                         version;
+    uint16_t                    version;
+
+    /**
+     * Number of lanes in each memory descriptor.
+     */
+    uint16_t                    num_lanes;
 
     /**
      * Number of entries in the memory descriptors array @a elems.
      */
-    uint32_t                         length;
+    uint32_t                    length;
 
-    uct_device_local_mem_list_elem_t mem_elements[0];
+    uct_device_local_mem_elem_t mem_elements[0];
 } ucp_device_local_mem_list_t;
 
 #endif /* UCP_DEVICE_TYPES_H */
diff --git a/src/ucp/core/ucp_device.c b/src/ucp/core/ucp_device.c
index 0fedc394803..8efc601fe15 100644
--- a/src/ucp/core/ucp_device.c
+++ b/src/ucp/core/ucp_device.c
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -23,9 +23,6 @@
 #include "ucp_mm.inl"
 
 
-#define UCP_DEVICE_MEM_LIST_MAX_EPS 2
-
-
 typedef struct {
     uct_allocated_memory_t mem;
     uint32_t               mem_list_length;
@@ -43,6 +40,13 @@ static ucs_spinlock_t ucp_device_handle_hash_lock;
 /* Size of temporary allocation for local sys_dev detection */
 #define UCP_DEVICE_LOCAL_SYS_DEV_DETECT_SIZE 64
 
+enum {
+    UCP_DEVICE_TL_TYPE_FIRST,
+    UCP_DEVICE_TL_TYPE_LKEY = UCP_DEVICE_TL_TYPE_FIRST,
+    UCP_DEVICE_TL_TYPE_NOLKEY,
+    UCP_DEVICE_TL_TYPE_LAST,
+};
+
 
 void ucp_device_init(void)
 {
@@ -138,135 +142,47 @@ ucp_device_detect_local_sys_dev(ucp_context_h context,
     return UCS_OK;
 }
 
-static ucp_md_map_t
-ucp_device_detect_local_md_map(const ucp_context_h context,
-                               ucs_sys_device_t local_sys_dev)
+static void
+ucp_device_get_tl_bitmap(const ucp_worker_h worker,
+                         ucp_tl_bitmap_t tl_bitmap[UCP_DEVICE_TL_TYPE_LAST],
+                         ucs_sys_device_t local_sys_dev)
 {
-    ucp_md_map_t local_md_map = 0;
-    ucp_md_index_t md_index;
-
-    /* Build MD map from MDs that can access the local_sys_dev */
-    for (md_index = 0; md_index < context->num_mds; md_index++) {
-        ucp_sys_dev_map_t sys_dev_map = context->tl_mds[md_index].sys_dev_map;
+    const ucp_worker_iface_t *wiface;
+    ucp_rsc_index_t tl_id;
+    int tl_type;
 
-        if (sys_dev_map & UCS_BIT(local_sys_dev)) {
-            local_md_map |= UCS_BIT(md_index);
-        }
+    /** TODO: Maybe cache results */
+    for (tl_type = UCP_DEVICE_TL_TYPE_FIRST;
+         tl_type < UCP_DEVICE_TL_TYPE_LAST; tl_type++) {
+        UCS_STATIC_BITMAP_RESET_ALL(&tl_bitmap[tl_type]);
     }
 
-    ucs_trace("detected local_md_map=0x%" PRIx64 " for local_sys_dev=%u",
-              local_md_map, local_sys_dev);
-    return local_md_map;
-}
-
-static void ucp_device_mem_list_lane_lookup(
-        ucp_ep_h ep, ucp_ep_config_t *ep_config, ucs_sys_device_t local_sys_dev,
-        ucp_md_map_t local_md_map, ucs_sys_device_t remote_sys_dev,
-        ucp_md_map_t remote_md_map,
-        ucp_lane_index_t lanes[UCP_DEVICE_MEM_LIST_MAX_EPS])
-{
-    double best_bw[UCP_DEVICE_MEM_LIST_MAX_EPS] = {-1., -1.};
-    ucp_lane_index_t lane;
-    double bandwidth;
-    ucp_ep_config_key_lane_t *lane_key;
-    ucs_sys_device_t src_sys_dev;
-    ucp_md_index_t src_md_index;
-
-    lanes[0] = UCP_NULL_LANE;
-    lanes[1] = UCP_NULL_LANE;
+    UCS_STATIC_BITMAP_FOR_EACH_BIT(tl_id, &worker->context->tl_bitmap) {
+        wiface = ucp_worker_iface(worker, tl_id);
 
-    for (lane = 0; lane < ep_config->key.num_lanes; ++lane) {
-        if (!(ep_config->key.lanes[lane].lane_types &
-              UCS_BIT(UCP_LANE_TYPE_DEVICE))) {
+        if (!(wiface->attr.cap.flags & UCT_IFACE_FLAG_DEVICE_EP)) {
             continue;
         }
 
-        lane_key = &ep_config->key.lanes[lane];
-        /* Check lane remote sys dev only when remote memory is not host */
-        if ((remote_sys_dev != UCS_SYS_DEVICE_ID_UNKNOWN) &&
-            (remote_sys_dev != lane_key->dst_sys_dev)) {
-            ucs_trace("lane[%u] wrong destination sys_dev: dst_sys_dev=%u",
-                      lane, lane_key->dst_sys_dev);
+        if ((wiface->attr.ctl_device != UCS_SYS_DEVICE_ID_UNKNOWN) &&
+            (wiface->attr.ctl_device != local_sys_dev)) {
             continue;
         }
 
-        if (!(remote_md_map & UCS_BIT(lane_key->dst_md_index))) {
-            ucs_trace("lane[%u] missing remote md: dst_md_index=%u", lane,
-                      lane_key->dst_md_index);
-            continue;
-        }
-
-        src_sys_dev = ucp_ep_get_tl_rsc(ep, lane)->sys_device;
-        if (src_sys_dev != local_sys_dev) {
-            ucs_trace("lane[%u] wrong source sys_dev: src_sys_dev=%u", lane,
-                      src_sys_dev);
-            continue;
-        }
-
-        src_md_index = ucp_ep_md_index(ep, lane);
-        if (!(local_md_map & UCS_BIT(src_md_index))) {
-            ucs_trace("lane[%u] missing local md: src_md_index=%u", lane,
-                      src_md_index);
-            continue;
-        }
-
-        bandwidth = ucp_worker_iface_bandwidth(ep->worker,
-                                               ucp_ep_get_rsc_index(ep, lane));
-        ucs_trace("checking lane[%u] src_md_index=%u dst_md_index=%u "
-                  "src_sys_dev=%u dst_sys_dev=%u bandwidth=%lfMB/s",
-                  lane, src_md_index, lane_key->dst_md_index, src_sys_dev,
-                  lane_key->dst_sys_dev, bandwidth / UCS_MBYTE);
-
-        UCS_STATIC_ASSERT(UCP_DEVICE_MEM_LIST_MAX_EPS == 2);
-        if (bandwidth > best_bw[0]) {
-            best_bw[1] = best_bw[0];
-            lanes[1]   = lanes[0];
-            best_bw[0] = bandwidth;
-            lanes[0]   = lane;
-        } else if (bandwidth > best_bw[1]) {
-            best_bw[1] = bandwidth;
-            lanes[1]   = lane;
+        if (wiface->attr.cap.flags & UCT_IFACE_FLAG_DEVICE_LKEY) {
+            tl_type = UCP_DEVICE_TL_TYPE_LKEY;
         } else {
-            continue;
-        }
-
-        ucs_trace("best lanes: lane[%u]=%lfMB/s lane[%u]=%lfMB/s", lanes[0],
-                  best_bw[0] / UCS_MBYTE, lanes[1], best_bw[1] / UCS_MBYTE);
-    }
-}
-
-
-static ucp_worker_iface_t *
-ucp_device_get_worker_iface_by_device_id(ucp_worker_h worker,
-                                         ucs_sys_device_t device_mem_id)
-{
-    const ucp_tl_resource_desc_t *resource;
-    const uct_md_attr_v2_t *md_attr;
-    ucp_md_index_t md_index;
-    unsigned i;
-
-    /** TODO: Maybe cache results */
-    for (i = 0; i < worker->num_ifaces; i++) {
-        resource = &worker->context->tl_rscs[worker->ifaces[i]->rsc_index];
-        md_index = resource->md_index;
-        md_attr  = &worker->context->tl_mds[md_index].attr;
-        if ((md_attr->flags & UCT_MD_FLAG_NEED_MEMH) &&
-            (resource->tl_rsc.sys_device == device_mem_id)) {
-            return worker->ifaces[i];
+            tl_type = UCP_DEVICE_TL_TYPE_NOLKEY;
         }
+        UCS_STATIC_BITMAP_SET(&tl_bitmap[tl_type], tl_id);
     }
-
-    return NULL;
 }
 
 static ucs_status_t ucp_device_local_mem_list_element_pack(
         const ucp_worker_h worker, const ucp_worker_iface_t *wiface,
         const ucp_device_mem_list_elem_t *element,
-        const ucs_memory_type_t mem_type,
-        uct_device_local_mem_list_elem_t *mem_element)
+        const ucs_memory_type_t mem_type, uct_device_mem_elem_t *mem_element)
 {
-    void *local_addr = UCS_PARAM_VALUE(UCP_DEVICE_MEM_LIST_ELEM_FIELD, element,
-                                       local_addr, LOCAL_ADDR, NULL);
     ucp_tl_resource_desc_t *resource;
     ucp_md_index_t md_index;
     ucp_mem_h memh;
@@ -274,7 +190,6 @@ static ucs_status_t ucp_device_local_mem_list_element_pack(
     ucp_tl_md_t *ucp_md;
     ucs_status_t status;
 
-    mem_element->addr = local_addr;
     if (wiface == NULL) {
         return UCS_OK;
     }
@@ -291,7 +206,7 @@ static ucs_status_t ucp_device_local_mem_list_element_pack(
     }
 
     status = uct_md_mem_elem_pack(ucp_md->md, uct_memh, UCT_INVALID_RKEY,
-                                  &mem_element->uct_mem_element);
+                                  mem_element);
     if (status != UCS_OK) {
         ucs_error("failed to pack local mem element for memh=%p", memh);
     }
@@ -328,47 +243,62 @@ static ucs_status_t ucp_device_local_mem_list_create_handle(
 {
     const ucp_worker_h worker   = UCS_PARAM_VALUE(
             UCP_DEVICE_MEM_LIST_PARAMS_FIELD, params, worker, WORKER, NULL);
-    const size_t uct_elem_size  = sizeof(uct_device_local_mem_list_elem_t);
-    size_t handle_size          = 0;
+    size_t uct_elem_size;
+    size_t handle_size;
+    int tl_type = UCP_DEVICE_TL_TYPE_LKEY;
+    ucp_tl_bitmap_t tl_bitmap[UCP_DEVICE_TL_TYPE_LAST] = {};
     const ucp_device_mem_list_elem_t *ucp_element;
     const ucp_worker_iface_t *wiface;
     ucp_device_local_mem_list_t *handle;
-    uct_device_local_mem_list_elem_t *uct_element;
-    size_t i;
+    uct_device_local_mem_elem_t *uct_element;
+    uct_device_mem_elem_t *tl_element;
+    size_t i, num_lanes;
     ucs_status_t status;
+    ucp_rsc_index_t tl_id;
+    void *local_addr;
+
+    ucp_device_get_tl_bitmap(worker, tl_bitmap, local_sys_dev);
+    num_lanes = UCS_STATIC_BITMAP_POPCOUNT(tl_bitmap[tl_type]);
 
-    handle_size = (uct_elem_size * params->num_elements) + sizeof(*handle);
-    handle      = ucs_calloc(1, handle_size, "ucp_device_local_mem_list_t");
+    uct_elem_size = sizeof(uct_device_local_mem_elem_t) +
+                    (sizeof(uct_device_mem_elem_t) * num_lanes);
+    handle_size   = (uct_elem_size * params->num_elements) + sizeof(*handle);
+    handle        = ucs_calloc(1, handle_size, "ucp_device_local_mem_list_t");
     if (handle == NULL) {
         ucs_error("failed to allocate ucp_device_local_mem_list_t");
         return UCS_ERR_NO_MEMORY;
     }
 
-    /* TODO: To support multi lanes we need to pack all memhs of ifaces that require memhs */
-    wiface = ucp_device_get_worker_iface_by_device_id(worker, local_sys_dev);
-    if (wiface == NULL) {
-        ucs_debug("no worker iface found for device_id=%u", local_sys_dev);
-    }
-
     /* Populate element specific parameters */
     ucp_element = params->elements;
-    uct_element = UCS_PTR_BYTE_OFFSET(handle, sizeof(*handle));
+    uct_element = UCS_PTR_TYPE_OFFSET(handle, *handle);
     for (i = 0; i < params->num_elements; i++) {
-        status = ucp_device_local_mem_list_element_pack(worker, wiface,
-                                                        ucp_element, mem_type,
-                                                        uct_element);
-        if (status != UCS_OK) {
-            ucs_error("failed to pack local mem list element for element=%zu",
-                      i);
-            goto out;
-        }
+        local_addr        = UCS_PARAM_VALUE(UCP_DEVICE_MEM_LIST_ELEM_FIELD,
+                                            ucp_element, local_addr, LOCAL_ADDR, NULL);
+        uct_element->addr = local_addr;
+        tl_element        = uct_element->tl;
+        UCS_STATIC_BITMAP_FOR_EACH_BIT(tl_id, &tl_bitmap[tl_type]) {
+            wiface = ucp_worker_iface(worker, tl_id);
+            status = ucp_device_local_mem_list_element_pack(worker, wiface,
+                                                            ucp_element,
+                                                            mem_type,
+                                                            tl_element);
+            if (status != UCS_OK) {
+                ucs_error("failed to pack local mem list element for "
+                          "element=%zu",
+                          i);
+                goto out;
+            }
 
+            tl_element = UCS_PTR_TYPE_OFFSET(tl_element, *tl_element);
+        }
+        uct_element = (void*)tl_element;
         ucp_element = UCS_PTR_BYTE_OFFSET(ucp_element, params->element_size);
-        uct_element = UCS_PTR_BYTE_OFFSET(uct_element, uct_elem_size);
     }
 
     handle->version = UCP_DEVICE_MEM_LIST_VERSION_V1;
     handle->length  = params->num_elements;
+    handle->num_lanes = num_lanes;
     status          = ucp_device_mem_list_export_handle(
             worker, handle, handle_size, mem_type, local_sys_dev, mem,
             "ucp_device_local_mem_list_handle_t");
@@ -466,22 +396,46 @@ ucp_device_local_mem_list_create(const ucp_device_mem_list_params_t *params,
     return status;
 }
 
+static ucp_lane_index_t ucp_device_ep_find_lane(const ucp_ep_h ep, ucp_rsc_index_t tl_id)
+{
+    ucp_lane_index_t lane;
+
+    for (lane = 0; lane < ucp_ep_num_lanes(ep); ++lane) {
+        if (ucp_ep_get_rsc_index(ep, lane) == tl_id) {
+            return lane;
+        }
+    }
+
+    return UCP_NULL_LANE;
+}
+
+static int
+ucp_device_ep_check_lanes(const ucp_ep_h ep, ucp_tl_bitmap_t *tl_bitmap)
+{
+    ucp_lane_index_t lane;
+    ucp_rsc_index_t tl_id;
+
+    if (UCS_STATIC_BITMAP_POPCOUNT(*tl_bitmap) == 0) {
+        return 0;
+    }
+
+    UCS_STATIC_BITMAP_FOR_EACH_BIT(tl_id, tl_bitmap) {
+        lane = ucp_device_ep_find_lane(ep, tl_id);
+        if (lane == UCP_NULL_LANE) {
+            return 0;
+        }
+    }
+
+    return 1;
+}
+
 static ucs_status_t ucp_device_remote_mem_list_element_pack(
-        const ucp_device_mem_list_elem_t *element,
-        const ucs_sys_device_t local_sys_dev, const ucs_memory_type_t mem_type,
-        uct_device_remote_mem_list_elem_t *mem_element)
+        const ucp_device_mem_list_elem_t *element, ucp_rsc_index_t tl_id,
+        uct_device_remote_tl_elem_t *mem_element)
 {
     const ucp_ep_h ep     = element->ep;
     const ucp_rkey_h rkey = element->rkey;
-    const ucp_md_map_t local_md_map =
-            ucp_device_detect_local_md_map(ep->worker->context, local_sys_dev);
-    const ucp_worker_cfg_index_t rkey_cfg_index = element->rkey->cfg_index;
-    const ucp_rkey_config_t rkey_config =
-            ucs_array_elem(&ep->worker->rkey_config, rkey_cfg_index);
-    const ucs_sys_device_t remote_sys_dev = rkey_config.key.sys_dev;
-    const ucp_md_map_t remote_md_map      = rkey_config.key.md_map;
-    ucp_ep_config_t *ep_config            = ucp_ep_config(ep);
-    ucp_lane_index_t lanes[UCP_DEVICE_MEM_LIST_MAX_EPS];
+    ucp_ep_config_t *ep_config = ucp_ep_config(ep);
     uint8_t rkey_index;
     uct_rkey_t uct_rkey;
     uct_ep_h uct_ep;
@@ -489,10 +443,7 @@ static ucs_status_t ucp_device_remote_mem_list_element_pack(
     ucs_status_t status;
     ucp_lane_index_t lane;
 
-    ucp_device_mem_list_lane_lookup(ep, ep_config, local_sys_dev, local_md_map,
-                                    remote_sys_dev, remote_md_map, lanes);
-    lane = lanes[0];
-
+    lane = ucp_device_ep_find_lane(ep, tl_id);
     if (lane == UCP_NULL_LANE) {
         ucs_error("no lane found for ep=%p", ep);
         return UCS_ERR_NO_DEVICE;
@@ -510,10 +461,9 @@ static ucs_status_t ucp_device_remote_mem_list_element_pack(
     uct_rkey   = ucp_rkey_get_tl_rkey(rkey, rkey_index);
     ucs_assert(uct_rkey != UCT_INVALID_RKEY);
 
-    mem_element->device_ep = device_ep;
-    mem_element->addr      = element->remote_addr;
-    status = uct_md_mem_elem_pack(ucp_ep_md(ep, lane), NULL, uct_rkey,
-                                  &mem_element->uct_mem_element);
+    mem_element->ep = device_ep;
+    status          = uct_md_mem_elem_pack(ucp_ep_md(ep, lane), NULL, uct_rkey,
+                                           &mem_element->uct);
     if (status != UCS_OK) {
         ucs_error("failed to pack uct memory element for lane=%u", lane);
     }
@@ -555,19 +505,49 @@ static ucp_ep_h ucp_device_remote_mem_list_get_first_ep(
     return NULL;
 }
 
+static ucs_status_t
+ucp_device_remote_mem_list_fill(const ucp_device_mem_list_elem_t *ucp_element,
+                                ucp_tl_bitmap_t *tl_bitmap, size_t num_lanes,
+                                uct_device_remote_mem_elem_t *uct_element)
+{
+    uct_device_remote_tl_elem_t *tl_element;
+    ucp_rsc_index_t tl_id;
+    ucs_status_t status;
+    size_t i;
+
+    uct_element->addr = ucp_element->remote_addr;
+    tl_element        = uct_element->tl;
+    for (i = 0; i < num_lanes;) {
+        UCS_STATIC_BITMAP_FOR_EACH_BIT(tl_id, tl_bitmap) {
+            status = ucp_device_remote_mem_list_element_pack(ucp_element, tl_id,
+                                                             tl_element);
+            if (status != UCS_OK) {
+                return status;
+            }
+
+            tl_element = UCS_PTR_TYPE_OFFSET(tl_element, *tl_element);
+            i++;
+        }
+    }
+
+    return UCS_OK;
+}
+
 static ucs_status_t ucp_device_remote_mem_list_create_handle(
         const ucp_device_mem_list_params_t *params, ucs_memory_type_t mem_type,
         uct_allocated_memory_t *mem)
 {
     const ucp_ep_h ep = ucp_device_remote_mem_list_get_first_ep(params);
-    const size_t uct_elem_size = sizeof(uct_device_remote_mem_list_elem_t);
+    size_t uct_elem_size;
     size_t handle_size         = 0;
+    ucp_tl_bitmap_t tl_bitmap[UCP_DEVICE_TL_TYPE_LAST] = {};
     const ucp_device_mem_list_elem_t *ucp_element;
     ucp_device_remote_mem_list_t *handle;
-    uct_device_remote_mem_list_elem_t *uct_element;
+    uct_device_remote_mem_elem_t *uct_element;
     ucs_sys_device_t local_sys_dev;
-    size_t i;
+    size_t i, num_lanes;
     ucs_status_t status;
+    int tl_type;
 
     if (ep == NULL) {
         ucs_error("no ep found in remote mem list");
@@ -580,34 +560,65 @@ static ucs_status_t ucp_device_remote_mem_list_create_handle(
         return status;
     }
 
-    handle_size = sizeof(*handle) + (uct_elem_size * params->num_elements);
+    ucp_device_get_tl_bitmap(ep->worker, tl_bitmap, local_sys_dev);
+
+    /* handle->num_lanes is the least common multiple of both lane types, so:
+     * - each lane is replicated num_lanes / popcount(tl_bitmap) times
+     * - channel_id % num_lanes maps to the correct lane, regardless of lane type
+     */
+    num_lanes = UCS_STATIC_BITMAP_POPCOUNT(tl_bitmap[UCP_DEVICE_TL_TYPE_LKEY]);
+    if (!num_lanes) {
+        if (!UCS_STATIC_BITMAP_POPCOUNT(tl_bitmap[UCP_DEVICE_TL_TYPE_NOLKEY])) {
+            ucs_error("failed to pack uct memory element for first element");
+            return UCS_ERR_INVALID_PARAM;
+        }
+
+        ucs_assert(UCS_STATIC_BITMAP_POPCOUNT(
+                           tl_bitmap[UCP_DEVICE_TL_TYPE_NOLKEY]) == 1);
+        num_lanes = 1;
+    }
+
+    ucp_element   = params->elements;
+    uct_elem_size = sizeof(uct_device_remote_mem_elem_t) +
+                    (sizeof(uct_device_remote_tl_elem_t) * num_lanes);
+    handle_size   = sizeof(*handle) + (params->num_elements * uct_elem_size);
     handle      = ucs_calloc(1, handle_size, "ucp_device_remote_mem_list_t");
     if (handle == NULL) {
         ucs_error("failed to allocate ucp_device_remote_mem_list_t");
         return UCS_ERR_NO_MEMORY;
     }
 
-    ucp_element = params->elements;
-    uct_element = UCS_PTR_BYTE_OFFSET(handle, sizeof(*handle));
+    uct_element = UCS_PTR_TYPE_OFFSET(handle, *handle);
     for (i = 0; i < params->num_elements; i++) {
         if (!UCP_DEVICE_MEM_ELEMENT_IS_GAP(ucp_element)) {
-            status = ucp_device_remote_mem_list_element_pack(ucp_element,
-                                                             local_sys_dev,
-                                                             mem_type,
-                                                             uct_element);
+            for (tl_type = UCP_DEVICE_TL_TYPE_FIRST;
+                 tl_type < UCP_DEVICE_TL_TYPE_LAST; tl_type++) {
+                if (ucp_device_ep_check_lanes(ucp_element->ep,
+                                              &tl_bitmap[tl_type])) {
+                    break;
+                }
+            }
+
+            if (tl_type == UCP_DEVICE_TL_TYPE_LAST) {
+                ucs_error("lane not found for element %zd", i);
+                status =  UCS_ERR_INVALID_PARAM;
+                goto out;
+            }
+
+            status = ucp_device_remote_mem_list_fill(ucp_element,
+                                                     &tl_bitmap[tl_type],
+                                                     num_lanes, uct_element);
             if (status != UCS_OK) {
-                ucs_error("failed to pack uct memory element for element=%zu",
-                          i);
                 goto out;
             }
         }
-
-        ucp_element = UCS_PTR_BYTE_OFFSET(ucp_element, params->element_size);
         uct_element = UCS_PTR_BYTE_OFFSET(uct_element, uct_elem_size);
+        ucp_element = UCS_PTR_BYTE_OFFSET(ucp_element, params->element_size);
     }
 
     handle->version = UCP_DEVICE_MEM_LIST_VERSION_V1;
     handle->length  = params->num_elements;
+    handle->num_lanes = num_lanes;
     status          = ucp_device_mem_list_export_handle(
             ep->worker, handle, handle_size, mem_type, local_sys_dev, mem,
             "ucp_device_remote_mem_list_handle_t");
@@ -716,7 +727,6 @@ ucp_device_remote_mem_list_create(const ucp_device_mem_list_params_t *params,
     return status;
 }
 
-
 uint32_t ucp_device_get_mem_list_length(const void *mem_list_h)
 {
     khiter_t iter;
diff --git a/src/ucp/wireup/address.c b/src/ucp/wireup/address.c
index 71b0ea1d0fc..70406b0766d 100644
--- a/src/ucp/wireup/address.c
+++ b/src/ucp/wireup/address.c
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2001-2016. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2001-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -315,6 +315,7 @@ ucp_address_get_device(ucp_context_h context, ucp_rsc_index_t rsc_index,
 
     dev = &devices[(*num_devices_p)++];
     memset(dev, 0, sizeof(*dev));
+    dev->num_paths = 1;
 out:
     return dev;
 }
@@ -453,9 +454,11 @@ ucp_address_gather_devices(ucp_worker_h worker, const ucp_ep_config_key_t *key,
             return UCS_ERR_UNSUPPORTED;
         }
 
-        dev->rsc_index  = rsc_index;
+        dev->rsc_index = rsc_index;
         UCS_STATIC_BITMAP_SET(&dev->tl_bitmap, rsc_index);
-        dev->num_paths  = ucs_min(max_num_paths, iface_attr->dev_num_paths);
+        if (!(iface_attr->cap.flags & UCT_IFACE_FLAG_DEVICE_EP)) {
+            dev->num_paths = ucs_min(max_num_paths, iface_attr->dev_num_paths);
+        }
     }
 
     *devices_p     = devices;
diff --git a/src/ucp/wireup/select.c b/src/ucp/wireup/select.c
index b6b1572aaa5..181c93fa79b 100644
--- a/src/ucp/wireup/select.c
+++ b/src/ucp/wireup/select.c
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2001-2016. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2001-2026. ALL RIGHTS RESERVED.
  * Copyright (C) Los Alamos National Security, LLC. 2019 ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
@@ -1821,7 +1821,13 @@ static int ucp_wireup_add_bw_lanes_pairwise(
 
         /* Account for possible path override */
         local_num_paths  = iface_attr->dev_num_paths;
-        remote_num_paths = ae->dev_num_paths;
+
+        if (bw_info->criteria.lane_type == UCP_LANE_TYPE_DEVICE) {
+            remote_num_paths = 1;
+        } else {
+            remote_num_paths = ae->dev_num_paths;
+        }
+
         if (allow_extra_path &&
             ((skip_dev_index != UCP_NULL_RESOURCE) && /* clang sanitizer */
              (skip_dev_index == dev_index))) {
diff --git a/src/uct/api/device/uct_device_impl.h b/src/uct/api/device/uct_device_impl.h
index 255a9d45632..310827a243b 100644
--- a/src/uct/api/device/uct_device_impl.h
+++ b/src/uct/api/device/uct_device_impl.h
@@ -69,12 +69,11 @@ union uct_device_completion {
  * @return Error code as defined by @ref ucs_status_t
  */
 template<ucs_device_level_t level>
-UCS_F_DEVICE ucs_status_t
-uct_device_ep_put(uct_device_ep_h device_ep,
-                  const uct_device_local_mem_list_elem_t *src_uct_elem,
-                  const uct_device_mem_element_t *mem_elem, const void *address,
-                  uint64_t remote_address, size_t length, unsigned channel_id,
-                  uint64_t flags, uct_device_completion_t *comp)
+UCS_F_DEVICE ucs_status_t uct_device_ep_put(
+        uct_device_ep_h device_ep, const uct_device_mem_elem_t *src_uct_elem,
+        const uct_device_mem_elem_t *mem_elem, const void *address,
+        uint64_t remote_address, size_t length, unsigned channel_id,
+        uint64_t flags, uct_device_completion_t *comp)
 {
 #if UCT_RC_MLX5_GDA_SUPPORTED
     if (device_ep->uct_tl_id == UCT_DEVICE_TL_RC_MLX5_GDA) {
@@ -123,7 +122,7 @@ uct_device_ep_put(uct_device_ep_h device_ep,
  */
 template<ucs_device_level_t level>
 UCS_F_DEVICE ucs_status_t uct_device_ep_atomic_add(
-        uct_device_ep_h device_ep, const uct_device_mem_element_t *mem_elem,
+        uct_device_ep_h device_ep, const uct_device_mem_elem_t *mem_elem,
         uint64_t inc_value, uint64_t remote_address, unsigned channel_id,
         uint64_t flags, uct_device_completion_t *comp)
 {
@@ -160,7 +159,7 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_atomic_add(
  * @return Error code as defined by @ref ucs_status_t
  */
 UCS_F_DEVICE ucs_status_t uct_device_ep_get_ptr(
-        uct_device_ep_h device_ep, const uct_device_mem_element_t *mem_elem,
+        uct_device_ep_h device_ep, const uct_device_mem_elem_t *mem_elem,
         uint64_t address, void **addr_p)
 {
 #if UCT_CUDA_IPC_SUPPORTED
diff --git a/src/uct/api/device/uct_device_types.h b/src/uct/api/device/uct_device_types.h
index c2104bd07fc..ec10a2370b4 100644
--- a/src/uct/api/device/uct_device_types.h
+++ b/src/uct/api/device/uct_device_types.h
@@ -72,22 +72,26 @@ typedef union uct_device_completion uct_device_completion_t;
 
 
 /* Union structure of all device memory elements types */
-union uct_device_mem_element {
+union uct_device_mem_elem {
     uct_ib_md_device_mem_element_t       ib_md_mem_element;
     uct_cuda_ipc_md_device_mem_element_t cuda_ipc_md_mem_element;
 };
 
 
-struct uct_device_local_mem_list_elem {
-    void                     *addr;
-    uct_device_mem_element_t uct_mem_element;
+struct uct_device_local_mem_elem {
+    void                  *addr;
+    uct_device_mem_elem_t tl[0];
 };
 
 
-struct uct_device_remote_mem_list_elem {
-    uct_device_ep_h          device_ep;
-    uint64_t                 addr;
-    uct_device_mem_element_t uct_mem_element;
+struct uct_device_remote_tl_elem {
+    uct_device_ep_h       ep;
+    uct_device_mem_elem_t uct;
+};
+
+struct uct_device_remote_mem_elem {
+    uint64_t                    addr;
+    uct_device_remote_tl_elem_t tl[0];
 };
 
 #endif
diff --git a/src/uct/api/uct.h b/src/uct/api/uct.h
index 9a18e3ea916..082f30d72bb 100644
--- a/src/uct/api/uct.h
+++ b/src/uct/api/uct.h
@@ -438,6 +438,7 @@ typedef enum uct_atomic_op {
         /* Interface capability */
 #define UCT_IFACE_FLAG_INTER_NODE      UCS_BIT(54) /**< Interface is inter-node capable */
 #define UCT_IFACE_FLAG_DEVICE_EP       UCS_BIT(55) /**< Interface supports device endpoint */
+#define UCT_IFACE_FLAG_DEVICE_LKEY     UCS_BIT(56) /**< Interface requires lkey for device operations */
 /**
  * @}
  */
@@ -1187,6 +1188,9 @@ struct uct_iface_attr {
                                                 achieve higher total bandwidth
                                                 compared to using only a single
                                                 endpoint. */
+
+    ucs_sys_device_t ctl_device;           /**< System device controlling this iface,
+                                                if it's not CPU. */
 };
 
 
diff --git a/src/uct/api/uct_def.h b/src/uct/api/uct_def.h
index 688d08689fc..0895ce7469d 100644
--- a/src/uct/api/uct_def.h
+++ b/src/uct/api/uct_def.h
@@ -111,9 +111,10 @@ typedef uint64_t                     uct_tag_t;  /* tag type - 64 bit */
 typedef int                          uct_worker_cb_id_t;
 typedef void*                        uct_conn_request_h;
 typedef struct uct_device_ep         *uct_device_ep_h;
-typedef union uct_device_mem_element uct_device_mem_element_t;
-typedef struct uct_device_local_mem_list_elem uct_device_local_mem_list_elem_t;
-typedef struct uct_device_remote_mem_list_elem uct_device_remote_mem_list_elem_t;
+typedef union uct_device_mem_elem uct_device_mem_elem_t;
+typedef struct uct_device_local_mem_elem uct_device_local_mem_elem_t;
+typedef struct uct_device_remote_tl_elem uct_device_remote_tl_elem_t;
+typedef struct uct_device_remote_mem_elem uct_device_remote_mem_elem_t;
 
 /**
  * @}
diff --git a/src/uct/api/v2/uct_v2.h b/src/uct/api/v2/uct_v2.h
index 19b64d1a52b..e4be474f0ef 100644
--- a/src/uct/api/v2/uct_v2.h
+++ b/src/uct/api/v2/uct_v2.h
@@ -1364,7 +1364,7 @@ ucs_status_t uct_rkey_unpack_v2(uct_component_h component,
  * @return UCS_OK on success or error code in case of failure.
  */
 ucs_status_t uct_md_mem_elem_pack(uct_md_h md, uct_mem_h memh, uct_rkey_t rkey,
-                                  uct_device_mem_element_t *mem_elem);
+                                  uct_device_mem_elem_t *mem_elem);
 
 END_C_DECLS
 
diff --git a/src/uct/base/uct_md.c b/src/uct/base/uct_md.c
index 8e8d5a71ecf..3df7ce1c0ab 100644
--- a/src/uct/base/uct_md.c
+++ b/src/uct/base/uct_md.c
@@ -660,7 +660,7 @@ ucs_status_t uct_md_dummy_mem_dereg(uct_md_h uct_md,
 }
 
 ucs_status_t uct_md_mem_elem_pack(uct_md_h md, uct_mem_h memh, uct_rkey_t rkey,
-                                  uct_device_mem_element_t *mem_elem)
+                                  uct_device_mem_elem_t *mem_elem)
 {
     return md->ops->mem_elem_pack(md, memh, rkey, mem_elem);
 }
diff --git a/src/uct/base/uct_md.h b/src/uct/base/uct_md.h
index 8219ac652d7..31909de3c22 100644
--- a/src/uct/base/uct_md.h
+++ b/src/uct/base/uct_md.h
@@ -128,7 +128,7 @@ typedef ucs_status_t (*uct_md_detect_memory_type_func_t)(uct_md_h md,
 
 typedef ucs_status_t (*uct_md_mem_elem_pack_func_t)(
         uct_md_h md, uct_mem_h memh, uct_rkey_t rkey,
-        uct_device_mem_element_t *mem_elem_p);
+        uct_device_mem_elem_t *mem_elem_p);
 
 /**
  * Memory domain operations
diff --git a/src/uct/cuda/cuda_ipc/cuda_ipc.cuh b/src/uct/cuda/cuda_ipc/cuda_ipc.cuh
index 47b41f9020d..421ec2bfe7e 100644
--- a/src/uct/cuda/cuda_ipc/cuda_ipc.cuh
+++ b/src/uct/cuda/cuda_ipc/cuda_ipc.cuh
@@ -292,7 +292,7 @@ void uct_cuda_ipc_copy_level<UCS_DEVICE_LEVEL_GRID>(void *dst, const void *src,
 
 template<ucs_device_level_t level = UCS_DEVICE_LEVEL_BLOCK>
 UCS_F_DEVICE ucs_status_t uct_cuda_ipc_ep_put(
-        uct_device_ep_h device_ep, const uct_device_mem_element_t *mem_elem,
+        uct_device_ep_h device_ep, const uct_device_mem_elem_t *mem_elem,
         const void *address, uint64_t remote_address, size_t length,
         uint64_t flags, uct_device_completion_t *comp)
 {
@@ -310,7 +310,7 @@ UCS_F_DEVICE ucs_status_t uct_cuda_ipc_ep_put(
 template<ucs_device_level_t level = UCS_DEVICE_LEVEL_BLOCK>
 UCS_F_DEVICE ucs_status_t
 uct_cuda_ipc_ep_atomic_add(uct_device_ep_h device_ep,
-                           const uct_device_mem_element_t *mem_elem,
+                           const uct_device_mem_elem_t *mem_elem,
                            uint64_t inc_value, uint64_t remote_address,
                            uint64_t flags, uct_device_completion_t *comp)
 {
@@ -332,7 +332,7 @@ uct_cuda_ipc_ep_atomic_add(uct_device_ep_h device_ep,
 }
 
 UCS_F_DEVICE ucs_status_t uct_cuda_ipc_ep_get_ptr(
-        uct_device_ep_h device_ep, const uct_device_mem_element_t *mem_elem,
+        uct_device_ep_h device_ep, const uct_device_mem_elem_t *mem_elem,
         uint64_t remote_address, void **addr_p)
 {
     auto cuda_ipc_mem_element =
diff --git a/src/uct/cuda/cuda_ipc/cuda_ipc_iface.c b/src/uct/cuda/cuda_ipc/cuda_ipc_iface.c
index 954f663a060..ef41ec5263c 100644
--- a/src/uct/cuda/cuda_ipc/cuda_ipc_iface.c
+++ b/src/uct/cuda/cuda_ipc/cuda_ipc_iface.c
@@ -268,6 +268,9 @@ static ucs_status_t uct_cuda_ipc_iface_query(uct_iface_h tl_iface,
                                           UCT_IFACE_FLAG_GET_ZCOPY        |
                                           UCT_IFACE_FLAG_PUT_ZCOPY        |
                                           UCT_IFACE_FLAG_DEVICE_EP;
+
+    iface_attr->ctl_device              = UCS_SYS_DEVICE_ID_UNKNOWN;
+
     if (md->enable_mnnvl) {
         iface_attr->cap.flags |= UCT_IFACE_FLAG_INTER_NODE;
     }
diff --git a/src/uct/cuda/cuda_ipc/cuda_ipc_md.c b/src/uct/cuda/cuda_ipc/cuda_ipc_md.c
index 1502740a8f5..3f6173def7d 100644
--- a/src/uct/cuda/cuda_ipc/cuda_ipc_md.c
+++ b/src/uct/cuda/cuda_ipc/cuda_ipc_md.c
@@ -561,7 +561,7 @@ static void uct_cuda_ipc_md_close(uct_md_h md)
 
 static ucs_status_t
 uct_cuda_ipc_md_mem_elem_pack(uct_md_h md, uct_mem_h memh, uct_rkey_t rkey,
-                              uct_device_mem_element_t *mem_elem_p)
+                              uct_device_mem_elem_t *mem_elem_p)
 {
     uct_cuda_ipc_unpacked_rkey_t *key = (uct_cuda_ipc_unpacked_rkey_t*)rkey;
     uct_cuda_ipc_md_device_mem_element_t *cuda_ipc_md_mem_element =
diff --git a/src/uct/ib/base/ib_md.c b/src/uct/ib/base/ib_md.c
index 405dffb60e2..262c84b24ca 100644
--- a/src/uct/ib/base/ib_md.c
+++ b/src/uct/ib/base/ib_md.c
@@ -120,10 +120,10 @@ ucs_config_field_t uct_ib_md_config_table[] = {
      "Use GPU Direct RDMA for HCA to access GPU pages directly\n",
      ucs_offsetof(uct_ib_md_config_t, enable_gpudirect_rdma), UCS_CONFIG_TYPE_TERNARY},
 
-    {"GDA_MAX_HCA_PER_GPU", "1",
+    {"GDA_MAX_HCA_PER_GPU", "auto",
      "Max number of HCA devices to use for GDA per one GPU device.",
      ucs_offsetof(uct_ib_md_config_t, ext.gda_max_hca_per_gpu),
-     UCS_CONFIG_TYPE_UINT},
+     UCS_CONFIG_TYPE_ULUNITS},
 
     {"GDA_DMABUF_ENABLE", "try",
      "Enable DMA-BUF in GDA.",
diff --git a/src/uct/ib/base/ib_md.h b/src/uct/ib/base/ib_md.h
index 40c37a57ea3..ca851796b20 100644
--- a/src/uct/ib/base/ib_md.h
+++ b/src/uct/ib/base/ib_md.h
@@ -112,7 +112,7 @@ typedef struct uct_ib_md_ext_config {
     unsigned long            reg_retry_cnt; /**< Memory registration retry count */
     unsigned                 smkey_block_size; /**< Mkey indexes in a symmetric block */
     int                      direct_nic; /**< Direct NIC with GPU functionality */
-    unsigned                 gda_max_hca_per_gpu; /**< Threshold of IB per GPU */
+    unsigned long            gda_max_hca_per_gpu; /**< Threshold of IB per GPU */
     int                      gda_dmabuf_enable; /**< Enable DMA-BUF in GDA */
     /**< Retain and use an inactive CUDA primary context to query device capabilities */
     int                      gda_retain_inactive_ctx;
diff --git a/src/uct/ib/mlx5/dv/ib_mlx5dv_md.c b/src/uct/ib/mlx5/dv/ib_mlx5dv_md.c
index 93d3a6a05b6..4993c5a8c70 100644
--- a/src/uct/ib/mlx5/dv/ib_mlx5dv_md.c
+++ b/src/uct/ib/mlx5/dv/ib_mlx5dv_md.c
@@ -3162,7 +3162,7 @@ uct_ib_mlx5_devx_md_open(struct ibv_device *ibv_device,
 static ucs_status_t
 uct_ib_md_mlx5_devx_md_mem_elem_pack(uct_md_h md, uct_mem_h memh,
                                      uct_rkey_t rkey,
-                                     uct_device_mem_element_t *mem_elem_p)
+                                     uct_device_mem_elem_t *mem_elem_p)
 {
     uct_ib_md_device_mem_element_t *mem_elem = (uct_ib_md_device_mem_element_t*)
             mem_elem_p;
diff --git a/src/uct/ib/mlx5/gdaki/gdaki.c b/src/uct/ib/mlx5/gdaki/gdaki.c
index 43a37b84ce5..c243327b682 100644
--- a/src/uct/ib/mlx5/gdaki/gdaki.c
+++ b/src/uct/ib/mlx5/gdaki/gdaki.c
@@ -788,6 +788,7 @@ uct_rc_gdaki_iface_query(uct_iface_h tl_iface, uct_iface_attr_t *iface_attr)
     iface_attr->cap.flags = UCT_IFACE_FLAG_CONNECT_TO_EP |
                             UCT_IFACE_FLAG_INTER_NODE |
                             UCT_IFACE_FLAG_DEVICE_EP |
+                            UCT_IFACE_FLAG_DEVICE_LKEY |
                             UCT_IFACE_FLAG_ERRHANDLE_PEER_FAILURE;
 
     iface_attr->ep_addr_len    = sizeof(uct_ib_uint24_t) * iface->num_channels;
@@ -797,6 +798,8 @@ uct_rc_gdaki_iface_query(uct_iface_h tl_iface, uct_iface_attr_t *iface_attr)
     iface_attr->cap.put.min_zcopy = 0;
     iface_attr->cap.put.max_zcopy =
             uct_ib_iface_port_attr(&iface->super.super.super)->max_msg_sz;
+    iface_attr->ctl_device    = uct_cuda_get_sys_dev(iface->cuda_dev);
+    iface_attr->dev_num_paths = 1;
 
     return UCS_OK;
 }
@@ -1199,7 +1202,7 @@ static int uct_gdaki_dev_matrix_score(const void *pa, const void *pb, void *arg)
 uct_gdaki_dev_matrix_elem_t *
 uct_gdaki_dev_matrix_init(const uct_ib_md_t *ib_md, size_t *dmat_length_p)
 {
-    unsigned ib_per_cuda              = ib_md->config.gda_max_hca_per_gpu;
+    unsigned long ib_per_cuda         = ib_md->config.gda_max_hca_per_gpu;
     uct_gdaki_dev_matrix_elem_t *dmat = NULL;
     ucs_status_t status;
     int ibdev_index, cudadev_index, ibdev_count, cudadev_count;
@@ -1278,6 +1281,10 @@ uct_gdaki_dev_matrix_init(const uct_ib_md_t *ib_md, size_t *dmat_length_p)
 
     ucs_assert(cudadev_count < UCT_GDAKI_MAX_CUDA_PER_IB);
 
+    if (ib_per_cuda == UCS_ULUNITS_AUTO) {
+        ib_per_cuda = ibdev_count / cudadev_count;
+    }
+
     /* Map each CUDA device to the best suited IB devices */
     for (cudadev_index = 0; cudadev_index < cudadev_count; cudadev_index++) {
         status = UCT_CUDADRV_FUNC_LOG_ERR(
@@ -1413,7 +1420,6 @@ uct_gdaki_query_tl_devices(uct_md_h tl_md,
     uct_tl_device_resource_t *tl_devices;
     ucs_status_t status;
     CUdevice device;
-    ucs_sys_device_t dev;
     int i;
     uct_gdaki_dev_matrix_elem_t *ibdesc;
 
@@ -1460,14 +1466,12 @@ uct_gdaki_query_tl_devices(uct_md_h tl_md,
             goto err;
         }
 
-        dev = uct_cuda_get_sys_dev(device);
-
         snprintf(tl_devices[num_tl_devices].name,
                  sizeof(tl_devices[num_tl_devices].name), "%s%d-%s:%d",
                  UCT_DEVICE_CUDA_NAME, device, uct_ib_device_name(&ib_md->dev),
                  ib_md->dev.first_port);
         tl_devices[num_tl_devices].type       = UCT_DEVICE_TYPE_NET;
-        tl_devices[num_tl_devices].sys_device = dev;
+        tl_devices[num_tl_devices].sys_device = ib_md->dev.sys_dev;
         num_tl_devices++;
     }
 
diff --git a/src/uct/ib/mlx5/gdaki/gdaki.cuh b/src/uct/ib/mlx5/gdaki/gdaki.cuh
index dff63d87d3c..948431ca561 100644
--- a/src/uct/ib/mlx5/gdaki/gdaki.cuh
+++ b/src/uct/ib/mlx5/gdaki/gdaki.cuh
@@ -263,7 +263,7 @@ uct_rc_mlx5_gda_fc(const uct_rc_gdaki_dev_ep_t *ep, uint16_t wqe_idx)
 
 template<ucs_device_level_t level>
 UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_single(
-        uct_rc_gdaki_dev_ep_t *ep, const uct_device_mem_element_t *tl_mem_elem,
+        uct_rc_gdaki_dev_ep_t *ep, const uct_device_mem_elem_t *tl_mem_elem,
         const void *address, uint32_t lkey, uint64_t remote_address,
         uint32_t rkey, size_t length, unsigned cid, uint64_t flags,
         uct_device_completion_t *tl_comp, uint32_t opcode, bool is_atomic,
@@ -310,9 +310,8 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_single(
 
 template<ucs_device_level_t level>
 UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put(
-        uct_device_ep_h tl_ep,
-        const uct_device_local_mem_list_elem_t *src_uct_elem,
-        const uct_device_mem_element_t *tl_mem_elem, const void *address,
+        uct_device_ep_h tl_ep, const uct_device_mem_elem_t *src_uct_elem,
+        const uct_device_mem_elem_t *tl_mem_elem, const void *address,
         uint64_t remote_address, size_t length, unsigned channel_id,
         uint64_t flags, uct_device_completion_t *comp)
 {
@@ -321,7 +320,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put(
             tl_mem_elem);
     auto local_mem_elem =
             reinterpret_cast<const uct_ib_md_device_mem_element_t*>(
-                    &src_uct_elem->uct_mem_element);
+                    src_uct_elem);
     auto cid = channel_id & ep->channel_mask;
 
     return uct_rc_mlx5_gda_ep_single<level>(ep, tl_mem_elem, address,
@@ -333,7 +332,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put(
 
 template<ucs_device_level_t level>
 UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_atomic_add(
-        uct_device_ep_h tl_ep, const uct_device_mem_element_t *tl_mem_elem,
+        uct_device_ep_h tl_ep, const uct_device_mem_elem_t *tl_mem_elem,
         uint64_t value, uint64_t remote_address, unsigned channel_id,
         uint64_t flags, uct_device_completion_t *comp)
 {
diff --git a/test/gtest/ucp/cuda/test_kernels.cu b/test/gtest/ucp/cuda/test_kernels.cu
index 309bb2fa265..18a5dc7ff9a 100644
--- a/test/gtest/ucp/cuda/test_kernels.cu
+++ b/test/gtest/ucp/cuda/test_kernels.cu
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -108,56 +108,6 @@ private:
     ucp_device_request_t *m_ptr;
 };
 
-UCS_F_DEVICE ucs_status_t
-ucp_test_kernel_get_state(const test_ucp_device_kernel_params_t &params,
-                          test_ucp_device_kernel_result_t &result)
-{
-    ucp_device_request_t *req_ptr = nullptr;
-    const uct_device_mem_element_t *uct_elem;
-    uct_device_ep_t *device_ep;
-    uint64_t remote_address;
-    uct_device_completion_t *comp;
-    ucs_status_t status = UCS_OK;
-
-    if (nullptr == params.remote_mem_list) {
-        return UCS_OK;
-    }
-
-    __syncthreads();
-    if (threadIdx.x == 0) {
-        for (unsigned i = 0; i < params.remote_mem_list->length; ++i) {
-            status = ucp_device_prepare_send_remote(params.remote_mem_list, i,
-                                                    remote_address, req_ptr,
-                                                    device_ep, uct_elem, comp);
-            if ((status == UCS_OK) && (device_ep != nullptr)) {
-                break;
-            }
-        }
-
-        result.producer_index = 0;
-        result.ready_index    = 0;
-#if HAVE_MLX5_DV
-        if ((status == UCS_OK) &&
-            (device_ep != nullptr) &&
-            (device_ep->uct_tl_id == UCT_DEVICE_TL_RC_MLX5_GDA)) {
-            uint16_t wqe_cnt;
-            uct_rc_gdaki_dev_ep_t *ep =
-                        reinterpret_cast<uct_rc_gdaki_dev_ep_t*>(device_ep);
-            unsigned i;
-
-            for (i = 0; i < params.num_channels; i++) {
-                result.producer_index +=
-                        uct_rc_mlx5_gda_parse_cqe(ep, i, &wqe_cnt, nullptr) + 1;
-                result.ready_index += ep->qps[i].sq_ready_index;
-            }
-        }
-#endif
-    }
-
-    __syncthreads();
-    return status;
-}
-
 template<ucs_device_level_t level>
 UCS_F_DEVICE void
 ucp_test_kernel_job(const test_ucp_device_kernel_params_t &params,
@@ -175,11 +125,6 @@ ucp_test_kernel_job(const test_ucp_device_kernel_params_t &params,
             shared_reqs[device_request<level>::num_shared_reqs()];
     device_request<level> req(shared_reqs);
 
-    status = ucp_test_kernel_get_state(params, *result_ptr);
-    if (status != UCS_OK) {
-        return;
-    }
-
     ucp_device_request_t *req_ptr = params.with_request ? req.ptr() : nullptr;
     uint64_t flags                = params.with_no_delay ?
                                                 UCT_DEVICE_FLAG_NODELAY : 0;
@@ -196,11 +141,6 @@ ucp_test_kernel_job(const test_ucp_device_kernel_params_t &params,
     // function to the API.
     status = ucp_test_kernel_do_operation<level>(params, UCT_DEVICE_FLAG_NODELAY,
                                                  req.ptr());
-    if (status != UCS_OK) {
-        return;
-    }
-
-    status = ucp_test_kernel_get_state(params, *result_ptr);
 }
 
 template<ucs_device_level_t level>
@@ -256,8 +196,6 @@ launch_test_ucp_device_kernel(const test_ucp_device_kernel_params_t &params)
 
     ucx_cuda::device_result_ptr<test_ucp_device_kernel_result_t> result;
     result->status         = UCS_ERR_NOT_IMPLEMENTED;
-    result->producer_index = 0;
-    result->ready_index    = 0;
 
     switch (params.level) {
     case UCS_DEVICE_LEVEL_THREAD:
diff --git a/test/gtest/ucp/cuda/test_kernels.h b/test/gtest/ucp/cuda/test_kernels.h
index fcac4d39312..b9d10f9a142 100644
--- a/test/gtest/ucp/cuda/test_kernels.h
+++ b/test/gtest/ucp/cuda/test_kernels.h
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -51,8 +51,6 @@ typedef struct {
 
 struct test_ucp_device_kernel_result_t {
     ucs_status_t status;
-    uint64_t     producer_index;
-    uint64_t     ready_index;
 };
 
 test_ucp_device_kernel_result_t
diff --git a/test/gtest/ucp/test_ucp_device.cc b/test/gtest/ucp/test_ucp_device.cc
index fe2f192d235..0e281e714f7 100644
--- a/test/gtest/ucp/test_ucp_device.cc
+++ b/test/gtest/ucp/test_ucp_device.cc
@@ -1,5 +1,5 @@
 /**
- * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -538,21 +538,6 @@ class test_ucp_device_kernel : public test_ucp_device {
         ASSERT_UCS_OK(result.status);
         return result;
     }
-
-    void check_result(const test_ucp_device_kernel_params_t &params,
-                      const test_ucp_device_kernel_result_t &result,
-                      unsigned count)
-    {
-        unsigned num_threads = params.num_threads;
-        if (params.level == UCS_DEVICE_LEVEL_WARP) {
-            num_threads /= UCS_DEVICE_NUM_THREADS_IN_WARP;
-        }
-
-        uint64_t expected = params.num_iters * num_threads * count;
-        EXPECT_UCS_OK(result.status);
-        EXPECT_EQ(expected, result.producer_index);
-        EXPECT_EQ(expected, result.ready_index);
-    }
 };
 
 UCS_TEST_P(test_ucp_device_kernel, local_counter)
@@ -781,11 +766,10 @@ UCS_TEST_SKIP_COND_P(test_ucp_device_xfer, put_stress_test,
     params.num_iters                    = 1000;
     params.num_blocks                   = 1;
     params.num_threads                  = MAX_THREADS;
-    auto result                         = launch_kernel(params);
+    launch_kernel(params);
 
     // Check proper index received data
     list.dst_pattern_check(mem_list_index, mem_list::SEED_SRC);
-    check_result(params, result, 1);
 }
 
 UCS_TEST_P(test_ucp_device_xfer, counter)
diff --git a/test/gtest/uct/cuda/test_cuda_ipc_device.cc b/test/gtest/uct/cuda/test_cuda_ipc_device.cc
index 0c1cff31d88..1f64979875e 100644
--- a/test/gtest/uct/cuda/test_cuda_ipc_device.cc
+++ b/test/gtest/uct/cuda/test_cuda_ipc_device.cc
@@ -189,7 +189,7 @@ UCS_TEST_P(test_cuda_ipc_rma, get_mem_elem_pack)
     mapped_buffer sendbuf(length, SEED1, *m_sender, 0, UCS_MEMORY_TYPE_CUDA);
     mapped_buffer recvbuf(length, SEED2, *m_receiver, 0, UCS_MEMORY_TYPE_CUDA);
 
-    uct_device_mem_element_t mem_elem_host;
+    uct_device_mem_elem_t mem_elem_host;
     EXPECT_UCS_OK(uct_md_mem_elem_pack(m_sender->md(), sendbuf.memh(),
                                        recvbuf.rkey(), &mem_elem_host));
 }
@@ -221,26 +221,22 @@ UCS_TEST_P(test_cuda_ipc_rma_device, put_device)
     mapped_buffer sendbuf(length, SEED1, *m_sender, 0, UCS_MEMORY_TYPE_CUDA);
     mapped_buffer recvbuf(length, SEED2, *m_receiver, 0, UCS_MEMORY_TYPE_CUDA);
 
-    uct_device_local_mem_list_elem_t src_elem_host;
+    uct_device_mem_elem_t src_elem_host;
     ASSERT_UCS_OK(uct_md_mem_elem_pack(m_sender->md(), sendbuf.memh(),
-                                       recvbuf.rkey(),
-                                       &src_elem_host.uct_mem_element));
+                                       recvbuf.rkey(), &src_elem_host));
 
 
-    uct_device_local_mem_list_elem_t *src_elem;
-    ASSERT_EQ(CUDA_SUCCESS,
-              cuMemAlloc((CUdeviceptr*)&src_elem,
-                         sizeof(uct_device_local_mem_list_elem_t)));
-    ASSERT_EQ(CUDA_SUCCESS,
-              cuMemcpyHtoD((CUdeviceptr)src_elem, &src_elem_host,
-                           sizeof(uct_device_local_mem_list_elem_t)));
+    uct_device_mem_elem_t *src_elem;
+    ASSERT_EQ(CUDA_SUCCESS, cuMemAlloc((CUdeviceptr*)&src_elem,
+                                       sizeof(uct_device_mem_elem_t)));
+    ASSERT_EQ(CUDA_SUCCESS, cuMemcpyHtoD((CUdeviceptr)src_elem, &src_elem_host,
+                                         sizeof(uct_device_mem_elem_t)));
 
-    uct_device_mem_element_t *mem_elem;
+    uct_device_mem_elem_t *mem_elem;
     ASSERT_EQ(CUDA_SUCCESS, cuMemAlloc((CUdeviceptr*)&mem_elem,
-                                       sizeof(uct_device_mem_element_t)));
-    ASSERT_EQ(CUDA_SUCCESS, cuMemcpyHtoD((CUdeviceptr)mem_elem,
-                                         &src_elem_host.uct_mem_element,
-                                         sizeof(uct_device_mem_element_t)));
+                                       sizeof(uct_device_mem_elem_t)));
+    ASSERT_EQ(CUDA_SUCCESS, cuMemcpyHtoD((CUdeviceptr)mem_elem, &src_elem_host,
+                                         sizeof(uct_device_mem_elem_t)));
 
     uct_device_ep_h device_ep;
     ASSERT_UCS_OK(uct_ep_get_device_ep(m_sender->ep(0), &device_ep));
@@ -262,7 +258,7 @@ UCS_TEST_P(test_cuda_ipc_rma_device, atomic_add_device)
     unsigned           num_threads   = get_num_threads();
     unsigned           num_blocks    = get_num_blocks();
     uct_device_ep_h device_ep;
-    uct_device_mem_element_t *mem_elem;
+    uct_device_mem_elem_t *mem_elem;
 
     if (device_level == UCS_DEVICE_LEVEL_GRID) {
         GTEST_SKIP() << "Grid level is not supported";
@@ -275,7 +271,7 @@ UCS_TEST_P(test_cuda_ipc_rma_device, atomic_add_device)
     mapped_buffer signal(sizeof(uint64_t), 0, *m_receiver, 0, UCS_MEMORY_TYPE_CUDA);
     ASSERT_UCS_OK(uct_ep_get_device_ep(m_sender->ep(0), &device_ep));
 
-    uct_device_mem_element_t mem_elem_host;
+    uct_device_mem_elem_t mem_elem_host;
     ASSERT_EQ(CUDA_SUCCESS, cuMemAlloc((CUdeviceptr*)&mem_elem, mem_elem_size));
     ASSERT_UCS_OK(uct_md_mem_elem_pack(m_sender->md(), nullptr, signal.rkey(),
                                        &mem_elem_host));
diff --git a/test/gtest/uct/cuda/test_kernels.cu b/test/gtest/uct/cuda/test_kernels.cu
index a3dfb7404f8..ca28709880a 100644
--- a/test/gtest/uct/cuda/test_kernels.cu
+++ b/test/gtest/uct/cuda/test_kernels.cu
@@ -13,9 +13,8 @@
 namespace ucx_cuda {
 
 static __global__ void
-uct_put_kernel(uct_device_ep_h ep,
-               const uct_device_local_mem_list_elem_t *src_elem,
-               const uct_device_mem_element_t *mem_elem, const void *va,
+uct_put_kernel(uct_device_ep_h ep, const uct_device_mem_elem_t *src_elem,
+               const uct_device_mem_elem_t *mem_elem, const void *va,
                uint64_t rva, size_t length, ucs_status_t *status_p)
 {
     uct_device_completion_t comp;
@@ -39,8 +38,8 @@ uct_put_kernel(uct_device_ep_h ep,
  * Basic single element put operation.
  */
 ucs_status_t launch_uct_put(uct_device_ep_h ep,
-                            const uct_device_local_mem_list_elem_t *src_elem,
-                            const uct_device_mem_element_t *mem_elem,
+                            const uct_device_mem_elem_t *src_elem,
+                            const uct_device_mem_elem_t *mem_elem,
                             const void *va, uint64_t rva, size_t length)
 {
     device_result_ptr<ucs_status_t> status = UCS_ERR_NOT_IMPLEMENTED;
@@ -52,7 +51,7 @@ ucs_status_t launch_uct_put(uct_device_ep_h ep,
 }
 
 static __global__ void
-uct_atomic_kernel(uct_device_ep_h ep, uct_device_mem_element_t *mem_elem,
+uct_atomic_kernel(uct_device_ep_h ep, uct_device_mem_elem_t *mem_elem,
                   uint64_t rva, uint64_t add, ucs_status_t *status_p)
 {
     uct_device_completion_t comp;
@@ -75,7 +74,7 @@ uct_atomic_kernel(uct_device_ep_h ep, uct_device_mem_element_t *mem_elem,
  * Atomic operation.
  */
 ucs_status_t launch_uct_atomic(uct_device_ep_h ep,
-                               uct_device_mem_element_t *mem_elem, uint64_t rva,
+                               uct_device_mem_elem_t *mem_elem, uint64_t rva,
                                uint64_t add)
 {
     device_result_ptr<ucs_status_t> status = UCS_ERR_NOT_IMPLEMENTED;
diff --git a/test/gtest/uct/cuda/test_kernels.h b/test/gtest/uct/cuda/test_kernels.h
index fa5c22607ed..ff23a4f2d28 100644
--- a/test/gtest/uct/cuda/test_kernels.h
+++ b/test/gtest/uct/cuda/test_kernels.h
@@ -14,12 +14,12 @@
 namespace ucx_cuda {
 
 ucs_status_t launch_uct_put(uct_device_ep_h ep,
-                            const uct_device_local_mem_list_elem_t *src_elem,
-                            const uct_device_mem_element_t *mem_elem,
+                            const uct_device_mem_elem_t *src_elem,
+                            const uct_device_mem_elem_t *mem_elem,
                             const void *va, uint64_t rva, size_t length);
 
 ucs_status_t launch_uct_atomic(uct_device_ep_h ep,
-                               uct_device_mem_element_t *mem_elem, uint64_t rva,
+                               uct_device_mem_elem_t *mem_elem, uint64_t rva,
                                uint64_t add);
 
 }; // namespace ucx_cuda
diff --git a/test/gtest/uct/cuda/test_kernels_uct.cu b/test/gtest/uct/cuda/test_kernels_uct.cu
index 0e1528d2876..4665f5cf5db 100644
--- a/test/gtest/uct/cuda/test_kernels_uct.cu
+++ b/test/gtest/uct/cuda/test_kernels_uct.cu
@@ -99,9 +99,8 @@ template<typename T> class device_result_ptr {
 
 template<ucs_device_level_t level>
 static __global__ void
-uct_put_kernel(uct_device_ep_h ep,
-               const uct_device_local_mem_list_elem_t *src_elem,
-               const uct_device_mem_element_t *mem_elem, const void *va,
+uct_put_kernel(uct_device_ep_h ep, const uct_device_mem_elem_t *src_elem,
+               const uct_device_mem_elem_t *mem_elem, const void *va,
                uint64_t rva, size_t length, ucs_status_t *status_p)
 {
     uct_device_completion_t comp;
@@ -119,12 +118,11 @@ uct_put_kernel(uct_device_ep_h ep,
     }
 }
 
-ucs_status_t launch_uct_put(uct_device_ep_h device_ep,
-                            const uct_device_local_mem_list_elem_t *src_elem,
-                            const uct_device_mem_element_t *mem_elem,
-                            const void *va, uint64_t rva, size_t length,
-                            ucs_device_level_t level, unsigned num_threads,
-                            unsigned num_blocks)
+ucs_status_t
+launch_uct_put(uct_device_ep_h device_ep, const uct_device_mem_elem_t *src_elem,
+               const uct_device_mem_elem_t *mem_elem, const void *va,
+               uint64_t rva, size_t length, ucs_device_level_t level,
+               unsigned num_threads, unsigned num_blocks)
 {
     device_result_ptr<ucs_status_t> status = UCS_ERR_NOT_IMPLEMENTED;
     cudaError_t st;
@@ -165,8 +163,7 @@ ucs_status_t launch_uct_put(uct_device_ep_h device_ep,
 
 template<ucs_device_level_t level>
 static __global__ void
-uct_atomic_kernel(uct_device_ep_h ep,
-                  const uct_device_mem_element_t *mem_elem,
+uct_atomic_kernel(uct_device_ep_h ep, const uct_device_mem_elem_t *mem_elem,
                   uint64_t rva, uint64_t add, ucs_status_t *status_p)
 {
     uct_device_completion_t comp;
@@ -179,11 +176,9 @@ uct_atomic_kernel(uct_device_ep_h ep,
 }
 
 ucs_status_t launch_uct_atomic(uct_device_ep_h device_ep,
-                               const uct_device_mem_element_t *mem_elem,
-                               uint64_t rva,
-                               uint64_t add,
-                               ucs_device_level_t level,
-                               unsigned num_threads,
+                               const uct_device_mem_elem_t *mem_elem,
+                               uint64_t rva, uint64_t add,
+                               ucs_device_level_t level, unsigned num_threads,
                                unsigned num_blocks)
 {
     device_result_ptr<ucs_status_t> status = UCS_ERR_NOT_IMPLEMENTED;
diff --git a/test/gtest/uct/cuda/test_kernels_uct.h b/test/gtest/uct/cuda/test_kernels_uct.h
index 886654a77cd..b62ecda67a6 100644
--- a/test/gtest/uct/cuda/test_kernels_uct.h
+++ b/test/gtest/uct/cuda/test_kernels_uct.h
@@ -15,20 +15,18 @@ namespace cuda_uct {
 
 int launch_memcmp(const void *s1, const void *s2, size_t size);
 
-ucs_status_t launch_uct_put(uct_device_ep_h device_ep,
-                            const uct_device_local_mem_list_elem_t *src_elem,
-                            const uct_device_mem_element_t *mem_elem,
-                            const void *va, uint64_t rva, size_t length,
-                            ucs_device_level_t level, unsigned num_threads,
-                            unsigned num_blocks);
+ucs_status_t
+launch_uct_put(uct_device_ep_h device_ep, const uct_device_mem_elem_t *src_elem,
+               const uct_device_mem_elem_t *mem_elem, const void *va,
+               uint64_t rva, size_t length, ucs_device_level_t level,
+               unsigned num_threads, unsigned num_blocks);
 
 ucs_status_t launch_uct_atomic(uct_device_ep_h device_ep,
-                               const uct_device_mem_element_t *mem_elem,
-                               uint64_t rva,
-                               uint64_t add,
-                               ucs_device_level_t level,
-                               unsigned num_threads,
+                               const uct_device_mem_elem_t *mem_elem,
+                               uint64_t rva, uint64_t add,
+                               ucs_device_level_t level, unsigned num_threads,
                                unsigned num_blocks);
+
 }; // namespace cuda_uct
 
 #endif
diff --git a/test/gtest/uct/test_device.cc b/test/gtest/uct/test_device.cc
index dad85a1388b..5dd8f7d650a 100644
--- a/test/gtest/uct/test_device.cc
+++ b/test/gtest/uct/test_device.cc
@@ -1,5 +1,5 @@
 /**
-* Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+* Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025-2026. ALL RIGHTS RESERVED.
 * See file LICENSE for terms.
 */
 
@@ -101,17 +101,6 @@ class test_device : public uct_test {
         ucs_status_t status;
 
         uct_test::init();
-        m_cuda_dev = uct_cuda_get_cuda_device(GetParam()->sys_device);
-        ASSERT_NE(m_cuda_dev, CU_DEVICE_INVALID) << " sys_device "
-                  << static_cast<int>(GetParam()->sys_device);
-
-        status = UCT_CUDADRV_FUNC_LOG_ERR(
-                cuDevicePrimaryCtxRetain(&ctx, m_cuda_dev));
-        ASSERT_UCS_OK(status);
-
-        status = UCT_CUDADRV_FUNC_LOG_ERR(cuCtxPushCurrent(ctx));
-        ASSERT_UCS_OK(status);
-
         m_receiver = uct_test::create_entity(0);
         m_entities.push_back(m_receiver);
 
@@ -119,11 +108,29 @@ class test_device : public uct_test {
         m_entities.push_back(m_sender);
 
         m_sender->connect(0, *m_receiver, 0);
+
+        m_cuda_dev = uct_cuda_get_cuda_device(
+                m_sender->iface_attr().ctl_device);
+        if (m_cuda_dev == CU_DEVICE_INVALID) {
+            return;
+        }
+
+        status = UCT_CUDADRV_FUNC_LOG_ERR(
+                cuDevicePrimaryCtxRetain(&ctx, m_cuda_dev));
+        ASSERT_UCS_OK(status);
+
+        status = UCT_CUDADRV_FUNC_LOG_ERR(cuCtxPushCurrent(ctx));
+        ASSERT_UCS_OK(status);
     }
 
     void cleanup()
     {
         uct_test::cleanup();
+
+        if (m_cuda_dev == CU_DEVICE_INVALID) {
+            return;
+        }
+
         (void)UCT_CUDADRV_FUNC_LOG_WARN(cuCtxPopCurrent(NULL));
         (void)UCT_CUDADRV_FUNC_LOG_WARN(cuDevicePrimaryCtxRelease(m_cuda_dev));
     }
@@ -147,29 +154,27 @@ UCS_TEST_P(test_device, put)
     mapped_buffer sendbuf(length, SEED1, *m_sender, 0, UCS_MEMORY_TYPE_CUDA);
     mapped_buffer recvbuf(length, SEED2, *m_receiver, 0, UCS_MEMORY_TYPE_CUDA);
 
-    uct_device_local_mem_list_elem_t src_elem_host;
-    src_elem_host.addr = nullptr;
+    uct_device_mem_elem_t src_elem_host;
     ASSERT_UCS_OK(uct_md_mem_elem_pack(m_sender->md(), sendbuf.memh(),
-                                       recvbuf.rkey(),
-                                       &src_elem_host.uct_mem_element));
+                                       recvbuf.rkey(), &src_elem_host));
 
-    mapped_buffer src_elembuf(sizeof(uct_device_local_mem_list_elem_t), 0,
-                              *m_sender, 0, UCS_MEMORY_TYPE_CUDA);
+    mapped_buffer src_elembuf(sizeof(src_elem_host), 0, *m_sender, 0,
+                              UCS_MEMORY_TYPE_CUDA);
     ASSERT_EQ(CUDA_SUCCESS,
               cuMemcpyHtoD((CUdeviceptr)src_elembuf.ptr(), &src_elem_host,
-                           sizeof(uct_device_local_mem_list_elem_t)));
+                           sizeof(src_elem_host)));
 
-    mapped_buffer rem_elembuf(sizeof(uct_device_mem_element_t), 0, *m_sender, 0,
+    mapped_buffer rem_elembuf(sizeof(src_elem_host), 0, *m_sender, 0,
                               UCS_MEMORY_TYPE_CUDA);
-    ASSERT_EQ(CUDA_SUCCESS, cuMemcpyHtoD((CUdeviceptr)rem_elembuf.ptr(),
-                                         &src_elem_host.uct_mem_element,
-                                         sizeof(uct_device_mem_element_t)));
+    ASSERT_EQ(CUDA_SUCCESS,
+              cuMemcpyHtoD((CUdeviceptr)rem_elembuf.ptr(), &src_elem_host,
+                           sizeof(src_elem_host)));
 
     uct_device_ep_h dev_ep;
     ASSERT_UCS_OK(uct_ep_get_device_ep(m_sender->ep(0), &dev_ep));
     ASSERT_UCS_OK(ucx_cuda::launch_uct_put(
-            dev_ep, (const uct_device_local_mem_list_elem_t*)src_elembuf.ptr(),
-            (const uct_device_mem_element_t*)rem_elembuf.ptr(), sendbuf.ptr(),
+            dev_ep, (const uct_device_mem_elem_t*)src_elembuf.ptr(),
+            (const uct_device_mem_elem_t*)rem_elembuf.ptr(), sendbuf.ptr(),
             (uintptr_t)recvbuf.ptr(), length));
 
     recvbuf.pattern_check(SEED1);
@@ -187,18 +192,17 @@ UCS_TEST_P(test_device, atomic)
     uint64_t signal_val = 0;
     size_t i;
 
-    mapped_buffer elembuf_host(sizeof(uct_device_mem_element_t), 0, *m_sender,
-                               0, UCS_MEMORY_TYPE_HOST);
-    mapped_buffer elembuf(sizeof(uct_device_mem_element_t), 0, *m_sender, 0,
+    mapped_buffer elembuf_host(sizeof(uct_device_mem_elem_t), 0, *m_sender, 0,
+                               UCS_MEMORY_TYPE_HOST);
+    mapped_buffer elembuf(sizeof(uct_device_mem_elem_t), 0, *m_sender, 0,
                           UCS_MEMORY_TYPE_CUDA);
-    uct_device_mem_element_t *mem_elem_host = (uct_device_mem_element_t*)
-                                                      elembuf_host.ptr();
-    uct_device_mem_element_t *mem_elem      = (uct_device_mem_element_t*)
-                                                      elembuf.ptr();
+    uct_device_mem_elem_t *mem_elem_host = (uct_device_mem_elem_t*)
+                                                   elembuf_host.ptr();
+    uct_device_mem_elem_t *mem_elem = (uct_device_mem_elem_t*)elembuf.ptr();
     ASSERT_UCS_OK(uct_md_mem_elem_pack(m_sender->md(), nullptr, signal.rkey(),
                                        mem_elem_host));
     ASSERT_EQ(CUDA_SUCCESS, cuMemcpyHtoD((CUdeviceptr)mem_elem, mem_elem_host,
-                                         sizeof(uct_device_mem_element_t)));
+                                         sizeof(uct_device_mem_elem_t)));
 
     uct_device_ep_h dev_ep;
     ASSERT_UCS_OK(uct_ep_get_device_ep(m_sender->ep(0), &dev_ep));