[OpenMP][AMDGPU] Add interop support for OpenMP AMD GPU plugin

[jz10]

Summary

Add interop related functionalities for OpenMP AMD GPU plugin, including get
async queue, get device reference and get backend runtime's ref ID.

Depends on D155692

[jdoerfert]

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

2698

No return after else.

2709

Reverse the cases, small case first.

[jz10]

Thanks Johannes, your comments were addressed in this revision

[jdoerfert]

Where are the tests?

openmp/libomptarget/include/omptargetplugin.h

167

Why don't we "just" return the int32_t? The indirection seems counterintuitive to me. Also, "DriverInfo" is not really descriptive. Can you specify what it returns, wrt. standard defined structs etc.

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

2689

Please don't cast to "int". If anything, pick a explicit size.

2698

No else after return (is what I meant). Also other places.

openmp/libomptarget/src/interop.cpp

91

This seems odd. Is the runtime and the vendor the same thing for us?

221

Do you know it's null or allocated here?

221–225

Why is the ready check happening after the rest, could you explain the problem here?

236

Why do we need this new function if we call init device info afterwards anyway? It seems we now have two functions we always call in sequence to get two numbers out of the plugin. If so, we should merge them.

237

Leftover?

[jz10]

Thanks Johannes, I addressed your comments as follows:

1. "Why don't we "just" return the int32_t? The indirection seems counterintuitive to me. Also, "DriverInfo" is not really descriptive. Can you specify what it returns, wrt. standard defined structs etc."

I revised the interface and API name

2. "Please don't cast to "int". If anything, pick a explicit size."

fixed

3. "No else after return (is what I meant). Also other places."

fixed

4. "This seems odd. Is the runtime and the vendor the same thing for us?"

this is the test code that was not removed, fixed

5. Do you know it's null or allocated here?

fixed, see below

6. "Why do we need this new function if we call init device info afterwards anyway? It seems we now have two functions we always call in sequence to get two numbers out of the plugin. If so, we should merge them."

I revised the implementation of "tgt_interop_init" based on its original workflow.
The original workflow for "tgt_interop_init" is :
i. handle task dependencies
ii. create omp_interop_val_t, but the interop type was constantly set as CUDA (this is why we need to retrieve the driver type from plugin)
iii. check if the device is ready, if not, set the error string but preserve the omp_interop_val_t ;
iv. call init_device_info to setup the Context and Device fields for tgt_device_info object, if this failed, then destroy omp_interop_val_t;
v. if the interop type is kmp_interop_type_tasksync, then call init_async_info to setup the async Queue for tgt_async_info object, if this failed, then destroy omp_interop_val_t.

based on these steps, we currently don't merge the omp_interop_val_t object creation together with init_device_info,

7. "Why is the ready check happening after the rest, could you explain the problem here?"

fixed

8. "Where is test case"

I added the test case as openmp/libomptarget/test/api/test_interop_amdgpu.c

[jdoerfert]

In D137607#3927376, @jz10 wrote:

6. "Why do we need this new function if we call init device info afterwards anyway? It seems we now have two functions we always call in sequence to get two numbers out of the plugin. If so, we should merge them."

I revised the implementation of "tgt_interop_init" based on its original workflow.
The original workflow for "tgt_interop_init" is :
i. handle task dependencies
ii. create omp_interop_val_t, but the interop type was constantly set as CUDA (this is why we need to retrieve the driver type from plugin)
iii. check if the device is ready, if not, set the error string but preserve the omp_interop_val_t ;
iv. call init_device_info to setup the Context and Device fields for tgt_device_info object, if this failed, then destroy omp_interop_val_t;
v. if the interop type is kmp_interop_type_tasksync, then call init_async_info to setup the async Queue for tgt_async_info object, if this failed, then destroy omp_interop_val_t.

So, keeping the original flow is not as important as designing it properly.
Right now, if the device is not ready, you still call into the device RTL.
Either we can always do that, and therefore can always call init_device_info to initialize the thing in the plugin, or we cannot always do that, and should consequently not call get_driver_type if the device is not ready.
To determine which, we need to look at the "non-ready" case and what that means. If the latter is the case, we can initialize the interop type with "unknown" instead of CUDA.
Does that make sense?

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

2708

When I said, no return after else, I mean, remove the "else" not the "return". You should take a look at the LLVM coding style guide.

if (...) {
  ...
  return A;
} else {
  ...
}
return B;

should be

if (...) {
  ...
  return A;
}
...
return B;

You still have the former pattern in your code and you now again have the complicated case first. It's easier to read if you have

if (...) {
  // short
  return;
}
//long

than

if (...) {
  // long
} else {
  // short
  return;
}

[jz10]

Thanks Johannes,

1. for the "else and return" thing, I fixed the code style based on you suggestion, please check if that works

2. for the workflow of __tgt_interop_init, I agree with you point and revised the workflow as:

i. handle task dependencies
ii. check if the device is ready, if not, set the interop object as omp_interop_none;
iii. create omp_interop_val_t that is corresponding the actual type of backend runtime, it has two types currently CUDA and HIP, other backend runtime support are not fully available yet;
iv. call init_device_info to setup the Context and Device fields for tgt_device_info object, if this failed, then destroy omp_interop_val_t;
v. if the interop type is kmp_interop_type_tasksync, then call init_async_info to setup the async Queue for tgt_async_info object, if this failed, then destroy omp_interop_val_t.

So the main changes is at step ii and iii, since the functionality of deviceIsReady is to check: 1. if the RTL object (backend runtime) was created or not ; 2. if the RTL initOnce was successfully invoked and the isInit flag was set, if neither of those 2 conditions satisfied, the interop can not provide any valuable information to user (e.g. the queue object or device information), so we just set interop object as omp_interop_none if the invocation of deviceIsReady failed;

3. "Either we can always do that, and therefore can always call init_device_info to initialize the thing in the plugin,"

This needs the plugin to be aware about the interop related data structure and APIs, I'm not sure if that is right way to manage interop related content

4. "or we cannot always do that, and should consequently not call get_driver_type if the device is not ready."

This is what we are doing so far.

[jdoerfert]

With the new changes, we return omp_interop_none a lot w/o telling the user what went wrong. Why not create an object and set the error string as we did before?
You also still have two entry points into the plugin, we do not need that. The omp_interop_val_t can be initialized by a single init call into the plugin. This will not only remove all the boilerplate of a new plugin API, but also cut down on the complexity of the initialization code in libomptarget.

[jz10]

Thanks Johannes

1. "With the new changes, we return omp_interop_none a lot w/o telling the user what went wrong. Why not create an object and set the error string as we did before?"

The check of deviceIsReady returns omp_interop_val_t object with error message if device was not ready

2. "You also still have two entry points into the plugin, we do not need that. The omp_interop_val_t can be initialized by a single init call into the plugin. This will not only remove all the boilerplate of a new plugin API, but also cut down on the complexity of the initialization code in libomptarget."

A new API "get_interop_val " creates omp_interop_val_t from plugin directly, so this also simplifies the code in "__tgt_interop_init". The API "get_driver_type" is removed.

[jdoerfert]

This is still too complex and inconsistent:

1. Remove the old plugin API.
2. Always allocate an interop object in one place, not in 3, default initialize it.
3. If something goes wrong, set an error message and return the object.
4. Pass the interop to the plugin to be initialized.

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

2692

Here you write *AsyncInfo but only below you allocate memory for it. Either there is an object and you overwrite it below or there is not and this will segfault.

openmp/libomptarget/src/interop.cpp

91

^^^

209

No error message here?

211

This API is now unused from the outside, right? Can we get rid of it?

[jz10]

Thanks Johannes

1. "Remove the old plugin API."

The "tgt_rtl_init_device_info" and "tgt_rtl_init_async_info" were from CUDA plugin, so I'm not sure if we should remove "__tgt_rtl_init_device_info" since it was not used ourside.

2. "Always allocate an interop object in one place, not in 3, default initialize it."

fixed

3. "If something goes wrong, set an error message and return the object."

fixed, but the checking of API symbols from dynamic lib still returns omp_interop_none,

4. "Pass the interop to the plugin to be initialized."

Made it, but here's one issue, the "vendor_id" and "backend_type_id" used to be const type, but we made them as non-constant now.

5. "Here you write *AsyncInfo but only below you allocate memory for it. Either there is an object and you overwrite it below or there is not and this will segfault."

fixed

6. openmp/libomptarget/src/interop.cpp ln;209 "No error message here?"

This checks if the dynamic lib was loaded successfully or not, if not, then no interop capacity is available, thus omp_interop_none should be enough

7. openmp/libomptarget/src/interop.cpp ln;211 "This API is now unused from the outside, right? Can we get rid of it?"

this is same as 1st issue above

[jdoerfert]

The "tgt_rtl_init_device_info" and "tgt_rtl_init_async_info" were from CUDA plugin, so I'm not sure if we should remove "__tgt_rtl_init_device_info" since it was not used ourside.

I'm now very confused what public APIs exist and which ones are actually needed/used. We should have a single one, but I'm fairly sure that's not the case, right? Why not. I don't understand the sentence I quoted. Maybe elaborate a little more if you think we should have more than one external plugin API.

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

2696

There is an assert followed by a check, that seems wrong. Either it's a pre-condition or it's something we want to handle gracefully.

2696

A nullptr queue is not an error code. It is a perfectly valid state meaning there is no outstanding async work. Check the CUDA impl, it will just return OFFLOAD_FAIL if the DeviceId is invalid, do the same.

2720–2726

There is an assert followed by a check, that seems wrong. Either it's a pre-condition or it's something we want to handle gracefully.

2729–2731

2742

We are actually not using hip at all. There is no backend or vendor part that is hip. We use HSA. Add a new enum value for both.

openmp/libomptarget/src/interop.cpp

69

Why the ( )?

openmp/libomptarget/test/api/omp_interop_amdgpu.c

78

Clang format this file please.

[jz10]

Thanks Johannes, here are replies

1. "I'm now very confused what public APIs exist and which ones are actually needed/used. We should have a single one, but I'm fairly sure that's not the case, right? Why not. I don't understand the sentence I quoted. Maybe elaborate a little more if you think we should have more than one external plugin API."

Sorry that brings this confusion. What I was trying to say is plugins/cuda has "tgt_rtl_init_async_info" and "tgt_rtl_init_device_info", so I added those two APIs to plugins/amdgpu accordingly. To my understanding of your suggestion (maybe wrong) , is to get rid of "__tgt_rtl_init_device_info" could be removed since it is not invoked at somewhere else

2. "There is an assert followed by a check, that seems wrong. Either it's a pre-condition or it's something we want to handle gracefully.

A nullptr queue is not an error code. It is a perfectly valid state meaning there is no outstanding async work. Check the CUDA impl, it will just return OFFLOAD_FAIL if the DeviceId is invalid, do the same.
There is an assert followed by a check, that seems wrong. Either it's a pre-condition or it's something we want to handle gracefully."

fixed

3. "We are actually not using hip at all. There is no backend or vendor part that is hip. We use HSA. Add a new enum value for both."

added a new enum as 'amdhsa', since 'hsa' has been used for namespace

4. "Why the ( )?"

fixed, but there are some other legacy code that still use ("...")

5. "Clang format this file please"

fixed

[jdoerfert]

, is to get rid of "__tgt_rtl_init_device_info" could be removed since it is not invoked at somewhere else

yes, please.

[jz10]

Thanks Johannes, I removed the __tgt_init_device_info API related definitions. Please check if that works.

[jdoerfert]

LG, assuming the test pass. Two minor nits below.

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

1873

This check (with the &) doesn't make much sense, did you mean w/o &?

1883

Why the parenthesis around InteropPtr and not the entire expression?

[mhalk]

Adaptation to plugins-nextgen by @jplehr.
Refactored test + rebase.

[jplehr]

The commit currently mentions me as an author, I believe that this should not be the case, as I did not do the original work and only part of the porting.

[ye-luo]

openmp/libomptarget/test/api/omp_interop_amdgpu.c

52

return value not used? No actual testing of interop?

openmp/libomptarget/include/interop.h

typedef struct kmp_tasking_flags { /* Total struct must be exactly 32 bits */
   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_hip_1 = 5,
+
+  omp_interop_backend_type_invalid_1 = 10
 } 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 {
+  invalid = 0,
   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) {}
+  omp_interop_val_t(intptr_t device_id, kmp_interop_type_t interop_type,
+                    omp_foreign_runtime_ids_t vendor_id,
+                    intptr_t backend_type_id)
+      : interop_type(interop_type), device_id(device_id), vendor_id(vendor_id),
+        backend_type_id(backend_type_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;
+  const omp_foreign_runtime_ids_t vendor_id;
+  const intptr_t backend_type_id;
 } omp_interop_val_t;
 
 #ifdef __cplusplus
 }
 #endif
 #endif

openmp/libomptarget/include/omptarget.h

int __tgt_target_kernel_nowait(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
                                int32_t ThreadLimit, void *HostPtr,
                                __tgt_kernel_arguments *Args, int32_t DepNum,
                                void *DepList, int32_t NoAliasDepNum,
                                void *NoAliasDepList);
 
 void __tgt_set_info_flag(uint32_t);
 
 int __tgt_print_device_info(int64_t DeviceId);
+  
 #ifdef __cplusplus
 }
 #endif
 
 #ifdef __cplusplus
 #define EXTERN extern "C"
 #else
 #define EXTERN extern
 #endif
 
 #endif // _OMPTARGET_H_

openmp/libomptarget/include/omptargetplugin.h

 //===-- omptargetplugin.h - Target dependent OpenMP Plugin API --*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines an interface between target independent OpenMP offload
 // runtime library libomptarget and target dependent plugin.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGETPLUGIN_H_
 #define _OMPTARGETPLUGIN_H_
 
+#include <interop.h>
 #include <omptarget.h>
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
 // First method called on the plugin
 int32_t __tgt_rtl_init_plugin();

 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);
 
 // Print the device information
 void __tgt_rtl_print_device_info(int32_t ID);
 
+// Retrieve the interop object
+omp_interop_val_t *__tgt_rtl_get_interop_val(int32_t DeviceI,
+                                             kmp_interop_type_t InteropType);
+
 // Event related interfaces. It is expected to use the interfaces in the
 // following way:
 // 1) Create an event on the target device (__tgt_rtl_create_event).
 // 2) Record the event based on the status of \p AsyncInfo->Queue at the moment
 // of function call to __tgt_rtl_record_event. An event becomes "meaningful"
 // once it is recorded, such that others can depend on it.
 // 3) Call __tgt_rtl_wait_event to set dependence on the event. Whether the
 // operation is blocking or non-blocking depends on the target. It is expected

openmp/libomptarget/include/rtl.h

 //===------------ rtl.h - Target independent OpenMP target RTL ------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Declarations for handling RTL plugins.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_RTL_H
 #define _OMPTARGET_RTL_H
 
 #include "omptarget.h"
+#include "interop.h"
 #include "llvm/ADT/DenseSet.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/Support/DynamicLibrary.h"
 
-#include "omptarget.h"
-
 #include <list>
 #include <map>
 #include <mutex>
 #include <string>
 
 // Forward declarations.
 struct DeviceTy;
 struct __tgt_bin_desc;
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 int32_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(print_device_info_ty)(int32_t);
+  typedef omp_interop_val_t *(get_interop_val_ty)(int32_t, kmp_interop_type_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 **);
#endif
   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;
   print_device_info_ty *print_device_info = nullptr;
+  get_interop_val_ty *get_interop_val = 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;

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

 
 void __tgt_rtl_print_device_info(int32_t DeviceId) {
   // TODO: Assertion to see if DeviceId is correct
   // NOTE: We don't need to set context for print device info.
 
   DeviceInfo().printDeviceInfo(DeviceId, DeviceInfo().HSAAgents[DeviceId]);
 }
 
+int32_t __tgt_rtl_init_async_info(int32_t DeviceId,
+                                  __tgt_async_info **AsyncInfo) {
+  assert(DeviceInfo().isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(AsyncInfo && "async_info is nullptr");
+
+  if (DeviceInfo().HSAQueueSchedulers.size() <= DeviceId) {
+    // If the cooresponding device is not available, then set the error code to
+    // Queue field
+    (*AsyncInfo)->Queue = nullptr;
+
+    return OFFLOAD_FAIL;
+  }
+
+  // Retrieve the queue from per−device based HSA queue scheduler
+  hsa_queue_t *Queue = DeviceInfo().HSAQueueSchedulers[DeviceId].next();
+  *AsyncInfo = new __tgt_async_info();
+  if (Queue != nullptr) {
+    (*AsyncInfo)->Queue = Queue;
+
+    return OFFLOAD_SUCCESS;
+  }
+
+  // If the queue object is not available , then set the error code to Queue
+  // field
+  (*AsyncInfo)->Queue = nullptr;
+
+  return OFFLOAD_FAIL;
+}
+
+int32_t __tgt_rtl_init_device_info(int32_t DeviceId,
+                                   __tgt_device_info *DeviceInfoPtr,
+                                   const char **ErrStr) {
+  assert(DeviceInfo().isValidDeviceId(DeviceId) && "device_id is invalid");
+  assert(DeviceInfoPtr && "device_info_ptr is nullptr");
+
+  DeviceInfoPtr->Context = nullptr;
+  if (DeviceInfo().HSAAgents.size() <= DeviceId) {
+    // If the cooresponding device is not available, then set device as null
+    DeviceInfoPtr->Device = nullptr;
+
+    return OFFLOAD_FAIL;
+  }
+
+  // Retrieve the HSA agent from per−device based HSA agent list
+  hsa_agent_t &Agent = DeviceInfo().HSAAgents[DeviceId];
+  if (!DeviceInfoPtr->Device) {
+    DeviceInfoPtr->Device = &Agent;
+  }
+
+  return OFFLOAD_SUCCESS;
+}
+
+omp_interop_val_t *__tgt_rtl_get_interop_val(int32_t DeviceId,
+                                             kmp_interop_type_t InteropType) {
+  // Create the HIP based interop value object
+  omp_interop_val_t *InteropPtr = new omp_interop_val_t(
+      DeviceId, InteropType, hip, omp_interop_backend_type_hip_1);
+
+  // Retrieve device information
+  if (__tgt_rtl_init_device_info(DeviceId, &(InteropPtr)->device_info,
+                                 &(InteropPtr)->err_str) != OFFLOAD_SUCCESS) {
+    // If the device information is not able to be retrieved, delete interop
+    // value object
+    delete InteropPtr;
+
+    InteropPtr = omp_interop_none;
+  }
+
+  return InteropPtr;
+}
+
 } // extern "C"

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

int32_t __tgt_rtl_init_device_info(int32_t DeviceId,
   assert(DeviceInfoPtr && "device_info_ptr is nullptr");
 
   if (DeviceRTL.setContext(DeviceId) != OFFLOAD_SUCCESS)
     return OFFLOAD_FAIL;
 
   return DeviceRTL.initDeviceInfo(DeviceId, DeviceInfoPtr, ErrStr);
 }
 
+omp_interop_val_t *__tgt_rtl_get_interop_val(int32_t DeviceId,
+                                             kmp_interop_type_t InteropType) {
+  // Create the CUDA based interop value object
+  omp_interop_val_t *InteropPtr = new omp_interop_val_t(
+      DeviceId, InteropType, cuda, omp_interop_backend_type_cuda_1);
+
+  // Retrieve device information
+  if (__tgt_rtl_init_device_info(DeviceId, &(InteropPtr)->device_info,
+                                 &(InteropPtr)->err_str) != OFFLOAD_SUCCESS) {
+    // If the device information is not able to be retrieved, delete interop
+    // value object
+    delete InteropPtr;
+
+    InteropPtr = omp_interop_none;
+  }
+
+  return InteropPtr;
+}
+
 #ifdef __cplusplus
 }
 #endif

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 "omptargetplugin.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:

 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 &InteropVal,
                                omp_interop_property_t Property, int *Err) {
   switch (Property) {
   case omp_ipr_fr_id:
     return InteropVal.backend_type_id;
   case omp_ipr_vendor:
     return InteropVal.vendor_id;
   case omp_ipr_device_num:
     return InteropVal.device_id;
   default:;
   }
   getTypeMismatch(Property, Err);
   return 0;
 }
 
 template <>
 const char *getProperty<const char *>(omp_interop_val_t &InteropVal,
                                       omp_interop_property_t Property,
                                       int *Err) {
   switch (Property) {
   case omp_ipr_fr_id:
     return InteropVal.interop_type == kmp_interop_type_tasksync
                ? "tasksync"
                : "device+context";
   case omp_ipr_vendor_name:
+  case omp_ipr_fr_name:
     return getVendorIdToStr(InteropVal.vendor_id);
   default:
     getTypeMismatch(Property, Err);
     return nullptr;
   }
 }
 
 template <>
if (DeviceId == -1) {
     DeviceId = omp_get_default_device();
   }
 
   if (InteropType == kmp_interop_type_tasksync) {
     __kmpc_omp_wait_deps(LocRef, Gtid, Ndeps, DepList, NdepsNoalias,
                          NoaliasDepList);
   }
 
-  InteropPtr = new omp_interop_val_t(DeviceId, InteropType);
+  DeviceTy &Device = *PM->Devices[DeviceId];
+  if (!Device.RTL || !Device.RTL->get_interop_val) {
+    InteropPtr = omp_interop_none;
+
+    return;
+  }
+
   if (!deviceIsReady(DeviceId)) {
+    // If the coresponding device is not ready yet, the interop object is
+    // initialized with invalid flag
+    InteropPtr = new omp_interop_val_t(DeviceId, InteropType, invalid,
+                                       omp_interop_backend_type_invalid_1);
     InteropPtr->err_str = "Device not ready!";
+
     return;
   }
 
-  DeviceTy &Device = *PM->Devices[DeviceId];
-  if (!Device.RTL || !Device.RTL->init_device_info ||
-      Device.RTL->init_device_info(DeviceId, &(InteropPtr)->device_info,
-                                   &(InteropPtr)->err_str)) {
-    delete InteropPtr;
-    InteropPtr = omp_interop_none;
-  }
+  // Retrieve the target specific interop value object
+  InteropPtr = Device.RTL->get_interop_val(DeviceId, InteropType);
+  if (InteropPtr == omp_interop_none)
+    return;
+
   if (InteropType == kmp_interop_type_tasksync) {
     if (!Device.RTL || !Device.RTL->init_async_info ||
         Device.RTL->init_async_info(DeviceId, &(InteropPtr)->async_info)) {
       delete InteropPtr;
       InteropPtr = omp_interop_none;
     }
   }
 }
 
 void __tgt_interop_use(ident_t *LocRef, kmp_int32 Gtid,
                        omp_interop_val_t *&InteropPtr, kmp_int32 DeviceId,
                        kmp_int32 Ndeps, kmp_depend_info_t *DepList,
                        kmp_int32 HaveNowait) {
   kmp_int32 NdepsNoalias = 0;
   kmp_depend_info_t *NoaliasDepList = NULL;
   assert(InteropPtr && "Cannot use nullptr!");
   omp_interop_val_t *InteropVal = InteropPtr;
   if (DeviceId == -1) {
     DeviceId = omp_get_default_device();
   }

openmp/libomptarget/src/rtl.cpp

#endif
   *((void **)&RTL.unregister_lib) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_unregister_lib");
   *((void **)&RTL.supports_empty_images) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_supports_empty_images");
   *((void **)&RTL.set_info_flag) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_set_info_flag");
   *((void **)&RTL.print_device_info) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_print_device_info");
+  *((void **)&RTL.get_interop_val) =
+      DynLibrary->getAddressOfSymbol("__tgt_rtl_get_interop_val");
   *((void **)&RTL.create_event) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_create_event");
   *((void **)&RTL.record_event) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_record_event");
   *((void **)&RTL.wait_event) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_wait_event");
   *((void **)&RTL.sync_event) =
       DynLibrary->getAddressOfSymbol("__tgt_rtl_sync_event");

openmp/libomptarget/test/api/omp_interop_amdgpu.c

+// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O1 -mllvm 
+// RUN: %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa                               
+// REQUIRES: amdgcn-amd-amdhsa  
+
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define N 16384
+
+void myVectorSet(int n, double s, double *x) {
+  for (int i = 0; i < n; ++ i) x[i] = s * (i + 1);
+}
+
+void myVectorCopy(int n, double *x, double *y) {
+  for (int i = 0; i < n; ++ i) y[i] = x[i];
+}
+
+void myDscal(int n, double s, double *x) {
+  for (int i = 0; i < n; ++ i) x[i] = s* x[i];
+}
+
+int main () {
+  const  double scalar = 2.0;
+  double x[N], y[N];
+  int dev;
+
+  omp_interop_t obj = omp_interop_none;
+  intptr_t type;
+
+  dev = omp_get_default_device();
+
+#pragma omp target nowait depend(out: x[0 : N]) map(from: x[0 : N]) device(dev)
+  myVectorSet(N, 1.0, x);
+
+#pragma omp task depend(out: y[0 : N])
+  myVectorSet(N , -1.0, y);
+
+#pragma omp interop init(targetsync: obj) device(dev) depend(in: x[0 : N]) depend(inout: y[0 : N]) // get obj 4 syncing
+
+  int   id      = (int)omp_get_interop_int(obj, omp_ipr_fr_id, NULL);
+  char* rt_name = (char*)omp_get_interop_str(obj, omp_ipr_fr_name, NULL);
+ 
+  if (obj != omp_interop_none && id == omp_ifr_hip) {
+    printf("OpenMP working with %s runtime to execute async memcpy.\n", rt_name);
+    int rc;
+    omp_get_interop_ptr(obj, omp_ipr_targetsync, &rc);
+
+    if (rc != omp_irc_success) {
+      fprintf(stderr, "ERROR: Failed to get %s stream, rt error = %d.\n", rt_name, rc);
+      if (rc == omp_irc_no_value)
+  fprintf(stderr, "Parameters valide, no meaningful value available.");
+      exit(1);
+    }
+
+    myVectorCopy(N, x, y);
+
+  } else { // Execute as OpenMP offload
+    printf("Notice: Offloading myCopy to perform memcpy.\n");
+
+#pragma omp target depend(inout: y[0 : N]) depend(in: x[0 : N]) nowait map(to: x[0 : N]) map(tofrom: y[0 : N]) device(dev)
+    myVectorCopy(N, x, y);
+  }
+
+  // This also ensures foreign tasks complete
+#pragma omp interop destroy(obj) nowait depend(out: y[0 : N])
+
+#pragma omp target depend(inout: x[0 : N])
+  myDscal(N, scalar, x);
+
+#pragma omp taskwait
+
+  printf("(1 : 16384) %f:%f\n", y[0], y[N-1]);
+  printf("(2 : 32768) %f:%f\n", x[0], x[N-1]);
+  
+  return 0;
+}