smsc/accelerator: add new smsc component

opal/mca/accelerator/rocm/accelerator_rocm_module.c

@@ -390,6 +390,12 @@ static int mca_accelerator_rocm_memcpy(int dest_dev_id, int src_dev_id, void *de
                                 "error during synchronous copy\n");
             return OPAL_ERROR;
         }
+        err = hipStreamSynchronize(0);
+        if (hipSuccess != err ) {
+            opal_output_verbose(10, opal_accelerator_base_framework.framework_output,
+                                "error synchronizing default stream after hipMemcpy\n");
+            return OPAL_ERROR;
+        }
     }
 
     return OPAL_SUCCESS;
@@ -939,4 +945,4 @@ static int mca_accelerator_rocm_get_mem_bw(int device, float *bw)
 
     *bw = opal_accelerator_rocm_mem_bw[device];
     return OPAL_SUCCESS;
-}
+}

opal/mca/btl/sm/btl_sm_component.c

@@ -340,6 +340,10 @@ mca_btl_sm_component_init(int *num_btls, bool enable_progress_threads, bool enab
         mca_btl_sm.super.btl_get = mca_btl_sm_get;
         mca_btl_sm.super.btl_put = mca_btl_sm_put;
 
+        if (mca_smsc_base_has_feature(MCA_SMSC_FEATURE_ACCELERATOR)) {
+            mca_btl_sm.super.btl_flags |= MCA_BTL_FLAGS_ACCELERATOR_GET;
+        }
+
         mca_btl_sm.super.btl_bandwidth = 40000; /* Mbs */
 
         if (mca_smsc_base_has_feature(MCA_SMSC_FEATURE_CAN_MAP)) {

opal/mca/smsc/accelerator/Makefile.am

@@ -0,0 +1,58 @@
+#
+# Copyright (c) 2004-2005 The Trustees of Indiana University and Indiana
+#                         University Research and Technology
+#                         Corporation.  All rights reserved.
+# Copyright (c) 2004-2009 The University of Tennessee and The University
+#                         of Tennessee Research Foundation.  All rights
+#                         reserved.
+# Copyright (c) 2004-2009 High Performance Computing Center Stuttgart,
+#                         University of Stuttgart.  All rights reserved.
+# Copyright (c) 2004-2005 The Regents of the University of California.
+#                         All rights reserved.
+# Copyright (c) 2009-2014 Cisco Systems, Inc.  All rights reserved.
+# Copyright (c) 2011-2014 Los Alamos National Security, LLC. All rights
+#                         reserved.
+# Copyright (c) 2017      IBM Corporation.  All rights reserved.
+# Copyright (c) 2020-2021 Google, LLC. All rights reserved.
+# Copyright (c) 2024      Advanced Micro Devices, Inc. All Rights reserved.
+# $COPYRIGHT$
+#
+# Additional copyrights may follow
+#
+# $HEADER$
+#
+
+EXTRA_DIST = post_configure.sh
+
+AM_CPPFLAGS = $(smsc_accelerator_CPPFLAGS)
+
+
+libmca_smsc_accelerator_la_sources = \
+    smsc_accelerator_component.c \
+    smsc_accelerator_module.c \
+    smsc_accelerator_internal.h \
+    smsc_accelerator.h
+
+# Make the output library in this directory, and name it either
+# mca_<type>_<name>.la (for DSO builds) or libmca_<type>_<name>.la
+# (for static builds).
+
+if MCA_BUILD_opal_smsc_accelerator_DSO
+component_noinst =
+component_install = mca_smsc_accelerator.la
+else
+component_noinst = libmca_smsc_accelerator.la
+component_install =
+endif
+
+mcacomponentdir = $(opallibdir)
+mcacomponent_LTLIBRARIES = $(component_install)
+mca_smsc_accelerator_la_SOURCES = $(libmca_smsc_accelerator_la_sources)
+mca_smsc_accelerator_la_LDFLAGS = -module -avoid-version $(smsc_accelerator_LDFLAGS)
+mca_smsc_accelerator_la_LIBADD = $(top_builddir)/opal/lib@[email protected] \
+	$(smsc_accelerator_LIBS)
+
+noinst_LTLIBRARIES = $(component_noinst)
+libmca_smsc_accelerator_la_SOURCES = $(libmca_smsc_accelerator_la_sources)
+libmca_smsc_accelerator_la_LIBADD = $(smsc_accelerator_LIBS)
+libmca_smsc_accelerator_la_LDFLAGS = -module -avoid-version $(smsc_accelerator_LDFLAGS)

opal/mca/smsc/accelerator/post_configure.sh

@@ -0,0 +1 @@
+DIRECT_CALL_HEADER="opal/mca/smsc/accelerator/smsc_accelerator.h"

opal/mca/smsc/accelerator/smsc_accelerator.h

@@ -0,0 +1,34 @@
+/* -*- Mode: C; c-basic-offset:4 ; indent-tabs-mode:nil -*- */
+/*
+ * Copyright (c) 2024      Advanced Micro Devices, Inc. All Rights reserved.
+ * $COPYRIGHT$
+ *
+ * Additional copyrights may follow
+ *
+ * $HEADER$
+ */
+
+#ifndef OPAL_MCA_SMSC_ACCELERATOR_H
+#define OPAL_MCA_SMSC_ACCELERATOR_H
+
+#include "opal_config.h"
+
+#include "opal/mca/smsc/smsc.h"
+
+mca_smsc_endpoint_t *mca_smsc_accelerator_get_endpoint(opal_proc_t *peer_proc);
+void mca_smsc_accelerator_return_endpoint(mca_smsc_endpoint_t *endpoint);
+
+int mca_smsc_accelerator_copy_to(mca_smsc_endpoint_t *endpoint, void *local_address, void *remote_address,
+                           size_t size, void *reg_handle);
+
+int mca_smsc_accelerator_copy_from(mca_smsc_endpoint_t *endpoint, void *local_address,
+                             void *remote_address, size_t size, void *reg_handle);
+
+void *mca_smsc_accelerator_map_peer_region(mca_smsc_endpoint_t *endpoint, uint64_t flags,
+                                     void *remote_ptr, size_t size, void **local_ptr);
+void mca_smsc_accelerator_unmap_peer_region(void *ctx);
+
+void *mca_smsc_accelerator_register_region(void *local_address, size_t size);
+void mca_smsc_accelerator_deregister_region(void *reg_data);
+
+#endif /* OPAL_MCA_SMSC_ACCELERATOR__H */

opal/mca/smsc/accelerator/smsc_accelerator_component.c

@@ -0,0 +1,106 @@
+/* -*- Mode: C; c-basic-offset:4 ; indent-tabs-mode:nil -*- */
+/*
+ * Copyright (c) 2024      Advanced Micro Devices, Inc. All Rights reserved.
+ * $COPYRIGHT$
+ *
+ * Additional copyrights may follow
+ *
+ * $HEADER$
+ */
+#include "opal_config.h"
+
+#include "opal/mca/smsc/base/base.h"
+#include "opal/mca/smsc/accelerator/smsc_accelerator_internal.h"
+#include "opal/mca/accelerator/accelerator.h"
+#include "opal/mca/accelerator/base/base.h"
+
+static int mca_smsc_accelerator_component_register(void);
+static int mca_smsc_accelerator_component_open(void);
+static int mca_smsc_accelerator_component_close(void);
+static int mca_smsc_accelerator_component_query(void);
+static mca_smsc_module_t *mca_smsc_accelerator_component_enable(void);
+
+#define MCA_SMSC_ACCELERATOR_DEFAULT_PRIORITY 0
+static const int mca_smsc_accelerator_default_priority = MCA_SMSC_ACCELERATOR_DEFAULT_PRIORITY;
+
+mca_smsc_accelerator_component_t mca_smsc_accelerator_component = {
+    .super = {
+        .smsc_version = {
+            MCA_SMSC_DEFAULT_VERSION("accelerator"),
+            .mca_open_component = mca_smsc_accelerator_component_open,
+            .mca_close_component = mca_smsc_accelerator_component_close,
+            .mca_register_component_params = mca_smsc_accelerator_component_register,
+        },
+        .priority = MCA_SMSC_ACCELERATOR_DEFAULT_PRIORITY,
+        .query = mca_smsc_accelerator_component_query,
+        .enable = mca_smsc_accelerator_component_enable,
+    },
+};
+
+static int mca_smsc_accelerator_component_register(void)
+{
+    mca_smsc_base_register_default_params(&mca_smsc_accelerator_component.super,
+                                          mca_smsc_accelerator_default_priority);
+    return OPAL_SUCCESS;
+}
+
+static int mca_smsc_accelerator_component_open(void)
+{
+    return OPAL_SUCCESS;
+}
+
+static int mca_smsc_accelerator_component_close(void)
+{
+    if (mca_smsc_accelerator_module.rcache) {
+        (void) mca_rcache_base_module_destroy(mca_smsc_accelerator_module.rcache);
+        mca_smsc_accelerator_module.rcache = NULL;
+    }
+
+    return OPAL_SUCCESS;
+}
+
+static int mca_smsc_accelerator_component_query(void)
+{
+    if (0 == strcmp(opal_accelerator_base_selected_component.base_version.mca_component_name,
+                    "null")) {
+        opal_output_verbose(10, opal_smsc_base_framework.framework_output,
+                          "smsc:accelerator:component_query: accelerator component is null: disqualifying myself");
+        return OPAL_ERROR;
+    }
+
+    if (!opal_accelerator.is_ipc_enabled()) {
+        opal_output_verbose(10, opal_smsc_base_framework.framework_output,
+                          "smsc:accelerator:component_query: accelerator component does not have support for IPC operations: disqualifying myself");
+        return OPAL_ERROR;
+    }
+
+    if (1 > mca_smsc_accelerator_component.super.priority) {
+        return OPAL_ERROR;
+    }
+
+    return OPAL_SUCCESS;
+}
+
+static mca_smsc_module_t *mca_smsc_accelerator_component_enable(void)
+{
+    if (0 > mca_smsc_accelerator_component.super.priority) {
+        return NULL;
+    }
+
+    mca_rcache_base_resources_t rcache_resources;
+    mca_smsc_accelerator_module.rcache = mca_rcache_base_module_create("gpusm", NULL, &rcache_resources);
+    if (NULL == mca_smsc_accelerator_module.rcache) {
+        return NULL;
+    }
+
+    /* Not set. Will initialize later */
+    mca_smsc_accelerator_module.device_id = MCA_ACCELERATOR_NO_DEVICE_ID;
+
+    mca_smsc_accelerator_module.prev_smsc = mca_smsc;
+    if ((NULL != mca_smsc_accelerator_module.prev_smsc) &&
+        (mca_smsc_accelerator_module.prev_smsc->features & MCA_SMSC_FEATURE_CAN_MAP)) {
+            mca_smsc_accelerator_module.super.features |= MCA_SMSC_FEATURE_CAN_MAP;
+    }
+
+    return &mca_smsc_accelerator_module.super;
+}

opal/mca/smsc/accelerator/smsc_accelerator_internal.h

@@ -0,0 +1,69 @@
+/* -*- Mode: C; c-basic-offset:4 ; indent-tabs-mode:nil -*- */
+/*
+ * Copyright (c) 2021      Google, Inc. All rights reserved.
+ * Copyright (c) 2024      Advanced Micro Devices, Inc. All Rights reserved.
+ * $COPYRIGHT$
+ *
+ * Additional copyrights may follow
+ *
+ * $HEADER$
+ */
+
+#ifndef OPAL_MCA_SMSC_ACCELERATOR_INTERNAL_H
+#define OPAL_MCA_SMSC_ACCELERATOR_INTERNAL_H
+
+#include "opal_config.h"
+
+#include "opal/mca/rcache/base/base.h"
+#include "opal/mca/rcache/rcache.h"
+#include "opal/mca/smsc/accelerator/smsc_accelerator.h"
+#include "opal/mca/accelerator/accelerator.h"
+#include "opal/include/opal/opal_gpu.h"
+
+#define SMSC_ACCELERATOR_HANDLE_SIZE IPC_MAX_HANDLE_SIZE
+struct mca_smsc_accelerator_registration_data_t {
+    uint64_t base_addr;
+    union {
+        uint8_t accelerator[SMSC_ACCELERATOR_HANDLE_SIZE];
+        void* host;
+    } handle;
+};
+typedef struct mca_smsc_accelerator_registration_data_t mca_smsc_accelerator_registration_data_t;
+
+struct mca_smsc_accelerator_registration_handle_t {
+    mca_opal_gpu_reg_t *gpu_reg;                   // contains mca_rcache_base_registration_t base
+    mca_smsc_accelerator_registration_data_t data;
+};
+typedef struct mca_smsc_accelerator_registration_handle_t mca_smsc_accelerator_registration_handle_t;
+OBJ_CLASS_DECLARATION(mca_smsc_accelerator_registration_handle_t);
+
+#define MCA_SMSC_ACCELERATOR_REG_DATA_TO_HANDLE(data_ptr)                                            \
+    ((mca_smsc_accelerator_registration_handle_t *) ((uintptr_t) data_ptr                            \
+                                              - offsetof(mca_smsc_accelerator_registration_handle_t, \
+                                                         data)))
+
+struct mca_smsc_accelerator_endpoint_t {
+    mca_smsc_endpoint_t super;
+    mca_smsc_endpoint_t *prev_endpoint;
+    mca_rcache_base_module_t *rcache;
+};
+typedef struct mca_smsc_accelerator_endpoint_t mca_smsc_accelerator_endpoint_t;
+OBJ_CLASS_DECLARATION(mca_smsc_accelerator_endpoint_t);
+
+struct mca_smsc_accelerator_component_t {
+    mca_smsc_component_t super;
+};
+typedef struct mca_smsc_accelerator_component_t mca_smsc_accelerator_component_t;
+
+struct mca_smsc_accelerator_module_t {
+    mca_smsc_module_t super;
+    mca_smsc_module_t *prev_smsc;
+    mca_rcache_base_module_t *rcache;
+    int device_id;
+};
+typedef struct mca_smsc_accelerator_module_t mca_smsc_accelerator_module_t;
+
+extern mca_smsc_accelerator_module_t mca_smsc_accelerator_module;
+extern mca_smsc_accelerator_component_t mca_smsc_accelerator_component;
+
+#endif /* OPAL_MCA_SMSC_ACCELERATOR_INTERNAL_H */