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

19

Probably not an option.

openmp/libomptarget/include/omptarget.h

21

Why is that?

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

76

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

1152–1153

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

13

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

1

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

65

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

109

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.

160

Remove the asserts.

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

76

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

45

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

177

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.

191

This assertion seems nonsensical, just return I guess.

204

Move the wait stuff to the beginning.

216

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

19

Okay. Any suggestions on how to overcome this?

109

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

177

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

177

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 ↗

(On Diff #361893)

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

13

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 ↗

(On Diff #362280)

ditto

1515 ↗

(On Diff #362280)

ditto

1527 ↗

(On Diff #362280)

ditto

1539 ↗

(On Diff #362280)

ditto

1550 ↗

(On Diff #362280)

same as below

1561 ↗

(On Diff #362280)

same as below

1570–1574 ↗

(On Diff #362280)

[sriharikrishna]

Runtime for Interop directive

[tianshilei1992]

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

1165–1171

openmp/libomptarget/src/interop.cpp

13

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 ↗

(On Diff #362524)

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

13

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 ↗

(On Diff #362524)

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 ↗

(On Diff #362524)

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

+#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
+
+/// 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;
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/*!
+ * 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);
+#ifdef __cplusplus
+}
+#endif
+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;
+
+extern "C" {
+
+/// 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) {
+    assert(interop_type != kmp_interop_type_unknown);
+    assert(!async_info);
+    assert(!device_info.Device);
+    assert(!device_info.Context);
+  }
+  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;
+}
+#endif

openmp/libomptarget/include/omptarget.h

 
 #ifndef _OMPTARGET_H_
 #define _OMPTARGET_H_
 
 #include <deque>
 #include <stddef.h>
 #include <stdint.h>
 
-#include <SourceInfo.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

 
 /// Keep entries table per device.
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
   std::vector<__tgt_offload_entry> Entries;
 };
 
 enum ExecutionModeType {
-  SPMD, // constructors, destructors,
+  SPMD,         // constructors, destructors,
                 // combined constructs (`teams distribute parallel for [simd]`)
-  GENERIC, // everything else
+  GENERIC,      // everything else
+  SPMD_GENERIC, // Generic kernel with SPMD execution
   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)
+  // 2 - SPMD mode execution with Generic mode semantics.
   int8_t ExecutionMode;
 
   /// Maximal number of threads per block for this kernel.
   int MaxThreadsPerBlock = 0;
 
   KernelTy(CUfunction _Func, int8_t _ExecutionMode)
       : Func(_Func), ExecutionMode(_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) {
 
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
         if (Err != CUDA_SUCCESS) {
           REPORT("Error when copying data from device to host. Pointers: "
                  "host = " DPxMOD ", device = " DPxMOD ", size = %zd\n",
                  DPxPTR(&ExecModeVal), DPxPTR(ExecModePtr), CUSize);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
-        if (ExecModeVal < 0 || ExecModeVal > 1) {
+        if (ExecModeVal < 0 || ExecModeVal > 2) {
           DP("Error wrong exec_mode value specified in cubin file: %d\n",
              ExecModeVal);
           return nullptr;
         }
       } else {
         REPORT("Loading global exec_mode '%s' - symbol missing, using default "
                "value GENERIC (1)\n",
                ExecModeName);
int runTargetTeamRegion(const int DeviceId, void *TgtEntryPtr, void **TgtArgs,
     if (TeamNum <= 0) {
       if (LoopTripCount > 0 && EnvNumTeams < 0) {
         if (KernelInfo->ExecutionMode == SPMD) {
           // We have a combined construct, i.e. `target teams distribute
           // parallel for [simd]`. We launch so many teams so that each thread
           // will execute one iteration of the loop. round up to the nearest
           // integer
           CudaBlocksPerGrid = ((LoopTripCount - 1) / CudaThreadsPerBlock) + 1;
-        } else {
+        } else if (KernelInfo->ExecutionMode == GENERIC) {
           // If we reach this point, then we have a non-combined construct, i.e.
           // `teams distribute` with a nested `parallel for` and each team is
           // assigned one iteration of the `distribute` loop. E.g.:
           //
           // #pragma omp target teams distribute
           // for(...loop_tripcount...) {
           //   #pragma omp parallel for
           //   for(...) {}
           // }
           //
           // Threads within a team will execute the iterations of the `parallel`
           // loop.
           CudaBlocksPerGrid = LoopTripCount;
+        } else if (KernelInfo->ExecutionMode == SPMD_GENERIC) {
+          // If we reach this point, then we are executing a kernel that was
+          // transformed from Generic-mode to SPMD-mode. This kernel has
+          // SPMD-mode execution, but needs its blocks to be scheduled
+          // differently because the current loop trip count only applies to the
+          // `teams distribute` region and will create var too few blocks using
+          // the regular SPMD-mode method.
+          CudaBlocksPerGrid = LoopTripCount;
+        } else {
+          REPORT("Unknown execution mode: %d\n", KernelInfo->ExecutionMode);
+          return OFFLOAD_FAIL;
         }
         DP("Using %d teams due to loop trip count %" PRIu32
            " and number of threads per block %d\n",
            CudaBlocksPerGrid, LoopTripCount, CudaThreadsPerBlock);
       } else {
         DP("Using default number of teams %d\n", DeviceData[DeviceId].NumTeams);
         CudaBlocksPerGrid = DeviceData[DeviceId].NumTeams;
       }
     } else if (TeamNum > DeviceData[DeviceId].BlocksPerGrid) {
       DP("Capping number of teams to team limit %d\n",
          DeviceData[DeviceId].BlocksPerGrid);
       CudaBlocksPerGrid = DeviceData[DeviceId].BlocksPerGrid;
     } else {
       DP("Using requested number of teams %d\n", TeamNum);
       CudaBlocksPerGrid = TeamNum;
     }
 
     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 
+             ? (KernelInfo->ExecutionMode == GENERIC ? "Generic" : "SPMD-Generic")
+             : "SPMD"));
 
     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;
+
+    __tgt_async_info *P = new __tgt_async_info;
+    *AsyncInfo = P;
+    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}/omptarget.cpp
+  ${CMAKE_CURRENT_SOURCE_DIR}/interop.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 {
+VERSION {
   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;
+    __kmpc_interop_init;
+    __kmpc_interop_use;
+    __kmpc_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"
+
+static omp_interop_rc_t
+__kmpc_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;
+}
+
+static void __kmpc_interop_type_mismatch(omp_interop_property_t property,
+                                         int *err) {
+  if (err)
+    *err = __kmpc_interop_get_property_err_type(property);
+}
+
+/// TODO Use vendor_id to look up the string
+static const char *__kmpc_interop_vendor_id_to_str(intptr_t vendor_id) {
+  return "CUDA_driver";
+}
+
+template <typename PropertyTy>
+static PropertyTy __kmpc_interop_get_property(omp_interop_val_t &interop_val,
+                                              omp_interop_property_t property,
+                                              int *err);
+
+template <>
+intptr_t __kmpc_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:;
+  }
+  __kmpc_interop_type_mismatch(property, err);
+  return 0;
+}
+
+template <>
+const char *__kmpc_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 __kmpc_interop_vendor_id_to_str(interop_val.vendor_id);
+  default:
+    __kmpc_interop_type_mismatch(property, err);
+    return nullptr;
+  }
+}
+
+template <>
+void *__kmpc_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:;
+  }
+  __kmpc_interop_type_mismatch(property, err);
+  return nullptr;
+}
+
+static bool __kmpc_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;
+}
+
+#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 (!__kmpc_interop_get_property_check(&interop_val, property_id, err)) {  \
+      return (RETURN_TYPE)(0);                                                 \
+    }                                                                          \
+    return __kmpc_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 (!__kmpc_interop_get_property_check(&interop_val, property_id, &err)) { \
+      return (RETURN_TYPE)(0);                                                 \
+    }                                                                          \
+    return nullptr;                                                            \
+    return __kmpc_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 __kmpc_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();
+  }
+  interop_ptr = new omp_interop_val_t(device_id, interop_type);
+
+  if (device_id == omp_get_initial_device()) {
+    assert(device_id != omp_get_initial_device());
+    return;
+  }
+
+  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) {
+    __kmpc_omp_wait_deps(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                         noalias_dep_list);
+  }
+
+  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;
+    }
+  }
+}
+
+void __kmpc_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 __kmpc_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
     omp_const_mem_space                    DATA
     omp_high_bw_mem_space                  DATA
     omp_low_lat_mem_space                  DATA
     # 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
 
+__kmpc_get_target_task_waiting_list 2900
+__kmpc_set_async_info 2901
+__kmpc_target_task_yield 2902
+__kmpc_get_target_task_npredecessors 2903
+
 %ifndef stub
     # Ordinals between 900 and 999 are reserved
 
     # Ordinals between 1000 and 1999 are reserved
     # for user-callable uppercase Fortran entries.
 
 
     # ATOMIC entries