UCT/ROCM: initial commit for device initiated ipc put

config/hip.am

@@ -0,0 +1,39 @@
+#
+# Copyright (c) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
+# See file LICENSE for terms.
+#
+
+SUFFIXES = .hip
+
+HIPCC ?= hipcc
+
+HIPCC_CMD = $(HIPCC) -DHAVE_CONFIG_H -DUCT_DEVICE_CODE_HIP -fPIE -I$(top_srcdir)/src -I$(top_builddir)/src $(BASE_CXXFLAGS) $(CXXFLAGS) $(HIP_CPPFLAGS) $(HIP_CXXFLAGS) $(HIPCC_EXTRA_FLAGS) -Wno-c++20-extensions -c $< -MT $@ -MF $(DEPDIR)/hip/$@.d -MMD -o $@
+HIPCC_LT_CMD = $(LIBTOOL) --tag=CXX --mode=compile $(HIPCC_CMD)
+
+define hipcc-build
+	@$(MKDIR_P) $(shell dirname $(DEPDIR)/hip/$@)
+	@$(if $(filter false,$(AM_V_P)),echo "  HIPCC    $@")
+	@$(if $(filter .o,$(suffix $@)),$(HIPCC_CMD),$(HIPCC_LT_CMD)) $($(1)) $(if $(filter false,$(AM_V_P)), >/dev/null)
+endef
+
+define hipcc-source
+	EXTRA_DIST += $(2)
+$(1): $(2)
+	$$(call hipcc-build,$(3))
+endef
+
+# Default rules when no target-specific compile flags are required
+.hip.o:
+	$(call hipcc-build)
+
+.hip.lo:
+	$(call hipcc-build)
+
+HIP_DEP_FILES := $(shell find $(DEPDIR)/hip/ -type f -name *.d 2>/dev/null)
+-include $(HIP_DEP_FILES)
+
+clean-local:
+	-rm -rf $(DEPDIR)/hip
+
+distclean-local:
+	-rm -rf $(DEPDIR)/hip

src/ucp/core/ucp_device.c

@@ -1,5 +1,6 @@
 /**
  * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (C) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -434,7 +435,11 @@ ucs_status_t
 ucp_device_local_mem_list_create(const ucp_device_mem_list_params_t *params,
                                  ucp_device_local_mem_list_h *mem_list_h)
 {
+#if HAVE_ROCM
+    const ucs_memory_type_t export_mem_type = UCS_MEMORY_TYPE_ROCM;
+#else
     const ucs_memory_type_t export_mem_type = UCS_MEMORY_TYPE_CUDA;
+#endif
     ucs_status_t status;
     uct_allocated_memory_t mem;
     ucs_sys_device_t local_sys_dev;
@@ -682,7 +687,11 @@ ucs_status_t
 ucp_device_remote_mem_list_create(const ucp_device_mem_list_params_t *params,
                                   ucp_device_remote_mem_list_h *mem_list_h)
 {
+#if HAVE_ROCM
+    const ucs_memory_type_t export_mem_type = UCS_MEMORY_TYPE_ROCM;
+#else
     const ucs_memory_type_t export_mem_type = UCS_MEMORY_TYPE_CUDA;
+#endif
     ucs_status_t status;
     uct_allocated_memory_t mem;
 

src/ucp/wireup/select.c

@@ -1,6 +1,7 @@
 /**
  * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2001-2016. ALL RIGHTS RESERVED.
  * Copyright (C) Los Alamos National Security, LLC. 2019 ALL RIGHTS RESERVED.
+ * Copyright (C) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -2504,6 +2505,14 @@ ucp_wireup_add_device_lanes(const ucp_wireup_select_params_t *select_params,
     found_lane = ucp_wireup_add_bw_lanes(select_params, &bw_info,
                                          mem_type_tl_bitmap, UCP_NULL_LANE,
                                          select_ctx, 0);
+
+    /* Add device lanes for ROCm memory */
+    ucp_wireup_memaccess_bitmap(context, UCS_MEMORY_TYPE_ROCM,
+                                &mem_type_tl_bitmap);
+    found_lane |= ucp_wireup_add_bw_lanes(select_params, &bw_info,
+                                          mem_type_tl_bitmap, UCP_NULL_LANE,
+                                          select_ctx, 0);
+
     if (!found_lane) {
         ucs_debug("ep %p: could not find device lanes", select_params->ep);
     }

src/ucs/sys/device_code.h

@@ -1,5 +1,6 @@
 /**
  * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (C) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -13,11 +14,11 @@
 /*
  * Declare GPU specific functions
  */
-#ifdef __NVCC__
+#if defined(__NVCC__) || defined(__HIPCC__)
 #define UCS_F_DEVICE __device__ __forceinline__ static
 #else
 #define UCS_F_DEVICE static inline
-#endif /* __NVCC__ */
+#endif /* __NVCC__ || __HIPCC__ */
 
 
 #ifndef UCP_DEVICE_ENABLE_PARAMS_CHECK

src/uct/api/device/uct_device_impl.h

@@ -1,5 +1,6 @@
 /**
  * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
+ * Copyright (C) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
  *
  * See file LICENSE for terms.
  */
@@ -10,10 +11,14 @@
 #include "uct_device_types.h"
 
 #include <uct/api/uct_def.h>
+#if HAVE_ROCM
+#include <uct/rocm/ipc/rocm_ipc.h>
+#else
 #include <uct/cuda/cuda_ipc/cuda_ipc.cuh>
+#endif
 #include <ucs/sys/device_code.h>
 
-#if __has_include(<uct/ib/mlx5/gdaki/gdaki.cuh>) && __has_include(<infiniband/mlx5dv.h>)
+#if defined(__NVCC__) && __has_include(<uct/ib/mlx5/gdaki/gdaki.cuh>) && __has_include(<infiniband/mlx5dv.h>)
 #  include <uct/ib/mlx5/gdaki/gdaki.cuh>
 #  define UCT_RC_MLX5_GDA_SUPPORTED 1
 #else
@@ -24,7 +29,11 @@ union uct_device_completion {
 #if UCT_RC_MLX5_GDA_SUPPORTED
     uct_rc_gda_completion_t   rc_gda;
 #endif
+#if HAVE_ROCM
+    uct_rocm_ipc_completion_t rocm_ipc;
+#else
     uct_cuda_ipc_completion_t cuda_ipc;
+#endif
 };
 
 
@@ -73,12 +82,20 @@ uct_device_ep_put(uct_device_ep_h device_ep,
                                              channel_id, flags, comp);
     } else
 #endif
+#if HAVE_ROCM
+            if (device_ep->uct_tl_id == UCT_DEVICE_TL_ROCM_IPC) {
+        return uct_rocm_ipc_ep_put<level>(device_ep, mem_elem, address,
+                                          remote_address, length, flags, comp);
+    } else
+#else
     if (device_ep->uct_tl_id == UCT_DEVICE_TL_CUDA_IPC) {
         return uct_cuda_ipc_ep_put<level>(device_ep, mem_elem, address,
                                           remote_address, length, flags, comp);
+    } else
+#endif
+    {
+        return UCS_ERR_UNSUPPORTED;
     }
-
-    return UCS_ERR_UNSUPPORTED;
 }
 
 
@@ -122,12 +139,20 @@ UCS_F_DEVICE ucs_status_t uct_device_ep_atomic_add(
                                                     channel_id, flags, comp);
     } else
 #endif
+#if HAVE_ROCM
+            if (device_ep->uct_tl_id == UCT_DEVICE_TL_ROCM_IPC) {
+        return uct_rocm_ipc_ep_atomic_add<level>(device_ep, mem_elem, inc_value,
+                                                 remote_address, flags, comp);
+    } else
+#else
     if (device_ep->uct_tl_id == UCT_DEVICE_TL_CUDA_IPC) {
         return uct_cuda_ipc_ep_atomic_add<level>(device_ep, mem_elem, inc_value,
                                                  remote_address, flags, comp);
+    } else
+#endif
+    {
+        return UCS_ERR_UNSUPPORTED;
     }
-
-    return UCS_ERR_UNSUPPORTED;
 }
 
 
@@ -149,14 +174,20 @@ 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,
         uint64_t address, void **addr_p)
 {
-    if (device_ep->uct_tl_id != UCT_DEVICE_TL_CUDA_IPC) {
+#if HAVE_ROCM
+    if (device_ep->uct_tl_id == UCT_DEVICE_TL_ROCM_IPC) {
+        return uct_rocm_ipc_ep_get_ptr(device_ep, mem_elem, address, addr_p);
+    } else
+#else
+    if (device_ep->uct_tl_id == UCT_DEVICE_TL_CUDA_IPC) {
+        return uct_cuda_ipc_ep_get_ptr(device_ep, mem_elem, address, addr_p);
+    } else
+#endif
+    {
         return UCS_ERR_UNSUPPORTED;
     }
-
-    return uct_cuda_ipc_ep_get_ptr(device_ep, mem_elem, address, addr_p);
 }
 
-
 /**
  * @ingroup UCT_DEVICE
  * @brief Progress all operations on device endpoint @a device_ep.

src/uct/api/device/uct_device_types.h

@@ -1,6 +1,7 @@
 /**
  * Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2025. ALL RIGHTS RESERVED.
- *
+ * Copyright (C) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
+ * 
  * See file LICENSE for terms.
  */
 
@@ -36,6 +37,21 @@ typedef struct {
 } uct_cuda_ipc_completion_t;
 
 
+/**
+ * @brief Device memory element for ROCm IPC.
+ */
+typedef struct {
+    ptrdiff_t mapped_offset;
+} uct_rocm_ipc_device_mem_element_t;
+
+
+/**
+ * @brief Completion object for device ROCm IPC operations.
+ */
+typedef struct {
+} uct_rocm_ipc_completion_t;
+
+
 /**
  * @brief Device memory element for GDAKI.
  */
@@ -57,6 +73,7 @@ typedef enum {
 typedef enum {
     UCT_DEVICE_TL_RC_MLX5_GDA,
     UCT_DEVICE_TL_CUDA_IPC,
+    UCT_DEVICE_TL_ROCM_IPC,
     UCT_DEVICE_TL_LAST
 } uct_device_tl_id_t;
 
@@ -75,6 +92,7 @@ typedef union uct_device_completion uct_device_completion_t;
 union uct_device_mem_element {
     uct_ib_md_device_mem_element_t       ib_md_mem_element;
     uct_cuda_ipc_md_device_mem_element_t cuda_ipc_md_mem_element;
+    uct_rocm_ipc_device_mem_element_t    rocm_ipc_mem_element;
 };
 
 

src/uct/rocm/Makefile.am

@@ -1,5 +1,6 @@
 #
 # Copyright (c) NVIDIA CORPORATION & AFFILIATES, 2001-2018. ALL RIGHTS RESERVED.
+# Copyright (C) Advanced Micro Devices, Inc. 2026. ALL RIGHTS RESERVED.
 # See file LICENSE for terms.
 #
 
@@ -14,6 +15,10 @@ libuct_rocm_la_LDFLAGS  = $(ROCM_LDFLAGS) $(ROCM_LIBS) -version-info $(SOVERSION
                           $(patsubst %, -Xlinker %, -L$(ROCM_ROOT)/lib -rpath $(ROCM_ROOT)/hip/lib -rpath $(ROCM_ROOT)/lib) \
                           $(patsubst %, -Xlinker %, --enable-new-dtags) \
                           $(patsubst %, -Xlinker %, -rpath $(ROCM_ROOT)/lib64)
+libuct_rocm_ladir       = $(includedir)/uct/rocm
+
+nobase_dist_libuct_rocm_la_HEADERS = \
+	ipc/rocm_ipc.h
 
 noinst_HEADERS = \
 	base/rocm_base.h \

src/uct/rocm/copy/rocm_copy_md.c

@@ -1,5 +1,5 @@
 /*
- * Copyright (C) Advanced Micro Devices, Inc. 2019-2023. ALL RIGHTS RESERVED.
+ * Copyright (C) Advanced Micro Devices, Inc. 2019-2026. ALL RIGHTS RESERVED.
  * See file LICENSE for terms.
  */
 
@@ -57,7 +57,8 @@ uct_rocm_copy_md_query(uct_md_h uct_md, uct_md_attr_v2_t *md_attr)
     md_attr->cache_mem_types  = UCS_BIT(UCS_MEMORY_TYPE_HOST) |
                                 UCS_BIT(UCS_MEMORY_TYPE_ROCM);
     md_attr->alloc_mem_types  = UCS_BIT(UCS_MEMORY_TYPE_ROCM);
-    md_attr->access_mem_types = UCS_BIT(UCS_MEMORY_TYPE_ROCM);
+    md_attr->access_mem_types = UCS_BIT(UCS_MEMORY_TYPE_HOST) |
+                                UCS_BIT(UCS_MEMORY_TYPE_ROCM);
     md_attr->detect_mem_types = UCS_BIT(UCS_MEMORY_TYPE_ROCM);
     if (md->have_dmabuf) {
         md_attr->dmabuf_mem_types |= UCS_BIT(UCS_MEMORY_TYPE_ROCM);