Runtime for Interop directive

[sriharikrishna]

Summary

This implements the runtime portion of the interop directive.
It expects the frontend and IRBuilder portions to be in place
for proper execution. It currently works only for GPUs
and has several TODOs that should be addressed going forward.

[RaviNarayanaswamy]

You should not be using kmpc in libomptarget

[AndreyChurbanov]

Looks like openmp/runtime/src/dllexports is edited by mistake. Added symbols are defined nowhere that causes build failure on Windows.

[tianshilei1992]

openmp/libomptarget/include/interop.h

20

Probably not an option.

openmp/libomptarget/include/omptarget.h

21

Why is that?

openmp/libomptarget/plugins/cuda/src/rtl.cpp

78

I think these series of changes regarding SPMD_GENERIC has nothing to do with interop. Please make a separate patch for them.

1137–1138

openmp/libomptarget/src/CMakeLists.txt

21

order please

openmp/libomptarget/src/exports

49

Like @RaviNarayanaswamy said, better to rename them to __tgt_xxx.

openmp/libomptarget/src/interop.cpp

14

There are some conventions in current libomptarget.

1. If a function is internal, use LLVM code style.
2. If a function is exported and exposed to compiler, it should be extern "C" and use code style similar to existing functions whose name prefix with __tgt.

So basically, if these functions are only visible to this file, please format it with LLVM code style, and use anonymous name space.

openmp/libomptarget/src/private.h

117

If this function is not defined in libomp, please don't use __kmpc.

[jdoerfert]

There are a lot of minor things that should be easily addressable.

openmp/libomptarget/include/interop.h

2

We need the llvm license header, copied from another file.

66

put the entire file content (after macros) into a single extern C.

110

The functions above need definitions in libomp and libomptarget.
For the former we need to edit the following two files, look for "get_device_num" (or "GET_DEVICE_NUM") in both
to see how to define it weak:
openmp/runtime/src/kmp_ftn_os.h
openmp/runtime/src/kmp_ftn_entry.h
just return 0 or nullptr for all of them in libomp.

In libomptarget already have the definitions and exports.

161

Remove the asserts.

openmp/libomptarget/plugins/cuda/src/rtl.cpp

78

This is a rebase problem. These changes are already upstream but somehow they managed to get into this commit/revision. Rebase locally and this should hopefully be gone.

openmp/libomptarget/src/exports

1

undo this.

openmp/libomptarget/src/interop.cpp

46

We nee a switch and return all values in Table 2.1 (https://www.openmp.org/wp-content/uploads/OpenMP-API-Additional-Definitions-2-0.pdf) or "unknown" as default.

openmp/libomptarget/src/private.h

117

It is in libomp, this is a forward decl. All good here.

openmp/runtime/src/dllexports

582

Maybe remove all these for now. That should also allow to not define omp_get_interop_XXX in libomp for now and it will just fail if we do not link libomptarget as a first step.

[jdoerfert]

openmp/libomptarget/src/interop.cpp

178

don't use references in this code with C bindings. Use a ** instead and change left-hand-side uses to *interop_ptr. This will then also match D105876 where all 3 functions are declared to take a void**. That means _use below should take a void** as well.

192

This assertion seems nonsensical, just return I guess.

205

Move the wait stuff to the beginning.

217

Again
interop_ptr = omp_interop_none;

[tianshilei1992]

openmp/libomptarget/src/private.h

117

Is it necessary to implement it in libomp? interop is for target.

[jdoerfert]

openmp/libomptarget/src/private.h

117

eventually it is. if you don't add -fopenmp-targets you would otherwise get link errors, which sounds pretty bad.

[tianshilei1992]

openmp/libomptarget/src/private.h

117

Well, we can implement a wrapper in libomp, just like omp_get_num_devices.

// Get number of NON-HOST devices.
// libomptarget, if loaded, provides this function in api.cpp.
int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)(void)
    KMP_WEAK_ATTRIBUTE_EXTERNAL;
int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)(void) {
#if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
  return 0;
#else
  int (*fptr)();
  if ((*(void **)(&fptr) = KMP_DLSYM("__tgt_get_num_devices"))) {
    return (*fptr)();
  } else if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_num_devices"))) {
    return (*fptr)();
  } else if ((*(void **)(&fptr) = KMP_DLSYM("_Offload_number_of_devices"))) {
    return (*fptr)();
  } else { // liboffload & libomptarget don't exist
    return 0;
  }
#endif // KMP_MIC || KMP_OS_DARWIN || KMP_OS_WINDOWS || defined(KMP_STUB)
}

[sriharikrishna]

openmp/libomptarget/include/interop.h

20

Okay. Any suggestions on how to overcome this?

110

Okay. While linking, if we use the order -lomptarget -lomp then we get the values from libomptarget. If we use the order -lomp -lomptarget we get the values from libomp. This seems fragile.

openmp/libomptarget/include/omptarget.h

21

Undid it.

openmp/libomptarget/src/interop.cpp

178

It is not clear how to fix this. In the user code's interop directive, the interop variable must be of type 'omp_interop_t'. So getting

[sriharikrishna]

Runtime for Interop directive

[sriharikrishna]

openmp/libomptarget/src/interop.cpp

178

Incorrect comment above. This is fixed.

[jdoerfert]

@AndreyChurbanov @tianshilei1992 Any concerns with merging this today if the two things below are addressed?

openmp/runtime/src/dllexports

553

Those values are taken by now, you need new values that are not taken yet.

openmp/runtime/src/kmp_ftn_entry.h

1546

We need the KMP_DLSYM_NEXT lookup in these, probably sufficient to do only that lookup though.

[tianshilei1992]

There are still many open comments.

[sriharikrishna]

Runtime for Interop directiv

[sriharikrishna]

I have fixed the latest comments.

[sriharikrishna]

Runtime for Interop directive

[sriharikrishna]

Runtime for Interop directive

[sriharikrishna]

Runtime for Interop directive

[tianshilei1992]

openmp/libomptarget/src/interop.cpp

14

I mean, this function doesn't have to start with __tgt because it is internal. Functions starting with __tgt are usually interface functions. From my perspective, it is better to name it as getPropertyErrorType, a.k.a. to use LLVM code style.

64

Same here

openmp/runtime/src/kmp_ftn_entry.h

1501

ditto

1515

ditto

1527

ditto

1539

ditto

1550

same as below

1561

same as below

1570–1574

[sriharikrishna]

Runtime for Interop directive

[tianshilei1992]

openmp/libomptarget/plugins/cuda/src/rtl.cpp

1150–1156

openmp/libomptarget/src/interop.cpp

14

Looks better. Can you check https://llvm.org/docs/CodingStandards.html for the code style?

44

Does it make sense to use enum here?

[RaviNarayanaswamy]

openmp/libomptarget/src/interop.cpp

195–197

Need to check if device is available for all 3 __tgt_interop_init/use/destroy

199–202

Interop object does not wait for any task from libomp.
Init just initializes the interop object.
The initialization of interop object should be done by the plugin as only the plugin knows what properties are supported

240–243

Need to flush the queue if interop object was created with targetsync

261–264

You don't wait for the omp tasks. Need to flush the queue associated with the interop through the plugin

openmp/runtime/src/kmp_ftn_entry.h

1449–1494

Why do you have all this in openmp/runtime. Openmp should call libomptarget to get interop properties. if libomptarget is not loaded it should return 0

[jdoerfert]

We need a runtime test.

Synchronizing the asyn info object and signal out dependences will happen in the next revision.
So will minor edits as requested.

I commented below on some things.

openmp/libomptarget/src/interop.cpp

14

It's unclear given the mixed nature. Let's go with this and figure out if we need to rename vars later or not.

199–202

Interop object does not wait for any task from libomp.

I don't know why you think we would not wait for libomp tasks. If we have dependences we need to wait for them.

The initialization of interop object should be done by the plugin as only the plugin knows what properties are supported.

It is, below. This is the generic part that then redirects to the plugin.

261–264

Waiting for omp task is necessary and the above should do it. Flushing and signaling out dependences will be added in the next revision.

openmp/runtime/src/kmp_ftn_entry.h

1449–1494

That is exactly what this code does, no? Look for libomptarget and redirect to that one, otherwise return 0.
Unsure I see your point. FWIW, this is copied from other functions we duplicate in libomp but that actually are part of libomptarget.

[RaviNarayanaswamy]

openmp/libomptarget/src/interop.cpp

199–202

Libomp would have not invoked the task which calls this routine if there are dependences. They must be executed before the task containing the interop creation is scheduled.

The interop_type should be passed to plugin and let it initialize the common for all interop-types and then add based on the interop_type

[jdoerfert]

openmp/libomptarget/src/interop.cpp

199–202

Libomp would have not invoked the task which calls this routine if there are dependences. They must be executed before the task containing the interop creation is scheduled.

To me it seems you are assuming that we create a task in which this routine is called. We do not, as far as I can tell. See D105876.

The interop_type should be passed to plugin and let it initialize the common for all interop-types and then add based on the interop_type

So what you are trying to say is that init_device_info should take the interop_type too? That makes sense to me. But as discussed in other reviews recently, we should not extend the API for "future use cases" but extend it as use cases become relevant. For now it seems we can simply set up the tgt_device_info part of the omp_interop_val_t without knowing the interop_type.

[RaviNarayanaswamy]

openmp/libomptarget/src/interop.cpp

199–202

Then need to change this code since the interop_type can be both target_sync and target which will not be handled correctly. target_sync and target have common initialization + additional property base don the interop_type requested

[jdoerfert]

openmp/libomptarget/src/interop.cpp

199–202

Then need to change this code since the interop_type can be both target_sync and target which will not be handled correctly. target_sync and target have common initialization + additional property base don the interop_type requested

Could you please elaborate what needs to be changed exactly. What information is currently not available in the setup as is? What properties would be different?

[RaviNarayanaswamy]

openmp/libomptarget/src/interop.cpp

199–202

Should be something like this

NEED to add _kmp_interop_type_target to represent interop target
interop_ptr->device_info would initialize the following: device handle, device_context, platform.
if (interop_type == kmp_interop_type_target) {
if (!Device.RTL || !Device.RTL->init_device_info ||

  Device.RTL->init_device_info(device_id, &(interop_ptr)->device_info,
                               &(interop_ptr)->err_str)) {
delete interop_ptr;
interop_ptr = omp_interop_none;

}
// Add target sync if request.
if (interop_type == kmp_interop_type_tasksync) {

  if (!Device.RTL || !Device.RTL->init_async_info ||
      Device.RTL->init_async_info(device_id, &(interop_ptr)->async_info)) {
    delete interop_ptr;
    interop_ptr = omp_interop_none;
}

[sriharikrishna]

openmp/libomptarget/src/interop.cpp

199–202

It looks like the code you have written for kmp_interop_type_target already exists in the else part. So I am still puzzled about this.

[sriharikrishna]

Runtime for Interop directive

[sriharikrishna]

Runtime for Interop directive

[sriharikrishna]

openmp/libomptarget/src/interop.cpp

240–243

A TODO has been left because flushing is not supported.

261–264

A TODOs been left because flushing is not supported.

openmp/runtime/src/kmp_ftn_entry.h

1449–1494

Im leaving it as is.

[RaviNarayanaswamy]

LGTM

[tianshilei1992]

rebase

[mstorsjo]

openmp/runtime/src/dllexports

553

This broke building for Windows:

generate-def.pl: (x) Error parsing file "llvm-project/openmp/runtime/src/dllexports" line 556:
generate-def.pl: (x)         omp_get_interop_int                     2514
generate-def.pl: (x) Ordinal of user-callable entry must be < 1000

[mstorsjo]

openmp/runtime/src/dllexports

553

... and even with those ordinals changed to something in the right range, the build later fails with linker errors:

ld.lld: error: <root>: undefined symbol: OMP_GET_INTEROP_INT
ld.lld: error: <root>: undefined symbol: OMP_GET_INTEROP_PTR
ld.lld: error: <root>: undefined symbol: OMP_GET_INTEROP_STR

So please fix the Windows build of OpenMP, before the 14.x branch is made early next week.

[AndreyChurbanov]

openmp/runtime/src/dllexports

553

As the warning says these should be < 1000, because lowercase are automatically replicated in uppercase adding +1000 to the ordinal and still be < 2000 (for some particular Fortran usages).
There are free ordinals below 1000 starting at 807. So, e.g. 807, 808, etc. should work fine.

Actually these interfaces described as "C/C++"-only in the specification, so there is no need to provide Fortran-specific versions, though it should not harm. Probably making them C-only does not worth efforts for now.

openmp/libomptarget/include/interop.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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _INTEROP_H_
+#define _INTEROP_H_
+
+#include "omptarget.h"
+#include <assert.h>
+
+#if defined(_WIN32)
+#define __KAI_KMPC_CONVENTION __cdecl
+#ifndef __KMP_IMP
+#define __KMP_IMP __declspec(dllimport)
+#endif
+#else
+#define __KAI_KMPC_CONVENTION
+#ifndef __KMP_IMP
+#define __KMP_IMP
+#endif
+#endif
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/// TODO: Include the `omp.h` of the current build
+/* OpenMP 5.1 interop */
+typedef intptr_t omp_intptr_t;
+
+/* 0..omp_get_num_interop_properties()-1 are reserved for implementation-defined
+ * properties */
+typedef enum omp_interop_property {
+  omp_ipr_fr_id = -1,
+  omp_ipr_fr_name = -2,
+  omp_ipr_vendor = -3,
+  omp_ipr_vendor_name = -4,
+  omp_ipr_device_num = -5,
+  omp_ipr_platform = -6,
+  omp_ipr_device = -7,
+  omp_ipr_device_context = -8,
+  omp_ipr_targetsync = -9,
+  omp_ipr_first = -9
+} omp_interop_property_t;
+
+#define omp_interop_none 0
+
+typedef enum omp_interop_rc {
+  omp_irc_no_value = 1,
+  omp_irc_success = 0,
+  omp_irc_empty = -1,
+  omp_irc_out_of_range = -2,
+  omp_irc_type_int = -3,
+  omp_irc_type_ptr = -4,
+  omp_irc_type_str = -5,
+  omp_irc_other = -6
+} omp_interop_rc_t;
+
+typedef enum omp_interop_fr {
+  omp_ifr_cuda = 1,
+  omp_ifr_cuda_driver = 2,
+  omp_ifr_opencl = 3,
+  omp_ifr_sycl = 4,
+  omp_ifr_hip = 5,
+  omp_ifr_level_zero = 6,
+  omp_ifr_last = 7
+} omp_interop_fr_t;
+
+typedef void *omp_interop_t;
+
+/*!
+ * The `omp_get_num_interop_properties` routine retrieves the number of
+ * implementation-defined properties available for an `omp_interop_t` object.
+ */
+int __KAI_KMPC_CONVENTION omp_get_num_interop_properties(const omp_interop_t);
+/*!
+ * The `omp_get_interop_int` routine retrieves an integer property from an
+ * `omp_interop_t` object.
+ */
+omp_intptr_t __KAI_KMPC_CONVENTION omp_get_interop_int(const omp_interop_t,
+                                                       omp_interop_property_t,
+                                                       int *);
+/*!
+ * The `omp_get_interop_ptr` routine retrieves a pointer property from an
+ * `omp_interop_t` object.
+ */
+void *__KAI_KMPC_CONVENTION omp_get_interop_ptr(const omp_interop_t,
+                                                omp_interop_property_t, int *);
+/*!
+ * The `omp_get_interop_str` routine retrieves a string property from an
+ * `omp_interop_t` object.
+ */
+const char *__KAI_KMPC_CONVENTION omp_get_interop_str(const omp_interop_t,
+                                                      omp_interop_property_t,
+                                                      int *);
+/*!
+ * The `omp_get_interop_name` routine retrieves a property name from an
+ * `omp_interop_t` object.
+ */
+const char *__KAI_KMPC_CONVENTION omp_get_interop_name(const omp_interop_t,
+                                                       omp_interop_property_t);
+/*!
+ * The `omp_get_interop_type_desc` routine retrieves a description of the type
+ * of a property associated with an `omp_interop_t` object.
+ */
+const char *__KAI_KMPC_CONVENTION
+omp_get_interop_type_desc(const omp_interop_t, omp_interop_property_t);
+/*!
+ * The `omp_get_interop_rc_desc` routine retrieves a description of the return
+ * code associated with an `omp_interop_t` object.
+ */
+extern const char *__KAI_KMPC_CONVENTION
+omp_get_interop_rc_desc(const omp_interop_t, omp_interop_rc_t);
+
+typedef struct kmp_tasking_flags { /* Total struct must be exactly 32 bits */
+  /* Compiler flags */             /* Total compiler flags must be 16 bits */
+  unsigned tiedness : 1;           /* task is either tied (1) or untied (0) */
+  unsigned final : 1;              /* task is final(1) so execute immediately */
+  unsigned merged_if0 : 1; // no __kmpc_task_{begin/complete}_if0 calls in if0
+  unsigned destructors_thunk : 1; // set if the compiler creates a thunk to
+  unsigned proxy : 1; // task is a proxy task (it will be executed outside the
+  unsigned priority_specified : 1; // set if the compiler provides priority
+  unsigned detachable : 1;         // 1 == can detach */
+  unsigned unshackled : 1;         /* 1 == unshackled task */
+  unsigned target : 1;             /* 1 == target task */
+  unsigned reserved : 7;           /* reserved for compiler use */
+  unsigned tasktype : 1;    /* task is either explicit(1) or implicit (0) */
+  unsigned task_serial : 1; // task is executed immediately (1) or deferred (0)
+  unsigned tasking_ser : 1; // all tasks in team are either executed immediately
+  unsigned team_serial : 1; // entire team is serial (1) [1 thread] or parallel
+  unsigned started : 1;     /* 1==started, 0==not started     */
+  unsigned executing : 1;   /* 1==executing, 0==not executing */
+  unsigned complete : 1;    /* 1==complete, 0==not complete   */
+  unsigned freed : 1;       /* 1==freed, 0==allocated        */
+  unsigned native : 1;      /* 1==gcc-compiled task, 0==intel */
+  unsigned reserved31 : 7;  /* reserved for library use */
+} kmp_tasking_flags_t;
+
+typedef enum omp_interop_backend_type_t {
+  // reserve 0
+  omp_interop_backend_type_cuda_1 = 1,
+} omp_interop_backend_type_t;
+
+typedef enum kmp_interop_type_t {
+  kmp_interop_type_unknown = -1,
+  kmp_interop_type_platform,
+  kmp_interop_type_device,
+  kmp_interop_type_tasksync,
+} kmp_interop_type_t;
+
+/// The interop value type, aka. the interop object.
+typedef struct omp_interop_val_t {
+  /// Device and interop-type are determined at construction time and fix.
+  omp_interop_val_t(intptr_t device_id, kmp_interop_type_t interop_type)
+      : interop_type(interop_type), device_id(device_id) {}
+  const char *err_str = nullptr;
+  __tgt_async_info *async_info = nullptr;
+  __tgt_device_info device_info;
+  const kmp_interop_type_t interop_type;
+  const intptr_t device_id;
+  const intptr_t vendor_id = 1; // LLVM?
+  const intptr_t backend_type_id = omp_interop_backend_type_cuda_1;
+} omp_interop_val_t;
+
+#ifdef __cplusplus
+}
+#endif
+#endif

openmp/libomptarget/include/omptarget.h

 
 #ifndef _OMPTARGET_H_
 #define _OMPTARGET_H_
 
 #include <deque>
 #include <stddef.h>
 #include <stdint.h>
 
 #include <SourceInfo.h>
 
 #define OFFLOAD_SUCCESS (0)
 #define OFFLOAD_FAIL (~0)
 
 #define OFFLOAD_DEVICE_DEFAULT -1
 
 // Don't format out enums and structs.
 // clang-format off

 
 /// This struct is a record of non-contiguous information
 struct __tgt_target_non_contig {
   uint64_t Offset;
   uint64_t Count;
   uint64_t Stride;
 };
 
+struct __tgt_device_info {
+  void *Context = nullptr;
+  void *Device = nullptr;
+};
+
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 int omp_get_num_devices(void);
 int omp_get_initial_device(void);
 void *omp_target_alloc(size_t size, int device_num);
 void omp_target_free(void *device_ptr, int device_num);

openmp/libomptarget/include/omptargetplugin.h

 
 // Device synchronization. In case of success, return zero. Otherwise, return an
 // error code.
 int32_t __tgt_rtl_synchronize(int32_t ID, __tgt_async_info *AsyncInfo);
 
 // Set plugin's internal information flag externally.
 void __tgt_rtl_set_info_flag(uint32_t);
 
+int32_t __tgt_rtl_init_async_info(int32_t ID, __tgt_async_info **AsyncInfoPtr);
+int32_t __tgt_rtl_init_device_info(int32_t ID, __tgt_device_info *DeviceInfoPtr,
+                                   const char **errStr);
+
 #ifdef __cplusplus
 }
 #endif
 
 #endif // _OMPTARGETPLUGIN_H_

openmp/libomptarget/plugins/cuda/src/rtl.cpp

 //===----RTLs/cuda/src/rtl.cpp - Target RTLs Implementation ------- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // RTL for CUDA machine
 //
 //===----------------------------------------------------------------------===//
 
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <string>
 #include <unordered_map>
 #include <vector>
 
 #include "Debug.h"

 };
 
 enum ExecutionModeType {
   SPMD, // constructors, destructors,
   // combined constructs (`teams distribute parallel for [simd]`)
   GENERIC, // everything else
   NONE
 };
 
 /// Use a single entity to encode a kernel and a set of flags.
 struct KernelTy {
   CUfunction Func;
 
   // execution mode of kernel
   // 0 - SPMD mode (without master warp)
   // 1 - Generic mode (with master warp)
   int8_t ExecutionMode;
class DeviceRTLTy {
   // Clear entries table for a device
   void clearOffloadEntriesTable(const int DeviceId) {
     DeviceData[DeviceId].FuncGblEntries.emplace_back();
     FuncOrGblEntryTy &E = DeviceData[DeviceId].FuncGblEntries.back();
     E.Entries.clear();
     E.Table.EntriesBegin = E.Table.EntriesEnd = nullptr;
   }
 
+public:
   CUstream getStream(const int DeviceId, __tgt_async_info *AsyncInfo) const {
     assert(AsyncInfo && "AsyncInfo is nullptr");
 
     if (!AsyncInfo->Queue)
       AsyncInfo->Queue = StreamManager->getStream(DeviceId);
 
     return reinterpret_cast<CUstream>(AsyncInfo->Queue);
   }
 
-public:
   // This class should not be copied
   DeviceRTLTy(const DeviceRTLTy &) = delete;
   DeviceRTLTy(DeviceRTLTy &&) = delete;
 
   DeviceRTLTy()
       : NumberOfDevices(0), EnvNumTeams(-1), EnvTeamLimit(-1),
         EnvTeamThreadLimit(-1), RequiresFlags(OMP_REQ_UNDEFINED) {
 
int runTargetTeamRegion(const int DeviceId, void *TgtEntryPtr, void **TgtArgs,
 
     INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
          "Launching kernel %s with %d blocks and %d threads in %s "
          "mode\n",
          (getOffloadEntry(DeviceId, TgtEntryPtr))
              ? getOffloadEntry(DeviceId, TgtEntryPtr)->name
              : "(null)",
          CudaBlocksPerGrid, CudaThreadsPerBlock,
-         (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic");
+   (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic");
 
     CUstream Stream = getStream(DeviceId, AsyncInfo);
     Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
                          /* gridDimZ */ 1, CudaThreadsPerBlock,
                          /* blockDimY */ 1, /* blockDimZ */ 1,
                          /* sharedMemBytes */ 0, Stream, &Args[0], nullptr);
     if (!checkResult(Err, "Error returned from cuLaunchKernel\n"))
       return OFFLOAD_FAIL;
int synchronize(const int DeviceId, __tgt_async_info *AsyncInfo) const {
     if (Err != CUDA_SUCCESS) {
       DP("Error when synchronizing stream. stream = " DPxMOD
          ", async info ptr = " DPxMOD "\n",
          DPxPTR(Stream), DPxPTR(AsyncInfo));
       CUDA_ERR_STRING(Err);
     }
     return (Err == CUDA_SUCCESS) ? OFFLOAD_SUCCESS : OFFLOAD_FAIL;
   }
+
+  int releaseAsyncInfo(int DeviceId, __tgt_async_info *AsyncInfo) const {
+    if (AsyncInfo->Queue) {
+      StreamManager->returnStream(
+          DeviceId, reinterpret_cast<CUstream>(AsyncInfo->Queue));
+      AsyncInfo->Queue = nullptr;
+    }
+
+    return OFFLOAD_SUCCESS;
+  }
+
+  int initAsyncInfo(int DeviceId, __tgt_async_info **AsyncInfo) const {
+    CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
+    if (!checkResult(Err, "error returned from cuCtxSetCurrent"))
+      return OFFLOAD_FAIL;
+
+    *AsyncInfo = new __tgt_async_info;
+    getStream(DeviceId, *AsyncInfo);
+    return OFFLOAD_SUCCESS;
+  }
+
+  int initDeviceInfo(const int DeviceId, __tgt_device_info *DeviceInfo,
+                     const char **errStr) const {
+    assert(DeviceInfo && "DeviceInfo is nullptr");
+
+    if (!DeviceInfo->Context)
+      DeviceInfo->Context = DeviceData[DeviceId].Context;
+    if (!DeviceInfo->Device) {
+      CUdevice Dev;
+      CUresult Err = cuDeviceGet(&Dev, DeviceId);
+      if (Err == CUDA_SUCCESS) {
+        DeviceInfo->Device = reinterpret_cast<void *>(Dev);
+      } else {
+        cuGetErrorString(Err, errStr);
+        return OFFLOAD_FAIL;
+      }
+    }
+    return OFFLOAD_SUCCESS;
+  }
 };
 
 DeviceRTLTy DeviceRTL;
 } // namespace
 
 // Exposed library API function
 #ifdef __cplusplus
 extern "C" {
int32_t __tgt_rtl_synchronize(int32_t device_id,
                               __tgt_async_info *async_info_ptr) {
   assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
   assert(async_info_ptr && "async_info_ptr is nullptr");
   assert(async_info_ptr->Queue && "async_info_ptr->Queue is nullptr");
 
   return DeviceRTL.synchronize(device_id, async_info_ptr);
 }
 
+int32_t __tgt_rtl_release_async_info(int32_t device_id,
+                                     __tgt_async_info *async_info) {
+  assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
+  assert(async_info && "async_info is nullptr");
+
+  return DeviceRTL.releaseAsyncInfo(device_id, async_info);
+}
+
+int32_t __tgt_rtl_init_async_info(int32_t device_id, __tgt_async_info **async_info) {
+  assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
+  assert(async_info && "async_info is nullptr");
+
+  return DeviceRTL.initAsyncInfo(device_id, async_info);
+}
+
 void __tgt_rtl_set_info_flag(uint32_t NewInfoLevel) {
   std::atomic<uint32_t> &InfoLevel = getInfoLevelInternal();
   InfoLevel.store(NewInfoLevel);
 }
 
+int32_t __tgt_rtl_init_device_info(int32_t device_id,
+                                   __tgt_device_info *device_info_ptr,
+                                   const char **errStr) {
+  assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
+  assert(device_info_ptr && "device_info_ptr is nullptr");
+
+  return DeviceRTL.initDeviceInfo(device_id, device_info_ptr, errStr);
+}
+
 #ifdef __cplusplus
 }
 #endif

openmp/libomptarget/plugins/exports

global:
     __tgt_rtl_run_target_team_region_async;
     __tgt_rtl_run_target_region;
     __tgt_rtl_run_target_region_async;
     __tgt_rtl_synchronize;
     __tgt_rtl_register_lib;
     __tgt_rtl_unregister_lib;
     __tgt_rtl_supports_empty_images;
     __tgt_rtl_set_info_flag;
+    __tgt_rtl_init_device_info;
+    __tgt_rtl_init_async_info;
   local:
     *;
 };

openmp/libomptarget/src/CMakeLists.txt

 ##===----------------------------------------------------------------------===##
 
 libomptarget_say("Building offloading runtime library libomptarget.")
 
 set(LIBOMPTARGET_SRC_FILES
   ${CMAKE_CURRENT_SOURCE_DIR}/api.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/device.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/interface.cpp
-  ${CMAKE_CURRENT_SOURCE_DIR}/rtl.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/interop.cpp
   ${CMAKE_CURRENT_SOURCE_DIR}/omptarget.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/rtl.cpp
 )
 
 set(LIBOMPTARGET_SRC_FILES ${LIBOMPTARGET_SRC_FILES} PARENT_SCOPE)
 
 include_directories(${LIBOMPTARGET_LLVM_INCLUDE_DIRS})
 
 # Build libomptarget library with libdl dependency.
 add_library(omptarget SHARED ${LIBOMPTARGET_SRC_FILES})

openmp/libomptarget/src/exports

 VERS1.0 {
   global:
     __tgt_register_requires;
     __tgt_register_lib;
     __tgt_unregister_lib;
     __tgt_target_data_begin;
     __tgt_target_data_end;
     __tgt_target_data_update;
     __tgt_target;
global:
     omp_target_memcpy;
     omp_target_memcpy_rect;
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;
     llvm_omp_target_alloc_host;
     llvm_omp_target_alloc_shared;
     llvm_omp_target_alloc_device;
     __tgt_set_info_flag;
+    omp_get_interop_ptr;
+    omp_get_interop_str;
+    omp_get_interop_int;
+    omp_get_interop_name;
+    omp_get_interop_type_desc;
+    omp_get_interop_rc_desc;
+    __tgt_interop_init;
+    __tgt_interop_use;
+    __tgt_interop_destroy;
   local:
     *;
 };
 

openmp/libomptarget/src/interop.cpp

+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "interop.h"
+#include "private.h"
+
+namespace {
+omp_interop_rc_t
+__tgt_interop_get_property_err_type(omp_interop_property_t property) {
+  switch (property) {
+  case omp_ipr_fr_id:
+    return omp_irc_type_int;
+  case omp_ipr_fr_name:
+    return omp_irc_type_str;
+  case omp_ipr_vendor:
+    return omp_irc_type_int;
+  case omp_ipr_vendor_name:
+    return omp_irc_type_str;
+  case omp_ipr_device_num:
+    return omp_irc_type_int;
+  case omp_ipr_platform:
+    return omp_irc_type_int;
+  case omp_ipr_device:
+    return omp_irc_type_ptr;
+  case omp_ipr_device_context:
+    return omp_irc_type_ptr;
+  case omp_ipr_targetsync:
+    return omp_irc_type_ptr;
+  };
+  return omp_irc_no_value;
+}
+
+void __tgt_interop_type_mismatch(omp_interop_property_t property, int *err) {
+  if (err)
+    *err = __tgt_interop_get_property_err_type(property);
+}
+
+const char *__tgt_interop_vendor_id_to_str(intptr_t vendor_id) {
+  switch (vendor_id) {
+  case 1:
+    return ("cuda");
+  case 2:
+    return ("cuda_driver");
+  case 3:
+    return ("opencl");
+  case 4:
+    return ("sycl");
+  case 5:
+    return ("hip");
+  case 6:
+    return ("level_zero");
+  default:
+    return ("unknown");
+  }
+}
+
+template <typename PropertyTy>
+PropertyTy __tgt_interop_get_property(omp_interop_val_t &interop_val,
+                                      omp_interop_property_t property,
+                                      int *err);
+
+template <>
+intptr_t __tgt_interop_get_property<intptr_t>(omp_interop_val_t &interop_val,
+                                              omp_interop_property_t property,
+                                              int *err) {
+  switch (property) {
+  case omp_ipr_fr_id:
+    return interop_val.backend_type_id;
+  case omp_ipr_vendor:
+    return interop_val.vendor_id;
+  case omp_ipr_device_num:
+    return interop_val.device_id;
+  default:;
+  }
+  __tgt_interop_type_mismatch(property, err);
+  return 0;
+}
+
+template <>
+const char *__tgt_interop_get_property<const char *>(
+    omp_interop_val_t &interop_val, omp_interop_property_t property, int *err) {
+  switch (property) {
+  case omp_ipr_fr_id:
+    return interop_val.interop_type == kmp_interop_type_tasksync
+               ? "tasksync"
+               : "device+context";
+  case omp_ipr_vendor_name:
+    return __tgt_interop_vendor_id_to_str(interop_val.vendor_id);
+  default:
+    __tgt_interop_type_mismatch(property, err);
+    return nullptr;
+  }
+}
+
+template <>
+void *__tgt_interop_get_property<void *>(omp_interop_val_t &interop_val,
+                                         omp_interop_property_t property,
+                                         int *err) {
+  switch (property) {
+  case omp_ipr_device:
+    if (interop_val.device_info.Device)
+      return interop_val.device_info.Device;
+    *err = omp_irc_no_value;
+    return const_cast<char *>(interop_val.err_str);
+  case omp_ipr_device_context:
+    return interop_val.device_info.Context;
+  case omp_ipr_targetsync:
+    return interop_val.async_info->Queue;
+  default:;
+  }
+  __tgt_interop_type_mismatch(property, err);
+  return nullptr;
+}
+
+bool __tgt_interop_get_property_check(omp_interop_val_t **interop_ptr,
+                                      omp_interop_property_t property,
+                                      int *err) {
+  if (err)
+    *err = omp_irc_success;
+  if (!interop_ptr) {
+    if (err)
+      *err = omp_irc_empty;
+    return false;
+  }
+  if (property >= 0 || property < omp_ipr_first) {
+    if (err)
+      *err = omp_irc_out_of_range;
+    return false;
+  }
+  if (property == omp_ipr_targetsync &&
+      (*interop_ptr)->interop_type != kmp_interop_type_tasksync) {
+    if (err)
+      *err = omp_irc_other;
+    return false;
+  }
+  if ((property == omp_ipr_device || property == omp_ipr_device_context) &&
+      (*interop_ptr)->interop_type == kmp_interop_type_tasksync) {
+    if (err)
+      *err = omp_irc_other;
+    return false;
+  }
+  return true;
+}
+
+} // namespace
+
+#define __OMP_GET_INTEROP_TY(RETURN_TYPE, SUFFIX)                              \
+  RETURN_TYPE omp_get_interop_##SUFFIX(const omp_interop_t interop,            \
+                                       omp_interop_property_t property_id,     \
+                                       int *err) {                             \
+    omp_interop_val_t *interop_val = (omp_interop_val_t *)interop;             \
+    assert((interop_val)->interop_type == kmp_interop_type_tasksync);          \
+    if (!__tgt_interop_get_property_check(&interop_val, property_id, err)) {   \
+      return (RETURN_TYPE)(0);                                                 \
+    }                                                                          \
+    return __tgt_interop_get_property<RETURN_TYPE>(*interop_val, property_id,  \
+                                                   err);                       \
+  }
+__OMP_GET_INTEROP_TY(intptr_t, int)
+__OMP_GET_INTEROP_TY(void *, ptr)
+__OMP_GET_INTEROP_TY(const char *, str)
+#undef __OMP_GET_INTEROP_TY
+
+#define __OMP_GET_INTEROP_TY3(RETURN_TYPE, SUFFIX)                             \
+  RETURN_TYPE omp_get_interop_##SUFFIX(const omp_interop_t interop,            \
+                                       omp_interop_property_t property_id) {   \
+    int err;                                                                   \
+    omp_interop_val_t *interop_val = (omp_interop_val_t *)interop;             \
+    if (!__tgt_interop_get_property_check(&interop_val, property_id, &err)) {  \
+      return (RETURN_TYPE)(0);                                                 \
+    }                                                                          \
+    return nullptr;                                                            \
+    return __tgt_interop_get_property<RETURN_TYPE>(*interop_val, property_id,  \
+                                                   &err);                      \
+  }
+__OMP_GET_INTEROP_TY3(const char *, name)
+__OMP_GET_INTEROP_TY3(const char *, type_desc)
+__OMP_GET_INTEROP_TY3(const char *, rc_desc)
+#undef __OMP_GET_INTEROP_TY3
+
+typedef int64_t kmp_int64;
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+void __tgt_interop_init(ident_t *loc_ref, kmp_int32 gtid,
+                        omp_interop_val_t *&interop_ptr,
+                        kmp_interop_type_t interop_type, kmp_int32 device_id,
+                        kmp_int64 ndeps, kmp_depend_info_t *dep_list,
+                        kmp_int32 have_nowait) {
+  kmp_int32 ndeps_noalias = 0;
+  kmp_depend_info_t *noalias_dep_list = NULL;
+  assert(interop_type != kmp_interop_type_unknown &&
+         "Cannot initialize with unknown interop_type!");
+  if (device_id == -1) {
+    device_id = omp_get_default_device();
+  }
+
+  if (interop_type == kmp_interop_type_tasksync) {
+    __kmpc_omp_wait_deps(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                         noalias_dep_list);
+  }
+
+  interop_ptr = new omp_interop_val_t(device_id, interop_type);
+  if (!device_is_ready(device_id)) {
+    interop_ptr->err_str = "Device not ready!";
+    return;
+  }
+
+  DeviceTy &Device = PM->Devices[device_id];
+  if (interop_type == kmp_interop_type_tasksync) {
+    if (!Device.RTL || !Device.RTL->init_async_info ||
+        Device.RTL->init_async_info(device_id, &(interop_ptr)->async_info)) {
+      delete interop_ptr;
+      interop_ptr = omp_interop_none;
+    }
+  } else {
+    if (!Device.RTL || !Device.RTL->init_device_info ||
+        Device.RTL->init_device_info(device_id, &(interop_ptr)->device_info,
+                                     &(interop_ptr)->err_str)) {
+      delete interop_ptr;
+      interop_ptr = omp_interop_none;
+    }
+  }
+}
+
+void __tgt_interop_use(ident_t *loc_ref, kmp_int32 gtid,
+                       omp_interop_val_t *interop_ptr, kmp_int32 device_id,
+                       kmp_int32 ndeps, kmp_depend_info_t *dep_list,
+                       kmp_int32 have_nowait) {
+  kmp_int32 ndeps_noalias = 0;
+  kmp_depend_info_t *noalias_dep_list = NULL;
+  assert(interop_ptr && "Cannot use nullptr!");
+  omp_interop_val_t *interop_val = interop_ptr;
+  assert(interop_val != omp_interop_none &&
+         "Cannot use uninitialized interop_ptr!");
+  assert((device_id == -1 || interop_val->device_id == device_id) &&
+         "Inconsistent device-id usage!");
+
+  if (interop_val->interop_type == kmp_interop_type_tasksync) {
+    __kmpc_omp_wait_deps(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                         noalias_dep_list);
+  }
+}
+
+void __tgt_interop_destroy(ident_t *loc_ref, kmp_int32 gtid,
+                           omp_interop_val_t *&interop_ptr, kmp_int32 device_id,
+                           kmp_int32 ndeps, kmp_depend_info_t *dep_list,
+                           kmp_int32 have_nowait) {
+  kmp_int32 ndeps_noalias = 0;
+  kmp_depend_info_t *noalias_dep_list = NULL;
+  assert(interop_ptr && "Cannot use nullptr!");
+  omp_interop_val_t *interop_val = interop_ptr;
+
+  if (interop_val == omp_interop_none)
+    return;
+
+  assert((device_id == -1 || interop_val->device_id == device_id) &&
+         "Inconsistent device-id usage!");
+
+  if (interop_val->interop_type == kmp_interop_type_tasksync) {
+    __kmpc_omp_wait_deps(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                         noalias_dep_list);
+  }
+
+  delete interop_ptr;
+  interop_ptr = omp_interop_none;
+}
+#ifdef __cplusplus
+} // extern "C"
+#endif

openmp/libomptarget/src/private.h

typedef int (*TargetDataFuncPtrTy)(ident_t *, DeviceTy &, int32_t, void **,
                                    void **, int64_t *, int64_t *,
                                    map_var_info_t *, void **, AsyncInfoTy &,
                                    bool);
 
 // Implemented in libomp, they are called from within __tgt_* functions.
 #ifdef __cplusplus
 extern "C" {
 #endif
+
+/*!
+ * The ident structure that describes a source location.
+ * The struct is identical to the one in the kmp.h file.
+ * We maintain the same data structure for compatibility.
+ */
+typedef int kmp_int32;
+typedef intptr_t kmp_intptr_t;
+
+// Compiler sends us this info:
+typedef struct kmp_depend_info {
+  kmp_intptr_t base_addr;
+  size_t len;
+  struct {
+    bool in : 1;
+    bool out : 1;
+    bool mtx : 1;
+  } flags;
+} kmp_depend_info_t;
+
 // functions that extract info from libomp; keep in sync
 int omp_get_default_device(void) __attribute__((weak));
 int32_t __kmpc_omp_taskwait(void *loc_ref, int32_t gtid) __attribute__((weak));
 int32_t __kmpc_global_thread_num(void *) __attribute__((weak));
 int __kmpc_get_target_offload(void) __attribute__((weak));
+void __kmpc_omp_wait_deps(ident_t *loc_ref, kmp_int32 gtid, kmp_int32 ndeps,
+                          kmp_depend_info_t *dep_list, kmp_int32 ndeps_noalias,
+                          kmp_depend_info_t *noalias_dep_list)
+    __attribute__((weak));
 #ifdef __cplusplus
 }
 #endif
 
 #define TARGET_NAME Libomptarget
 #define DEBUG_PREFIX GETNAME(TARGET_NAME)
 
 ////////////////////////////////////////////////////////////////////////////////

openmp/libomptarget/src/rtl.h

typedef int32_t(run_team_region_async_ty)(int32_t, void *, void **,
                                             ptrdiff_t *, int32_t, int32_t,
                                             int32_t, uint64_t,
                                             __tgt_async_info *);
   typedef int64_t(init_requires_ty)(int64_t);
   typedef int64_t(synchronize_ty)(int32_t, __tgt_async_info *);
   typedef int32_t (*register_lib_ty)(__tgt_bin_desc *);
   typedef int32_t(supports_empty_images_ty)();
   typedef void(set_info_flag_ty)(uint32_t);
+  typedef int32_t(release_async_info_ty)(int32_t, __tgt_async_info *);
+  typedef int32_t(init_async_info_ty)(int32_t, __tgt_async_info **);
+  typedef int64_t(init_device_into_ty)(int64_t, __tgt_device_info *,
+                                       const char **);
 
   int32_t Idx = -1;             // RTL index, index is the number of devices
                                 // of other RTLs that were registered before,
                                 // i.e. the OpenMP index of the first device
                                 // to be registered with this RTL.
   int32_t NumberOfDevices = -1; // Number of devices this RTL deals with.
 
   void *LibraryHandler = nullptr;
#endif
   run_team_region_ty *run_team_region = nullptr;
   run_team_region_async_ty *run_team_region_async = nullptr;
   init_requires_ty *init_requires = nullptr;
   synchronize_ty *synchronize = nullptr;
   register_lib_ty register_lib = nullptr;
   register_lib_ty unregister_lib = nullptr;
   supports_empty_images_ty *supports_empty_images = nullptr;
   set_info_flag_ty *set_info_flag = nullptr;
+  init_async_info_ty *init_async_info = nullptr;
+  init_device_into_ty *init_device_info = nullptr;
+  release_async_info_ty *release_async_info = nullptr;
 
   // Are there images associated with this RTL.
   bool isUsed = false;
 
   // Mutex for thread-safety when calling RTL interface functions.
   // It is easier to enforce thread-safety at the libomptarget level,
   // so that developers of new RTLs do not have to worry about it.
   std::mutex Mtx;

openmp/libomptarget/src/rtl.cpp

#endif
     *((void **)&R.register_lib) =
         dlsym(dynlib_handle, "__tgt_rtl_register_lib");
     *((void **)&R.unregister_lib) =
         dlsym(dynlib_handle, "__tgt_rtl_unregister_lib");
     *((void **)&R.supports_empty_images) =
         dlsym(dynlib_handle, "__tgt_rtl_supports_empty_images");
     *((void **)&R.set_info_flag) =
         dlsym(dynlib_handle, "__tgt_rtl_set_info_flag");
+    *((void **)&R.release_async_info) =
+        dlsym(dynlib_handle, "__tgt_rtl_release_async_info");
+    *((void **)&R.init_async_info) = dlsym(dynlib_handle, "__tgt_rtl_init_async_info");
+    *((void **)&R.init_device_info) = dlsym(dynlib_handle, "__tgt_rtl_init_device_info");
+
   }
 
   DP("RTLs loaded!\n");
 
   return;
 }
 
 ////////////////////////////////////////////////////////////////////////////////

openmp/runtime/src/dllexports

kmp_set_disp_num_buffers                    890
     omp_fulfill_event                       759
     omp_set_num_teams                       802
     omp_get_max_teams                       803
     omp_set_teams_thread_limit              804
     omp_get_teams_thread_limit              805
     omp_display_env                         733
     omp_calloc                              776
     omp_realloc                             777
+    omp_get_interop_int                     760
+    omp_get_interop_ptr                     761
+    omp_get_interop_str                     762
 
     omp_null_allocator                     DATA
     omp_default_mem_alloc                  DATA
     omp_large_cap_mem_alloc                DATA
     omp_const_mem_alloc                    DATA
     omp_high_bw_mem_alloc                  DATA
     omp_low_lat_mem_alloc                  DATA
     omp_cgroup_mem_alloc                   DATA
kmp_set_disp_num_buffers                    890
     # Preview of target memory support
     llvm_omp_target_host_mem_space         DATA
     llvm_omp_target_shared_mem_space       DATA
     llvm_omp_target_device_mem_space       DATA
 
 %ifndef stub
     # Ordinals between 900 and 999 are reserved
 
     # Ordinals between 1000 and 1999 are reserved
     # for user-callable uppercase Fortran entries.
 
 
     # ATOMIC entries
 
     %ifdef HAVE_QUAD
     __kmpc_atomic_cmplx16_div              2000
     %endif

openmp/runtime/src/kmp_ftn_entry.h

 /*
  * kmp_ftn_entry.h -- Fortran entry linkage support for OpenMP.
  */
 
 //===----------------------------------------------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef FTN_STDCALL
 #error The support file kmp_ftn_entry.h should not be compiled by itself.
 #endif
 
 #ifdef KMP_STUB
 #include "kmp_stub.h"
 #endif
 
 #include "kmp_i18n.h"
 

 // APPEND Fortran entries in this file. The GOMP versioned symbols of the C API
 // will take place where the ompc_* functions are defined.
 #if KMP_FTN_ENTRIES == KMP_FTN_APPEND
 #define KMP_EXPAND_NAME_IF_APPEND(name) KMP_EXPAND_NAME(name)
 #else
 #define KMP_EXPAND_NAME_IF_APPEND(name) name
 #endif
 
 void FTN_STDCALL FTN_SET_STACKSIZE(int KMP_DEREF arg) {
 #ifdef KMP_STUB
   __kmps_set_stacksize(KMP_DEREF arg);
 #else
   // __kmp_aux_set_stacksize initializes the library if needed
   __kmp_aux_set_stacksize((size_t)KMP_DEREF arg);
 #endif
 }
 
public:
   const char *get() const { return buf; }
 };
 #endif // KMP_STUB
 
 /*
  * Set the value of the affinity-format-var ICV on the current device to the
  * format specified in the argument.
  */
 void FTN_STDCALL KMP_EXPAND_NAME_IF_APPEND(FTN_SET_AFFINITY_FORMAT)(
     char const *format, size_t size) {
 #ifdef KMP_STUB
   return;
 #else
   if (!__kmp_init_serial) {
     __kmp_serial_initialize();
   }
   ConvertedString cformat(format, size);
__kmp_entry_thread()->th.th_current_task->td_icvs.default_device =
       KMP_DEREF arg;
 #endif
 }
 
 // Get number of NON-HOST devices.
 // libomptarget, if loaded, provides this function in api.cpp.
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)(void)
     KMP_WEAK_ATTRIBUTE_EXTERNAL;
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)(void) {
 #if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
   return 0;
 #else
   int (*fptr)();
   if ((*(void **)(&fptr) = KMP_DLSYM("__tgt_get_num_devices"))) {
     return (*fptr)();
   } else if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_num_devices"))) {
     return (*fptr)();
   } else if ((*(void **)(&fptr) = KMP_DLSYM("_Offload_number_of_devices"))) {
     return (*fptr)();
   } else { // liboffload & libomptarget don't exist
     return 0;
   }
 #endif // KMP_MIC || KMP_OS_DARWIN || KMP_OS_WINDOWS || defined(KMP_STUB)
 }
 
 // This function always returns true when called on host device.
 // Compiler/libomptarget should handle when it is called inside target region.
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_IS_INITIAL_DEVICE)(void)
     KMP_WEAK_ATTRIBUTE_EXTERNAL;
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_IS_INITIAL_DEVICE)(void) {
   return 1; // This is the host
 }
 
 // libomptarget, if loaded, provides this function
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)(void)
     KMP_WEAK_ATTRIBUTE_EXTERNAL;
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)(void) {
   // same as omp_get_num_devices()
   return KMP_EXPAND_NAME(FTN_GET_NUM_DEVICES)();
 }
 
 #if defined(KMP_STUB)
 // Entries for stubs library
 // As all *target* functions are C-only parameters always passed by value
 void *FTN_STDCALL FTN_TARGET_ALLOC(size_t size, int device_num) { return 0; }
#else
   return __kmp_max_task_priority;
 #endif
 }
 
 // This function will be defined in libomptarget. When libomptarget is not
 // loaded, we assume we are on the host and return KMP_HOST_DEVICE.
 // Compiler/libomptarget will handle this if called inside target.
 int FTN_STDCALL FTN_GET_DEVICE_NUM(void) KMP_WEAK_ATTRIBUTE_EXTERNAL;
 int FTN_STDCALL FTN_GET_DEVICE_NUM(void) {
   return KMP_EXPAND_NAME(FTN_GET_INITIAL_DEVICE)();
 }
 
 // Compiler will ensure that this is only called from host in sequential region
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_PAUSE_RESOURCE)(kmp_pause_status_t kind,
                                                     int device_num) {
 #ifdef KMP_STUB
   return 1; // just fail

 #else
   if (!__kmp_init_serial) {
     __kmp_serial_initialize();
   }
   return __kmp_get_teams_thread_limit();
 #endif
 }
 
+/// TODO: Include the `omp.h` of the current build
+/* OpenMP 5.1 interop */
+typedef intptr_t omp_intptr_t;
+
+/* 0..omp_get_num_interop_properties()-1 are reserved for implementation-defined
+ * properties */
+typedef enum omp_interop_property {
+  omp_ipr_fr_id = -1,
+  omp_ipr_fr_name = -2,
+  omp_ipr_vendor = -3,
+  omp_ipr_vendor_name = -4,
+  omp_ipr_device_num = -5,
+  omp_ipr_platform = -6,
+  omp_ipr_device = -7,
+  omp_ipr_device_context = -8,
+  omp_ipr_targetsync = -9,
+  omp_ipr_first = -9
+} omp_interop_property_t;
+
+#define omp_interop_none 0
+
+typedef enum omp_interop_rc {
+  omp_irc_no_value = 1,
+  omp_irc_success = 0,
+  omp_irc_empty = -1,
+  omp_irc_out_of_range = -2,
+  omp_irc_type_int = -3,
+  omp_irc_type_ptr = -4,
+  omp_irc_type_str = -5,
+  omp_irc_other = -6
+} omp_interop_rc_t;
+
+typedef enum omp_interop_fr {
+  omp_ifr_cuda = 1,
+  omp_ifr_cuda_driver = 2,
+  omp_ifr_opencl = 3,
+  omp_ifr_sycl = 4,
+  omp_ifr_hip = 5,
+  omp_ifr_level_zero = 6,
+  omp_ifr_last = 7
+} omp_interop_fr_t;
+
+typedef void *omp_interop_t;
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return 0.
+// Compiler/libomptarget will handle this if called inside target.
+int FTN_STDCALL FTN_GET_NUM_INTEROP_PROPERTIES(const omp_interop_t) {
+  return 0;
+} 
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return 0.
+// Compiler/libomptarget will handle this if called inside target.
+omp_intptr_t FTN_STDCALL FTN_GET_INTEROP_INT(const omp_interop_t,
+                                                       omp_interop_property_t property_id,
+                                                       int *err) {
+  return 0;
+}
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return nullptr.
+// Compiler/libomptarget will handle this if called inside target.
+void * FTN_STDCALL FTN_GET_INTEROP_PTR(const omp_interop_t,
+                                                omp_interop_property_t property_id, int *err) {
+  return nullptr;
+} 
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return nullptr.
+// Compiler/libomptarget will handle this if called inside target.
+const char *FTN_STDCALL FTN_GET_INTEROP_STR(const omp_interop_t,
+                                                      omp_interop_property_t property_id,
+                                                      int *err) {
+  return nullptr;
+}
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return nullptr.
+// Compiler/libomptarget will handle this if called inside target.
+const char * FTN_STDCALL FTN_GET_INTEROP_NAME(const omp_interop_t,
+                                                       omp_interop_property_t property_id) { 
+  return nullptr;
+}
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return nullptr.
+// Compiler/libomptarget will handle this if called inside target.
+const char * FTN_STDCALL FTN_GET_INTEROP_TYPE_DESC(const omp_interop_t, omp_interop_property_t property_id) {
+  return nullptr;
+}
+
+// This function is defined in libomptarget. When libomptarget is not
+// loaded, we assume we are on the host and return nullptr.
+// Compiler/libomptarget will handle this if called inside target.
+const char * FTN_STDCALL FTN_GET_INTEROP_RC_DESC(const omp_interop_t, omp_interop_rc_t property_id) {
+  return nullptr;
+}
+
 // display environment variables when requested
 void FTN_STDCALL FTN_DISPLAY_ENV(int verbose) {
 #ifndef KMP_STUB
   __kmp_omp_display_env(verbose);
 #endif
 }
 
 // GCC compatibility (versioned symbols)
 #ifdef KMP_USE_VERSION_SYMBOLS
 
 /* These following sections create versioned symbols for the
    omp_* routines. The KMP_VERSION_SYMBOL macro expands the API name and
    then maps it to a versioned symbol.
    libgomp ``versions'' its symbols (OMP_1.0, OMP_2.0, OMP_3.0, ...) while also
    retaining the default version which libomp uses: VERSION (defined in
    exports_so.txt). If you want to see the versioned symbols for libgomp.so.1
    then just type:
 
    objdump -T /path/to/libgomp.so.1 | grep omp_
 
    Example:
    Step 1) Create __kmp_api_omp_set_num_threads_10_alias which is alias of
      __kmp_api_omp_set_num_threads
    Step 2) Set __kmp_api_omp_set_num_threads_10_alias to version:
      omp_set_num_threads@OMP_1.0
    Step 2B) Set __kmp_api_omp_set_num_threads to default version:
      omp_set_num_threads@@VERSION
 */
 
 // OMP_1.0 versioned symbols
 KMP_VERSION_SYMBOL(FTN_SET_NUM_THREADS, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_GET_NUM_THREADS, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_GET_MAX_THREADS, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_GET_THREAD_NUM, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_GET_NUM_PROCS, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_IN_PARALLEL, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_SET_DYNAMIC, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_GET_DYNAMIC, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_SET_NESTED, 10, "OMP_1.0");
 KMP_VERSION_SYMBOL(FTN_GET_NESTED, 10, "OMP_1.0");

openmp/runtime/src/kmp_ftn_os.h

 #define FTN_GET_SUPPORTED_ACTIVE_LEVELS omp_get_supported_active_levels
 #define FTN_DISPLAY_ENV omp_display_env
 #define FTN_FULFILL_EVENT omp_fulfill_event
 #define FTN_SET_NUM_TEAMS omp_set_num_teams
 #define FTN_GET_MAX_TEAMS omp_get_max_teams
 #define FTN_SET_TEAMS_THREAD_LIMIT omp_set_teams_thread_limit
 #define FTN_GET_TEAMS_THREAD_LIMIT omp_get_teams_thread_limit
 
+#define FTN_GET_NUM_INTEROP_PROPERTIES omp_get_num_interop_properties
+#define FTN_GET_INTEROP_INT omp_get_interop_int
+#define FTN_GET_INTEROP_PTR omp_get_interop_ptr
+#define FTN_GET_INTEROP_STR omp_get_interop_str
+#define FTN_GET_INTEROP_NAME omp_get_interop_name
+#define FTN_GET_INTEROP_TYPE_DESC omp_get_interop_type_desc
+#define FTN_GET_INTEROP_RC_DESC omp_get_interop_rc_desc
+
 #endif /* KMP_FTN_PLAIN */
 
 /* ------------------------------------------------------------------------ */
 
 #if KMP_FTN_ENTRIES == KMP_FTN_APPEND
 
 #define FTN_SET_STACKSIZE kmp_set_stacksize_
 #define FTN_SET_STACKSIZE_S kmp_set_stacksize_s_