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

79

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

1238–1239

openmp/libomptarget/src/CMakeLists.txt

21

order please

openmp/libomptarget/src/exports

51

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

113

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

79

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

113

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

openmp/runtime/src/dllexports

587

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

113

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

[jdoerfert]

openmp/libomptarget/src/private.h

113

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

113

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

556

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

1251–1257

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

556

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

556

... 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

556

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;
+
+typedef enum omp_foreign_runtime_ids {
+  cuda = 1,
+  cuda_driver = 2,
+  opencl = 3,
+  sycl = 4,
+  hip = 5,
+  level_zero = 6,
+} omp_foreign_runtime_ids_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 omp_foreign_runtime_ids_t vendor_id = cuda;
+  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

 int32_t __tgt_rtl_wait_event(int32_t ID, void *Event,
                              __tgt_async_info *AsyncInfo);
 
 int32_t __tgt_rtl_sync_event(int32_t ID, void *Event);
 
 int32_t __tgt_rtl_destroy_event(int32_t ID, void *Event);
 // }
 
+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/include/rtl.h

struct RTLInfoTy {
   typedef int32_t(supports_empty_images_ty)();
   typedef void(print_device_info_ty)(int32_t);
   typedef void(set_info_flag_ty)(uint32_t);
   typedef int32_t(create_event_ty)(int32_t, void **);
   typedef int32_t(record_event_ty)(int32_t, void *, __tgt_async_info *);
   typedef int32_t(wait_event_ty)(int32_t, void *, __tgt_async_info *);
   typedef int32_t(sync_event_ty)(int32_t, void *);
   typedef int32_t(destroy_event_ty)(int32_t, void *);
+  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
   supports_empty_images_ty *supports_empty_images = nullptr;
   set_info_flag_ty *set_info_flag = nullptr;
   print_device_info_ty *print_device_info = nullptr;
   create_event_ty *create_event = nullptr;
   record_event_ty *record_event = nullptr;
   wait_event_ty *wait_event = nullptr;
   sync_event_ty *sync_event = nullptr;
   destroy_event_ty *destroy_event = 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/plugins/cuda/src/rtl.cpp

 /// Keep entries table per device.
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
   std::vector<__tgt_offload_entry> Entries;
 };
 
 /// Use a single entity to encode a kernel and a set of flags.
 struct KernelTy {
   CUfunction Func;
 
   // execution mode of kernel
   llvm::omp::OMPTgtExecModeFlags ExecutionMode;
 
   /// Maximal number of threads per block for this kernel.
   int MaxThreadsPerBlock = 0;
 
   KernelTy(CUfunction _Func, llvm::omp::OMPTgtExecModeFlags _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) {
       CUstream S;
       if (StreamPool[DeviceId]->acquire(S) != OFFLOAD_SUCCESS)
         return nullptr;
 
       AsyncInfo->Queue = S;
     }
 
     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),
         DynamicMemorySize(0), EventPool(AllocatorTy<CUevent>()) {
void printDeviceInfo(int32_t device_id) {
     checkResult(cuDeviceGet(&Device, device_id),
                 "Error returned from cuCtxGetDevice\n");
 
     cuDriverGetVersion(&TmpInt);
     printf("    CUDA Driver Version: \t\t%d \n", TmpInt);
     printf("    CUDA Device Number: \t\t%d \n", device_id);
     checkResult(cuDeviceGetName(TmpChar, 1000, Device),
                 "Error returned from cuDeviceGetName\n");
     printf("    Device Name: \t\t\t%s \n", TmpChar);
     checkResult(cuDeviceTotalMem(&TmpSt, Device),
                 "Error returned from cuDeviceTotalMem\n");
     printf("    Global Memory Size: \t\t%zu bytes \n", TmpSt);
     checkResult(cuDeviceGetAttribute(
                     &TmpInt, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, Device),
                 "Error returned from cuDeviceGetAttribute\n");
     printf("    Number of Multiprocessors: \t\t%d \n", TmpInt);
     checkResult(
         cuDeviceGetAttribute(&TmpInt, CU_DEVICE_ATTRIBUTE_GPU_OVERLAP, Device),
         "Error returned from cuDeviceGetAttribute\n");
     printf("    Concurrent Copy and Execution: \t%s \n", BOOL2TEXT(TmpInt));
     checkResult(cuDeviceGetAttribute(
                     &TmpInt, CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, Device),
                 "Error returned from cuDeviceGetAttribute\n");
     printf("    Total Constant Memory: \t\t%d bytes\n", TmpInt);
     checkResult(
         cuDeviceGetAttribute(
             &TmpInt, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK, Device),
         "Error returned from cuDeviceGetAttribute\n");
     printf("    Max Shared Memory per Block: \t%d bytes \n", TmpInt);
     checkResult(
         cuDeviceGetAttribute(
             &TmpInt, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK, Device),
         "Error returned from cuDeviceGetAttribute\n");
     printf("    Registers per Block: \t\t%d \n", TmpInt);
     checkResult(
         cuDeviceGetAttribute(&TmpInt, CU_DEVICE_ATTRIBUTE_WARP_SIZE, Device),
if (Err != CUDA_SUCCESS) {
       DP("Error when waiting event. stream = " DPxMOD ", event = " DPxMOD "\n",
          DPxPTR(Stream), DPxPTR(Event));
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }
 
     return OFFLOAD_SUCCESS;
   }
+
+  int releaseAsyncInfo(int DeviceId, __tgt_async_info *AsyncInfo) const {
+    if (AsyncInfo->Queue) {
+      StreamPool[DeviceId]->release(
+          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(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_destroy_event(int32_t device_id, void *event_ptr) {
   assert(event_ptr && "event is nullptr");
 
   return DeviceRTL.destroyEvent(event_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);
+}
+
+int32_t __tgt_rtl_init_device_info(int32_t device_id,
+                                   __tgt_device_info *device_info_ptr,
+                                   const char **err_str) {
+  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, err_str);
+}
+
 #ifdef __cplusplus
 }
 #endif

openmp/libomptarget/plugins/exports

global:
     __tgt_rtl_supports_empty_images;
     __tgt_rtl_set_info_flag;
     __tgt_rtl_print_device_info;
     __tgt_rtl_create_event;
     __tgt_rtl_record_event;
     __tgt_rtl_wait_event;
     __tgt_rtl_sync_event;
     __tgt_rtl_destroy_event;
+    __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_init_all_rtls;
     __tgt_target_data_begin;
     __tgt_target_data_end;
     __tgt_target_data_update;
global:
     omp_target_associate_ptr;
     omp_target_disassociate_ptr;
     llvm_omp_target_alloc_host;
     llvm_omp_target_alloc_shared;
     llvm_omp_target_alloc_device;
     llvm_omp_get_dynamic_shared;
     __tgt_set_info_flag;
     __tgt_print_device_info;
+    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

+//===---------------interop.cpp - Implementation of interop directive -----===//
+//
+// 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 getPropertyErrorType(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 getTypeMismatch(omp_interop_property_t Property, int *Err) {
+  if (Err)
+    *Err = getPropertyErrorType(Property);
+}
+
+const char *getVendorIdToStr(const omp_foreign_runtime_ids_t VendorId) {
+  switch (VendorId) {
+  case cuda:
+    return ("cuda");
+  case cuda_driver:
+    return ("cuda_driver");
+  case opencl:
+    return ("opencl");
+  case sycl:
+    return ("sycl");
+  case hip:
+    return ("hip");
+  case level_zero:
+    return ("level_zero");
+  }
+  return ("unknown");
+}
+
+template <typename PropertyTy>
+PropertyTy getProperty(omp_interop_val_t &InteropVal,
+                       omp_interop_property_t Property, int *Err);
+
+template <>
+intptr_t getProperty<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:;
+  }
+  getTypeMismatch(property, err);
+  return 0;
+}
+
+template <>
+const char *getProperty<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 getVendorIdToStr(interop_val.vendor_id);
+  default:
+    getTypeMismatch(property, err);
+    return nullptr;
+  }
+}
+
+template <>
+void *getProperty<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:;
+  }
+  getTypeMismatch(property, err);
+  return nullptr;
+}
+
+bool getPropertyCheck(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 (!getPropertyCheck(&interop_val, property_id, err)) {                   \
+      return (RETURN_TYPE)(0);                                                 \
+    }                                                                          \
+    return getProperty<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 (!getPropertyCheck(&interop_val, property_id, &err)) {                  \
+      return (RETURN_TYPE)(0);                                                 \
+    }                                                                          \
+    return nullptr;                                                            \
+    return getProperty<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 (!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;
+  }
+  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;
+    }
+  }
+}
+
+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;
+  if (device_id == -1) {
+    device_id = omp_get_default_device();
+  }
+  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 (!device_is_ready(device_id)) {
+    interop_ptr->err_str = "Device not ready!";
+    return;
+  }
+
+  if (interop_val->interop_type == kmp_interop_type_tasksync) {
+    __kmpc_omp_wait_deps(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                         noalias_dep_list);
+  }
+  // TODO Flush the queue associated with the interop through the plugin
+}
+
+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 (device_id == -1) {
+    device_id = omp_get_default_device();
+  }
+
+  if (interop_val == omp_interop_none)
+    return;
+
+  assert((device_id == -1 || interop_val->device_id == device_id) &&
+         "Inconsistent device-id usage!");
+  if (!device_is_ready(device_id)) {
+    interop_ptr->err_str = "Device not ready!";
+    return;
+  }
+
+  if (interop_val->interop_type == kmp_interop_type_tasksync) {
+    __kmpc_omp_wait_deps(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                         noalias_dep_list);
+  }
+  // TODO Flush the queue associated with the interop through the plugin
+  // TODO Signal out dependences
+
+  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_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.cpp

#endif
     *((void **)&R.create_event) =
         dlsym(dynlib_handle, "__tgt_rtl_create_event");
     *((void **)&R.record_event) =
         dlsym(dynlib_handle, "__tgt_rtl_record_event");
     *((void **)&R.wait_event) = dlsym(dynlib_handle, "__tgt_rtl_wait_event");
     *((void **)&R.sync_event) = dlsym(dynlib_handle, "__tgt_rtl_sync_event");
     *((void **)&R.destroy_event) =
         dlsym(dynlib_handle, "__tgt_rtl_destroy_event");
+    *((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/libomptarget/test/offloading/interop.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: nvptx64-nvidia-cuda
+
+#include <assert.h>
+#include <omp.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+typedef void *cudaStream_t;
+
+int main() {
+
+  int device_id = omp_get_default_device();
+
+#pragma omp parallel master
+  {
+
+    double D0, D2;
+    omp_interop_t interop;
+
+#pragma omp interop init(targetsync : interop) device(device_id) nowait
+    assert(interop != NULL);
+
+    int err;
+    for (int i = omp_ipr_first; i < 0; i++) {
+      const char *n =
+          omp_get_interop_name(interop, (omp_interop_property_t)(i));
+      long int li =
+          omp_get_interop_int(interop, (omp_interop_property_t)(i), &err);
+      const void *p =
+          omp_get_interop_ptr(interop, (omp_interop_property_t)(i), &err);
+      const char *s =
+          omp_get_interop_str(interop, (omp_interop_property_t)(i), &err);
+      const char *n1 =
+          omp_get_interop_type_desc(interop, (omp_interop_property_t)(i));
+    }
+#pragma omp interop use(interop) depend(in : D0, D2)
+
+    cudaStream_t stream =
+        (omp_get_interop_ptr(interop, omp_ipr_targetsync, NULL));
+    assert(stream != NULL);
+
+#pragma omp interop destroy(interop) depend(in : D0, D2) device(device_id)
+  }
+  printf("PASS\n");
+}
+// CHECK: PASS

openmp/runtime/src/dllexports

kmp_set_disp_num_buffers                    890
     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_aligned_alloc                       778
     omp_aligned_calloc                      806
+    omp_get_interop_int                     2514
+    omp_get_interop_ptr                     2515
+    omp_get_interop_str                     2516
 
     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

 #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;
+
+// libomptarget, if loaded, provides this function
+int FTN_STDCALL FTN_GET_NUM_INTEROP_PROPERTIES(const omp_interop_t interop) {
+#if KMP_MIC || KMP_OS_DARWIN || defined(KMP_STUB)
+  return 0;
+#else
+  int (*fptr)(const omp_interop_t);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_num_interop_properties")))
+    return (*fptr)(interop);
+  return 0;
+#endif // KMP_MIC || KMP_OS_DARWIN || KMP_OS_WINDOWS || defined(KMP_STUB)
+}
+
+/// TODO Convert FTN_GET_INTEROP_XXX functions into a macro like interop.cpp
+// libomptarget, if loaded, provides this function
+intptr_t FTN_STDCALL FTN_GET_INTEROP_INT(const omp_interop_t interop,
+                                         omp_interop_property_t property_id,
+                                         int *err) {
+  intptr_t (*fptr)(const omp_interop_t, omp_interop_property_t, int *);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_interop_int")))
+    return (*fptr)(interop, property_id, err);
+  return 0;
+}
+
+// libomptarget, if loaded, provides this function
+void *FTN_STDCALL FTN_GET_INTEROP_PTR(const omp_interop_t interop,
+                                      omp_interop_property_t property_id,
+                                      int *err) {
+  void *(*fptr)(const omp_interop_t, omp_interop_property_t, int *);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_interop_ptr")))
+    return (*fptr)(interop, property_id, err);
+  return nullptr;
+}
+
+// libomptarget, if loaded, provides this function
+const char *FTN_STDCALL FTN_GET_INTEROP_STR(const omp_interop_t interop,
+                                            omp_interop_property_t property_id,
+                                            int *err) {
+  const char *(*fptr)(const omp_interop_t, omp_interop_property_t, int *);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_interop_str")))
+    return (*fptr)(interop, property_id, err);
+  return nullptr;
+}
+
+// libomptarget, if loaded, provides this function
+const char *FTN_STDCALL FTN_GET_INTEROP_NAME(
+    const omp_interop_t interop, omp_interop_property_t property_id) {
+  const char *(*fptr)(const omp_interop_t, omp_interop_property_t);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_interop_name")))
+    return (*fptr)(interop, property_id);
+  return nullptr;
+}
+
+// libomptarget, if loaded, provides this function
+const char *FTN_STDCALL FTN_GET_INTEROP_TYPE_DESC(
+    const omp_interop_t interop, omp_interop_property_t property_id) {
+  const char *(*fptr)(const omp_interop_t, omp_interop_property_t);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_interop_type_desc")))
+    return (*fptr)(interop, property_id);
+  return nullptr;
+}
+
+// libomptarget, if loaded, provides this function
+const char *FTN_STDCALL FTN_GET_INTEROP_RC_DESC(
+    const omp_interop_t interop, omp_interop_property_t property_id) {
+  const char *(*fptr)(const omp_interop_t, omp_interop_property_t);
+  if ((*(void **)(&fptr) = KMP_DLSYM_NEXT("omp_get_interop_rec_desc")))
+    return (*fptr)(interop, 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_
 

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_