[OpenMP] fix issues defining external emissary APIs by defining weak references for EmissaryMPI, EmissaryReserved, and EmissaryHDF5

gregrodgers · gregrodgers · commit 3b134218629d · 2025-03-03T14:21:25.000-06:00
diff --git a/offload/CMakeLists.txt b/offload/CMakeLists.txt
@@ -436,6 +436,8 @@ if(OFFLOAD_ENABLE_EMISSARY_APIS)
   add_definitions(-DOFFLOAD_ENABLE_EMISSARY_APIS)
   install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/DeviceRTL/include/EmissaryIds.h
     DESTINATION ${CMAKE_INSTALL_PREFIX}/lib/clang/${LLVM_VERSION_MAJOR}/include)
+  install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/DeviceRTL/include/EmissaryMPI.h
+    DESTINATION ${CMAKE_INSTALL_PREFIX}/lib/clang/${LLVM_VERSION_MAJOR}/include)
 endif()
 
 # Build offloading plugins and device RTLs if they are available.
diff --git a/offload/DeviceRTL/include/EmissaryIds.h b/offload/DeviceRTL/include/EmissaryIds.h
@@ -20,6 +20,7 @@ typedef enum {
   EMIS_ID_PRINT,
   EMIS_ID_MPI,
   EMIS_ID_HDF5,
+  EMIS_ID_RESERVE,
 } offload_emis_id_t;
 
 typedef enum {
@@ -29,12 +30,6 @@ typedef enum {
   _ockl_asan_report_idx,
 } offload_emis_print_t;
 
-typedef enum {
-  _MPI_INVALID,
-  _MPI_Send_idx,
-  _MPI_Recv_idx,
-} offload_emis_mpi_t;
-
 /// The vargs function used by emissary API device stubs
 unsigned long long _emissary_exec(unsigned long long, ...);
 
@@ -59,28 +54,23 @@ typedef enum {
   _FortranAStopStatementText_idx,
 } offload_emis_fortrt_idx;
 
-//  mpi.h (needed for MPI types) will not compile while building DeviceRTL,
-//  So emissary stubs for MPI functions can NOT be in libomptarget.bc.
-//  These are skipped whild building DeviceRTL because compilation of DeviceRTL
-//  does not have include mpi.h. The user will build these stubs on their
-//  device pass when they include EmissaryIds.h.
+/// This structure is created by emisExtractArgBuf to make it easier
+/// to get values from the data buffer passed by rpc.
+typedef struct {
+  unsigned int DataLen;
+  unsigned int NumArgs;
+  unsigned int emisid;
+  unsigned int emisfnid;
+  unsigned long long data_not_used;
+  char *keyptr;
+  char *argptr;
+  char *strptr;
+} emisArgBuf_t;
+
+typedef unsigned long long emis_return_t;
+typedef unsigned long long emis_argptr_t;
+typedef emis_return_t emisfn_t(void *, ...);
 
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-#if defined(__has_include)
-#if __has_include("mpi.h")
-#include "mpi.h"
-extern "C" int MPI_Send(const void *buf, int count, MPI_Datatype datatype,
-                        int dest, int tag, MPI_Comm comm) {
-  return (int)_emissary_exec(_PACK_EMIS_IDS(EMIS_ID_MPI, _MPI_Send_idx), buf,
-                             count, datatype, dest, tag, comm);
-}
-extern "C" int MPI_Recv(void *buf, int count, MPI_Datatype datatype, int source,
-                        int tag, MPI_Comm comm, MPI_Status *st) {
-  return (int)_emissary_exec(_PACK_EMIS_IDS(EMIS_ID_MPI, _MPI_Recv_idx), buf,
-                             count, datatype, source, tag, comm, st);
-}
-#endif
-#endif
-#endif
+#define MAXVARGS 32
 
 #endif // OFFLOAD_EMISSARY_IDS_H
diff --git a/offload/DeviceRTL/include/EmissaryMPI.h b/offload/DeviceRTL/include/EmissaryMPI.h
@@ -1,37 +1,38 @@
-//===---- offload/plugins-nextgen/common/src/EmissaryFortrt.cpp  ----------===//
+//===--------------- offload/DeviceRTL/include/EmissaryMPI.h --------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
-// Host support for Fortran runtime Emissary API
+// EmissaryMPI.h This include must be included by MPI application
 //
 //===----------------------------------------------------------------------===//
-#include "PluginInterface.h"
-#include "RPC.h"
-#include "Shared/Debug.h"
-#include "Shared/RPCOpcodes.h"
-#include "shared/rpc.h"
-#include "shared/rpc_opcodes.h"
-#include "../../../DeviceRTL/include/EmissaryIds.h"
-#include "Emissary.h"
-#include <assert.h>
-#include <cstring>
-#include <ctype.h>
-#include <list>
-#include <mpi.h>
+#include "EmissaryIds.h"
 #include <stdarg.h>
-#include <stdbool.h>
-#include <stddef.h>
-#include <stdint.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <tuple>
-#include <vector>
 
+typedef enum {
+  _MPI_INVALID,
+  _MPI_Send_idx,
+  _MPI_Recv_idx,
+} offload_emis_mpi_t;
+
+///  Device stubs that call _emissary_exec using identical host API interface
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+extern "C" int MPI_Send(const void *buf, int count, MPI_Datatype datatype,
+                        int dest, int tag, MPI_Comm comm) {
+  return (int)_emissary_exec(_PACK_EMIS_IDS(EMIS_ID_MPI, _MPI_Send_idx), buf,
+                             count, datatype, dest, tag, comm);
+}
+extern "C" int MPI_Recv(void *buf, int count, MPI_Datatype datatype, int source,
+                        int tag, MPI_Comm comm, MPI_Status *st) {
+  return (int)_emissary_exec(_PACK_EMIS_IDS(EMIS_ID_MPI, _MPI_Recv_idx), buf,
+                             count, datatype, source, tag, comm, st);
+}
+#endif
+
+/// Host variadic wrapper functions.
 extern "C" {
 extern int V_MPI_Send(void *fnptr, ...) {
   va_list args;
@@ -61,11 +62,8 @@ extern int V_MPI_Recv(void *fnptr, ...) {
   return rval;
 }
 
-emis_return_t EmissaryMPI(char *data, emisArgBuf_t *ab) {
-  uint64_t *a[MAXVARGS];
-  if (EmissaryBuildVargs(ab->NumArgs, ab->keyptr, ab->argptr, ab->strptr,
-                         &ab->data_not_used, &a[0]) != _RC_SUCCESS)
-    return (emis_return_t)0;
+/// EmissaryMPI function selector
+emis_return_t EmissaryMPI(char *data, emisArgBuf_t *ab, emis_argptr_t *a[]) {
 
   switch (ab->emisfnid) {
   case _MPI_Send_idx: {
diff --git a/offload/libomptarget/exports b/offload/libomptarget/exports
@@ -97,6 +97,10 @@ VERS1.0 {
     __llvmPushCallConfiguration;
     __llvmPopCallConfiguration;
     llvmLaunchKernel;
+    EmissaryBuildVargs;
+    EmissaryHDF5;
+    EmissaryReserve;
+    EmissaryMPI;
   local:
     *;
 };
diff --git a/offload/plugins-nextgen/common/CMakeLists.txt b/offload/plugins-nextgen/common/CMakeLists.txt
@@ -5,13 +5,6 @@ if(OFFLOAD_ENABLE_EMISSARY_APIS)
     src/EmissaryFortrt.cpp
     src/EmissaryPrint.cpp
   )
-  set(OFFLOAD_EMISSARY_MPI_INCLUDE ""
-	  CACHE STRING "MPI include directory for building MPI Emissary API")
-  # dont enable the MPI Emissary api unless we have an MPI include dir
-  # that contains mpi.h needed to build the variadic wrappers.
-  if(OFFLOAD_EMISSARY_MPI_INCLUDE)
-    list(APPEND emissary_sources src/EmissaryMPI.cpp)
-  endif()
 endif()
 
 # NOTE: Don't try to build `PluginInterface` using `add_llvm_library` because we
@@ -44,9 +37,6 @@ target_link_libraries(PluginCommon PRIVATE llvm-libc-common-utilities)
 if(OFFLOAD_ENABLE_EMISSARY_APIS)
   target_link_libraries(PluginCommon PRIVATE flang_rt.runtime
   -L${CMAKE_BINARY_DIR}/../../lib  -L${CMAKE_INSTALL_PREFIX}/lib)
-  if(OFFLOAD_EMISSARY_MPI_INCLUDE)
-    target_include_directories(PluginCommon PUBLIC ${OFFLOAD_EMISSARY_MPI_INCLUDE})
-  endif()
 endif()
 if(TARGET llvmlibc_rpc_server AND ${LIBOMPTARGET_GPU_LIBC_SUPPORT})
 	target_link_libraries(PluginCommon PRIVATE llvmlibc_rpc_server)
diff --git a/offload/plugins-nextgen/common/include/Emissary.h b/offload/plugins-nextgen/common/include/Emissary.h
@@ -14,25 +14,7 @@
 #ifndef OFFLOAD_EMISSARY_H
 #define OFFLOAD_EMISSARY_H
 
-/// This structure is created by emisExtractArgBuf to make it easier
-/// to get values from the data buffer passed by rpc.
-typedef struct {
-  unsigned int DataLen;
-  unsigned int NumArgs;
-  unsigned int emisid;
-  unsigned int emisfnid;
-  size_t data_not_used;
-  char *keyptr;
-  char *argptr;
-  char *strptr;
-} emisArgBuf_t;
-
-typedef unsigned long long emis_return_t;
-typedef emis_return_t emisfn_t(void *, ...);
-
-// MAXVARGS is the maximum number of args in an emissary function
-// To increase this number, update EmissaryCallFnptr below
-#define MAXVARGS 32
+#include "../../../DeviceRTL/include/EmissaryIds.h"
 
 extern "C" {
 
@@ -46,10 +28,16 @@ emis_return_t EmissaryFortrt(char *data, emisArgBuf_t *ab);
 emis_return_t EmissaryPrint(char *data, emisArgBuf_t *ab);
 
 /// Called by Emissary for all MPI emissary API functions
-emis_return_t EmissaryMPI(char *data, emisArgBuf_t *ab);
+__attribute((weak)) emis_return_t EmissaryMPI(char *data, emisArgBuf_t *ab,
+                                              emis_argptr_t *arg[MAXVARGS]);
 
 /// Called by Emissary for all HDF5 Emissary API functions
-emis_return_t EmissaryHDF5(char *data, emisArgBuf_t *ab);
+__attribute((weak)) emis_return_t EmissaryHDF5(char *data, emisArgBuf_t *ab,
+                                               emis_argptr_t *arg[MAXVARGS]);
+
+/// Support externally supplied emissary API
+__attribute((weak)) emis_return_t EmissaryReserve(char *data, emisArgBuf_t *ab,
+                                                  emis_argptr_t *arg[MAXVARGS]);
 
 /// Called by Emissary to build the emisArgBuf_t structure from the emissary
 /// data buffer sent to the CPU by rpc. This buffer is created by clang CodeGen
@@ -66,15 +54,15 @@ void *getfnptr(char *val);
 
 /// Builds the array of pointers passed to V_ functions
 uint32_t EmissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr,
-                            char *strptr, size_t *data_not_used,
-                            uint64_t *a[MAXVARGS]);
+                            char *strptr, unsigned long long *data_not_used,
+                            emis_argptr_t *a[MAXVARGS]);
 
 } // end extern "C"
 
 /// Call the associated V_ function
 template <typename T, typename FT>
 extern T EmissaryCallFnptr(uint32_t NumArgs, void *fnptr,
-                           uint64_t *a[MAXVARGS]);
+                           emis_argptr_t *a[MAXVARGS]);
 
 // Error return codes (deprecated)
 typedef enum service_rc {
@@ -140,7 +128,7 @@ enum TypeID {
 
 template <typename T, typename FT>
 extern T EmissaryCallFnptr(uint32_t NumArgs, void *fnptr,
-                           uint64_t *a[MAXVARGS]) {
+                           emis_argptr_t *a[MAXVARGS]) {
   T rv;
   FT *vfnptr = (FT *)fnptr;
   switch (NumArgs) {
diff --git a/offload/plugins-nextgen/common/src/Emissary.cpp b/offload/plugins-nextgen/common/src/Emissary.cpp
@@ -23,6 +23,7 @@ extern "C" emis_return_t Emissary(char *data) {
   emisArgBuf_t ab;
   emisExtractArgBuf(data, &ab);
   emis_return_t result = 0;
+  emis_argptr_t *args[MAXVARGS]; // FIXME use malloc here
 
   switch (ab.emisid) {
   case EMIS_ID_INVALID: {
@@ -39,13 +40,24 @@ extern "C" emis_return_t Emissary(char *data) {
     break;
   }
   case EMIS_ID_MPI: {
-    result = EmissaryMPI(data, &ab);
+    if (EmissaryBuildVargs(ab.NumArgs, ab.keyptr, ab.argptr, ab.strptr,
+                           &ab.data_not_used, &args[0]) != _RC_SUCCESS)
+      return (emis_return_t)0;
+    result = EmissaryMPI(data, &ab, args);
     break;
   }
   case EMIS_ID_HDF5: {
-    // result = EmissaryHDF5(data, &ab);
-    result = 0;
-    fprintf(stderr, "Support for HDF5 Emissary API is in development\n");
+    if (EmissaryBuildVargs(ab.NumArgs, ab.keyptr, ab.argptr, ab.strptr,
+                           &ab.data_not_used, &args[0]) != _RC_SUCCESS)
+      return (emis_return_t)0;
+    result = EmissaryHDF5(data, &ab, args);
+    break;
+  }
+  case EMIS_ID_RESERVE: {
+    if (EmissaryBuildVargs(ab.NumArgs, ab.keyptr, ab.argptr, ab.strptr,
+                           &ab.data_not_used, &args[0]) != _RC_SUCCESS)
+      return (emis_return_t)0;
+    result = EmissaryReserve(data, &ab, args);
     break;
   }
   default:
@@ -114,8 +126,9 @@ extern "C" void *getfnptr(char *val) {
 
 // build argument array
 extern "C" uint32_t EmissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr,
-                                       char *strptr, size_t *data_not_used,
-                                       uint64_t *a[MAXVARGS]) {
+                                       char *strptr,
+                                       unsigned long long *data_not_used,
+                                       emis_argptr_t *a[MAXVARGS]) {
   size_t num_bytes;
   size_t bytes_consumed;
   size_t strsz;
@@ -145,9 +158,9 @@ extern "C" uint32_t EmissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr,
         return _RC_DATA_USED_ERROR;
 
       if (num_bytes == 4)
-        a[argcount] = (uint64_t *)getuint32(dataptr);
+        a[argcount] = (emis_argptr_t *)getuint32(dataptr);
       else
-        a[argcount] = (uint64_t *)getuint64(dataptr);
+        a[argcount] = (emis_argptr_t *)getuint64(dataptr);
 
       break;
 
@@ -163,9 +176,9 @@ extern "C" uint32_t EmissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr,
         return _RC_DATA_USED_ERROR;
 
       if (num_bytes == 4)
-        a[argcount] = (uint64_t *)getuint32(dataptr);
+        a[argcount] = (emis_argptr_t *)getuint32(dataptr);
       else
-        a[argcount] = (uint64_t *)getuint64(dataptr);
+        a[argcount] = (emis_argptr_t *)getuint64(dataptr);
 
       break;
 
@@ -176,7 +189,7 @@ extern "C" uint32_t EmissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr,
         strsz = (size_t)*(unsigned int *)dataptr;
         if ((*data_not_used) < bytes_consumed)
           return _RC_DATA_USED_ERROR;
-        a[argcount] = (uint64_t *)((char *)strptr);
+        a[argcount] = (emis_argptr_t *)((char *)strptr);
 
       } else {
         num_bytes = 8;
@@ -189,7 +202,7 @@ extern "C" uint32_t EmissaryBuildVargs(int NumArgs, char *keyptr, char *dataptr,
         if ((*data_not_used) < bytes_consumed)
           return _RC_DATA_USED_ERROR;
 
-        a[argcount] = (uint64_t *)getuint64(dataptr);
+        a[argcount] = (emis_argptr_t *)getuint64(dataptr);
       }
       break;
 
diff --git a/offload/plugins-nextgen/common/src/EmissaryFortrt.cpp b/offload/plugins-nextgen/common/src/EmissaryFortrt.cpp
@@ -232,7 +232,7 @@ extern "C" emis_return_t EmissaryFortrt(char *data, emisArgBuf_t *ab) {
   if (ab->NumArgs <= 0)
     return _RC_ERROR_INVALID_REQUEST;
 
-  uint64_t *a[MAXVARGS];
+  emis_argptr_t *a[MAXVARGS];
   if (EmissaryBuildVargs(ab->NumArgs, ab->keyptr, ab->argptr, ab->strptr,
                          &ab->data_not_used, a) != _RC_SUCCESS)
     return _RC_ERROR_INVALID_REQUEST;
@@ -254,7 +254,7 @@ extern "C" emis_return_t EmissaryFortrt(char *data, emisArgBuf_t *ab) {
       fprintf(stderr, "MALLOC FAILED for c_ptr size:%ld \n", slen);
     std::strncpy(c_ptr, (char *)a[5], slen - 1);
     c_ptr[slen - 1] = (char)0;
-    a[5] = (uint64_t *)c_ptr;
+    a[5] = (emis_argptr_t *)c_ptr;
     break;
   }
   case _FortranAioOutputAscii_idx: {
@@ -266,7 +266,7 @@ extern "C" emis_return_t EmissaryFortrt(char *data, emisArgBuf_t *ab) {
       fprintf(stderr, "MALLOC FAILED for c_ptr size:%ld \n", slen);
     std::strncpy(c_ptr, (char *)a[5], slen - 1);
     c_ptr[slen - 1] = (char)0;
-    a[5] = (uint64_t *)c_ptr;
+    a[5] = (emis_argptr_t *)c_ptr;
 
     break;
   }
@@ -378,7 +378,7 @@ extern "C" emis_return_t EmissaryFortrt(char *data, emisArgBuf_t *ab) {
         for (auto q : *_deferred_fns_ptr) {
           if ((thread_num == q->thread_num) && (team_num == q->team_num)) {
             for (uint32_t i = 0; i < q->NumArgs; i++)
-              a[i] = (uint64_t *)q->arg_array[i];
+              a[i] = (emis_argptr_t *)q->arg_array[i];
             q->return_value = EmissaryCallFnptr<emis_return_t, emisfn_t>(
                 q->NumArgs, q->fnptr, a);
           }
diff --git a/offload/plugins-nextgen/common/src/EmissaryPrint.cpp b/offload/plugins-nextgen/common/src/EmissaryPrint.cpp
@@ -212,7 +212,7 @@ static service_rc emissary_pfAddFloat(emissary_ValistExt_t *valist,
 static service_rc emissary_pfBuildValist(emissary_ValistExt_t *valist,
                                          int NumArgs, char *keyptr,
                                          char *dataptr, char *strptr,
-                                         size_t *data_not_used) {
+                                         unsigned long long *data_not_used) {
   emissary_pfRegSaveArea_t *regs;
   size_t regs_size = sizeof(*regs);
   regs = (emissary_pfRegSaveArea_t *)malloc(regs_size);
