UCT/GDA: Use gpunetio submodule

Makefile.am

@@ -81,6 +81,7 @@ EXTRA_DIST += debian/copyright
 EXTRA_DIST += debian/control
 EXTRA_DIST += debian/ucx.install
 EXTRA_DIST += debian/ucx-cuda.install
+EXTRA_DIST += debian/ucx-ib-mlx5-gda.install
 EXTRA_DIST += debian/rules
 EXTRA_DIST += debian/source/format
 EXTRA_DIST += ucx.pc.in

debian/rules.in

@@ -9,7 +9,7 @@
 # Uncomment this to turn on verbose mode.
 #export DH_VERBOSE=1
 
-CUDA_OPT = $(shell echo ${DEB_BUILD_PROFILES} | grep -qw cuda \
+CUDA_OPT = $(shell echo ${DEB_BUILD_PROFILES} | grep -qw cuda\\\|gda \
 	&& echo --with-cuda=/usr/local/cuda || echo --without-cuda)
 
 XPMEM_OPT = $(shell echo ${DEB_BUILD_PROFILES} | grep -qw xpmem \

debian/ucx-ib-mlx5-gda.install

@@ -1 +1,2 @@
-+usr/lib/ucx/libuct_ib_mlx5_gda.*
+usr/lib/ucx/libuct_ib_mlx5_gda.*
+usr/include/uct/ib/mlx5/gdaki/*

src/tools/perf/cuda/Makefile.am

@@ -16,7 +16,8 @@ libucx_perftest_cuda_la_SOURCES  = cuda_alloc.c
 if HAVE_NVCC
 if HAVE_GPUNETIO # temporarily all device code depends on doca
 
-NVCC_EXTRA_FLAGS = $(CUDA_CPPFLAGS) $(GPUNETIO_CFLAGS)
+NVCC_EXTRA_FLAGS = $(CUDA_CPPFLAGS) \
+    -I$(top_srcdir)/src/uct/ib/mlx5/gdaki/gpunetio/include
 
 libucx_perftest_cuda_la_LIBADD  += \
     $(abs_top_builddir)/src/tools/perf/lib/libucxperf.la \

src/uct/ib/configure.m4

@@ -257,6 +257,9 @@ AS_IF([test "x$with_ib" = "xyes"],
        AC_CHECK_MEMBERS([struct mlx5_cqe64.ib_stride_index],
                         [], [], [[#include <infiniband/$mlx5_include>]])
 
+       AC_CHECK_DECLS([mlx5dv_devx_umem_reg_ex],
+                      [have_mlx5dv_devx_umem=yes], [], [[#include <infiniband/$mlx5_include>]])
+
        AC_DEFINE([HAVE_IB], 1, [IB support])
 
        AS_IF([test "x$with_dc" != xno -a "x$have_dc_dv" = xyes -a "x$have_mlx5" = "xyes"], [

src/uct/ib/mlx5/gdaki/Makefile.am

@@ -9,7 +9,7 @@ module_LTLIBRARIES             = libuct_ib_mlx5_gda.la
 libuct_ib_mlx5_gda_la_CPPFLAGS = $(BASE_CPPFLAGS) $(IBVERBS_CPPFLAGS) \
                                  $(CUDA_CPPFLAGS)
 libuct_ib_mlx5_gda_la_CFLAGS   = $(BASE_CFLAGS) $(LT_CFLAGS) $(CUDA_CFLAGS) \
-                                 $(GPUNETIO_CFLAGS)
+                                 -I$(top_srcdir)/src/uct/ib/mlx5/gdaki/gpunetio/include
 libuct_ib_mlx5_gda_la_LDFLAGS  = $(IBVERBS_LDFLAGS) $(CUDA_LDFLAGS) \
                                  -version-info $(SOVERSION)
 libuct_ib_mlx5_gda_la_LIBADD   = $(top_builddir)/src/ucs/libucs.la \
@@ -27,6 +27,17 @@ nobase_dist_libuct_ib_mlx5_gda_la_HEADERS = \
 	gdaki.cuh \
 	gdaki_dev.h
 
+gpunetio_commondir = $(includedir)/uct/ib/mlx5/gdaki/common
+dist_gpunetio_common_HEADERS = \
+	gpunetio/include/common/doca_gpunetio_verbs_def.h \
+	gpunetio/include/common/doca_gpunetio_verbs_dev.h
+
+gpunetio_devicedir = $(includedir)/uct/ib/mlx5/gdaki/device
+dist_gpunetio_device_HEADERS = \
+	gpunetio/include/device/doca_gpunetio_dev_verbs_common.cuh \
+	gpunetio/include/device/doca_gpunetio_dev_verbs_qp.cuh \
+	gpunetio/include/device/doca_gpunetio_dev_verbs_cq.cuh
+
 noinst_HEADERS = \
 	gdaki.h
 

src/uct/ib/mlx5/gdaki/configure.m4

@@ -3,55 +3,46 @@
 # See file LICENSE for terms.
 #
 
-AC_ARG_WITH([doca-gpunetio],
-            [AS_HELP_STRING([--with-doca-gpunetio=(DIR)],
-                            [Use DOCA gpunetio (Default is guess)])],
-            [with_doca_gpunetio=$withval],
-            [with_doca_gpunetio=guess])
-
 UCX_CHECK_CUDA
 
-AS_IF([test "x$cuda_happy" = "xyes"],
+save_NVCCFLAGS="$NVCCFLAGS"
+NVCCFLAGS="$NVCCFLAGS -DCCCL_IGNORE_DEPRECATED_CPP_DIALECT"
+AC_MSG_CHECKING([checking cuda/atomic support])
+AC_LANG_PUSH([CUDA])
+AC_COMPILE_IFELSE([AC_LANG_SOURCE([[
+#include <cuda/atomic>
+int v;
+cuda::atomic_ref<int> ref{v};
+]])],
+[AC_MSG_RESULT([yes])
+have_cuda_atomic_support=yes],
+[AC_MSG_RESULT([no])])
+AC_LANG_POP
+NVCCFLAGS="$save_NVCCFLAGS"
+
+AS_IF([test "x$cuda_happy" = "xyes"] && [test "x$have_mlx5" = "xyes"] &&
+      ([test "$CUDA_MAJOR_VERSION" -eq 12 -a "$CUDA_MINOR_VERSION" -ge 2] ||
+       [test "$CUDA_MAJOR_VERSION" -ge 13]) &&
+      [test "x$have_mlx5dv_devx_umem" = "xyes"] &&
+      [test "x$have_cuda_atomic_support" = "xyes"],
       [
-       # Default value
-       GPUNETIO_CFLAGS=""
-       AS_IF([test "x$with_doca_gpunetio" != "xno"],
-             [
-              AS_IF([test "x$with_doca_gpunetio" = "xguess"],
-                    [
-                     AS_IF([$PKG_CONFIG --exists doca-gpunetio],
-                           [GPUNETIO_CFLAGS=$(pkg-config --cflags doca-gpunetio)])
-                    ],
-                    [
-                     GPUNETIO_CFLAGS="-I${with_doca_gpunetio}/include"
-                    ]) # "x$with_doca_gpunetio" != "xguess"
-             ]) # "x$with_doca_gpunetio" != "xno"
-
-       AS_IF([test "x$CUDA_MAJOR_VERSION" = "x12" -a "x$CUDA_MINOR_VERSION" = "x9"],
+       NVCCFLAGS="$NVCCFLAGS -DCCCL_IGNORE_DEPRECATED_CPP_DIALECT"
+
+       AS_IF([test "$CUDA_MAJOR_VERSION" -eq 12 -a "$CUDA_MINOR_VERSION" -ge 9] ||
+             [test "$CUDA_MAJOR_VERSION" -ge 13],
              [
-              GPUNETIO_CFLAGS="$GPUNETIO_CFLAGS -D_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE"
+              NVCCFLAGS="$NVCCFLAGS -D_LIBCUDACXX_ATOMIC_UNSAFE_AUTOMATIC_STORAGE"
              ])
 
-       save_CPPFLAGS="$CPPFLAGS"
-       CPPFLAGS="$CPPFLAGS $CUDA_CFLAGS $GPUNETIO_CFLAGS"
+       # TODO check submodule version
 
        gpunetio_happy=yes
-       AC_CHECK_HEADERS([doca_gpunetio.h], [], [gpunetio_happy=no])
-
-       CPPFLAGS="$save_CPPFLAGS"
-       LDFLAGS="$save_LDFLAGS"
       ],
       [gpunetio_happy=no])
 
 AS_IF([test "x$gpunetio_happy" = "xyes"],
       [
        uct_ib_mlx5_modules="${uct_ib_mlx5_modules}:gda"
-       AC_SUBST(GPUNETIO_CFLAGS)
-      ],
-      [
-       # gpunetio was requested but not found
-       AS_IF([test "x$with_doca_gpunetio" != "xno" -a "x$with_doca_gpunetio" != "xguess"],
-             [AC_MSG_ERROR([doca_gpunetio not found (cuda found: $cuda_happy)])])
       ])
 
 AM_CONDITIONAL([HAVE_GPUNETIO], [test x$gpunetio_happy = xyes])

src/uct/ib/mlx5/gdaki/gdaki.cuh

@@ -9,13 +9,22 @@
 
 #include "gdaki_dev.h"
 
-#include <doca_gpunetio_dev_verbs_qp.cuh>
+#include <infiniband/mlx5dv.h>
+#include <cuda.h>
+#include "device/doca_gpunetio_dev_verbs_qp.cuh"
 #include <cooperative_groups.h>
+#include <ucs/sys/device_code.h>
+#include <uct/api/uct_def.h>
 
 #define UCT_RC_GDA_RESV_WQE_NO_RESOURCE -1ULL
 #define UCT_RC_GDA_WQE_ERR              UCS_BIT(63)
 #define UCT_RC_GDA_WQE_MASK             UCS_MASK(63)
 
+UCS_F_DEVICE uint32_t uct_rc_mlx5_gda_wqe_idx_inc_mask(uint32_t wqe_idx,
+                                                       uint32_t increment)
+{
+    return (wqe_idx + increment) & DOCA_GPUNETIO_VERBS_WQE_PI_MASK;
+}
 
 UCS_F_DEVICE uct_rc_gdaki_dev_qp_t *
 uct_rc_mlx5_gda_get_qp(uct_rc_gdaki_dev_ep_t *ep, unsigned cid)
@@ -30,7 +39,7 @@ UCS_F_DEVICE void *uct_rc_mlx5_gda_get_wqe_ptr(uct_rc_gdaki_dev_ep_t *ep,
     const uintptr_t wqe_addr  = __ldg((uintptr_t*)&ep->sq_wqe_daddr);
     const uint32_t idx        = wqe_idx & (wqe_num - 1);
     const uint32_t full_idx   = idx + cid * wqe_num;
-    return (void*)(wqe_addr + (full_idx << DOCA_GPUNETIO_MLX5_WQE_SQ_SHIFT));
+    return (void*)(wqe_addr + (full_idx << DOCA_GPUNETIO_IB_MLX5_WQE_SQ_SHIFT));
 }
 
 UCS_F_DEVICE void uct_rc_mlx5_gda_ring_db(uct_rc_gdaki_dev_ep_t *ep,
@@ -320,7 +329,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_single(
     if (lane_id == 0) {
         uint16_t wqe_idx = (uint16_t)wqe_base;
         if ((comp != nullptr) || uct_rc_mlx5_gda_fc(ep, wqe_idx)) {
-            cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE;
+            cflag = DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_CQ_UPDATE;
             if (comp != nullptr) {
                 comp->wqe_idx = wqe_base;
                 comp->channel_id = cid;
@@ -423,7 +432,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi(
         return UCS_ERR_NO_RESOURCE;
     }
 
-    wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_base, lane_id);
+    wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_base, lane_id);
     for (uint32_t i = lane_id; i < count; i += num_lanes) {
         if (i == counter_index) {
             atomic         = true;
@@ -445,7 +454,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi(
         cflag = 0;
         if (((comp != nullptr) && (i == count - 1)) ||
             ((comp == nullptr) && uct_rc_mlx5_gda_fc(ep, wqe_idx))) {
-            cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE;
+            cflag = DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_CQ_UPDATE;
             if (comp != nullptr) {
                 comp->wqe_idx = wqe_base;
                 comp->channel_id = cid;
@@ -459,7 +468,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi(
                 ep, wqe_ptr, wqe_idx, opcode, cflag, remote_address, rkey,
                 reinterpret_cast<uint64_t>(address), lkey, length, atomic,
                 counter_inc_value, cid);
-        wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_idx, num_lanes);
+        wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_idx, num_lanes);
     }
 
     uct_rc_mlx5_gda_sync<level>();
@@ -517,7 +526,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial(
         return UCS_ERR_NO_RESOURCE;
     }
 
-    wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_base, lane_id);
+    wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_base, lane_id);
     for (uint32_t i = lane_id; i < count; i += num_lanes) {
         if (i == mem_list_count) {
             idx            = counter_index;
@@ -541,7 +550,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial(
         cflag = 0;
         if (((comp != nullptr) && (i == count - 1)) ||
             ((comp == nullptr) && uct_rc_mlx5_gda_fc(ep, wqe_idx))) {
-            cflag = DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE;
+            cflag = DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_CQ_UPDATE;
             if (comp != nullptr) {
                 comp->wqe_idx = wqe_base;
                 comp->channel_id = cid;
@@ -555,7 +564,7 @@ UCS_F_DEVICE ucs_status_t uct_rc_mlx5_gda_ep_put_multi_partial(
                 ep, wqe_ptr, wqe_idx, opcode, cflag, remote_address, rkey,
                 reinterpret_cast<uint64_t>(address), lkey, length, atomic,
                 counter_inc_value, cid);
-        wqe_idx = doca_gpu_dev_verbs_wqe_idx_inc_mask(wqe_idx, num_lanes);
+        wqe_idx = uct_rc_mlx5_gda_wqe_idx_inc_mask(wqe_idx, num_lanes);
     }
 
     uct_rc_mlx5_gda_sync<level>();

src/uct/ib/mlx5/gdaki/ucx-ib-mlx5-gda.pc.in

@@ -9,8 +9,8 @@ libdir = @libdir@/ucx
 Name: @PACKAGE@-ib-mlx5-gda
 Description: Unified Communication X Library GPU Direct Async module
 Version: @VERSION@
-Cflags: @GPUNETIO_CFLAGS@
+Cflags:
 Libs:
-Libs.private: -L${libdir} -luct_ib_mlx5_gda -Wl,--undefined=uct_rc_gda_init @IBVERBS_LDFLAGS@ @CUDA_LDFLAGS@ @GPUNETIO_LDFLAGS@
+Libs.private: -L${libdir} -luct_ib_mlx5_gda -Wl,--undefined=uct_rc_gda_init @IBVERBS_LDFLAGS@ @CUDA_LDFLAGS@
 Requires.private: doca-gpunetio
 

test/gtest/Makefile.am

@@ -273,9 +273,9 @@ if HAVE_NVCC
 if HAVE_GPUNETIO # temporarily all device code depends on doca
 NVCC_EXTRA_FLAGS = \
 	$(CUDA_CPPFLAGS) \
-	$(GPUNETIO_CFLAGS) \
 	-I$(top_srcdir)/test \
-	-I$(top_srcdir)/test/gtest
+	-I$(top_srcdir)/test/gtest \
+	-I$(top_srcdir)/src/uct/ib/mlx5/gdaki/gpunetio/include
 include $(top_srcdir)/config/cuda.am
 
 gtest_SOURCES += \

ucx.spec.in

@@ -75,9 +75,6 @@ BuildRequires: valgrind-devel
 %if %{with mad}
 BuildRequires: libibmad-devel libibumad-devel
 %endif
-%if %{with gda}
-BuildRequires: doca-sdk-gpunetio-devel
-%endif
 
 %description
 UCX is an optimized communication framework for high-performance distributed
@@ -438,7 +435,6 @@ Provides oneAPI Level Zero (ZE) Runtime support for UCX.
 %package ib-mlx5-gda
 Requires: %{name}-cuda%{?_isa} = %{version}-%{release}
 Requires: %{name}-ib%{?_isa} = %{version}-%{release}
-Requires: doca-sdk-gpunetio-devel
 Summary: UCX GPU Direct Async support
 Group: System Environment/Libraries