[OpenMP][libomptarget] Enable requires flags for target libraries.

[gtbercea]

Summary

Target link variables are currently implemented by creating a copy of the variables on the device side and unified memory never gets exploited.

When the prgram uses the:

#pragma omp requires unified_shared_memory

directive in conjunction with a declare target link, the linked variable is no longer allocated on the device and the host version is used instead.

This behavior is overridden by performing an explicit mapping.

A Clang side patch is required.

[gtbercea]

Remove line.

[gtbercea]

• Fix comment.

[gtbercea]

• Fix comment.

[ABataev]

libomptarget/include/omptarget.h

51

Do you really need this flag? Can you use something different? For example, OMP_TGT_MAPTYPE_LITERAL or OMP_TGT_MAPTYPE_PRIVATE?

libomptarget/plugins/cuda/src/rtl.cpp

454

8->sizeof(void*)

[gtbercea]

• Fix size.

[ABataev]

libomptarget/src/rtl.cpp

181

What should we do instead for other platforms?

[AlexEichenberger]

see comments

libomptarget/include/omptarget.h

51

I don't believe it is needed, as all is handled while registering the device library. If anything, you could just use the literal to naturally move the host pointer address to the device. This is if you want to avoid going through the indirection already set in place for globals mapped with "link" clause

libomptarget/plugins/cuda/src/rtl.cpp

454

Shouldn't this code be guarded by some "if the runtime is in unified memory mode," and if we see TARGET_LINK and the runtime is not in unified memory mode, should't there be some error as things are likely to go seriously wrong?

libomptarget/src/rtl.cpp

180

Same comment here... if the runtime is in non-unified memory mode, should't it be an error to see a link?

[gtbercea]

• Update patch.

[AlexEichenberger]

See suggested changes, should be pretty straightforward to do, let me know if you need help

libomptarget/include/omptarget.h

63

I wonder if it would not be wise to add an "undefined" flag. Initially, the internal state would be in undefined state. Then when the RT receives a "register_requires", one of two actions can happen.

If the RT state is in undefined state, it looks if the new requires are compatible with the current device; if yes, it accepts it, if not it either silently "degrade" the support to something supported, or abort.

If the RT state is not in undefined state, then it looks if the new require is compatible with the prior "register_required" emitted state. If incompatibilities are found, it abort.

Alternatively, one can add an additional bit in the RT state that indicates if one or more register requires have been reported. That second approach is cleaner but need one more variable.

libomptarget/plugins/cuda/src/rtl.cpp

234

See comment above, either set to unitialized, or let the state 0, and add additional variable to state "not initialized yet".

libomptarget/src/rtl.cpp

193

compatibility test required... minimally: accept if not previously initialized; refuse if different for any target related flags

[gtbercea]

• Fix checks and default flag.

[gtbercea]

• Update patch.

[ABataev]

libomptarget/include/omptargetplugin.h

35

Maybe it is better to make it target-independent rather than do this only for NVPTX target?

[AlexEichenberger]

LGTM, just made one minor suggestion

[AlexEichenberger]

response to Alexey's comment

libomptarget/include/omptargetplugin.h

35

I thought about this too, but say for one device, there is only unified memory... so we would actually not care about the requires unified.

So it's probably more flexible to have it there, but I am not opposed to take a stricter approach "flag must be consistent regardless of implementation"

[ABataev]

Also, it would be good to try to add the tests for this new functionality

[gtbercea]

Update

[ABataev]

libomptarget/plugins/cuda/src/rtl.cpp

115

I think you need to add some (at least basic) support for this in other plugins - generic-elf-64bit.

[gtbercea]

• Add function to generic elf.

[ABataev]

libomptarget/plugins/generic-elf-64bit/src/rtl.cpp

155

Where are you going to check that the provided RequiresFlags does not match the previously specified one? In this case, I think, the runtime should terminate execution of the program.

[gtbercea]

• Add basic functionality.

[gtbercea]

• Fix conditions.

[ABataev]

Still, I think you can add some basic tests at least to test error messages and the conditions. See llvm/projects/openmp/runtime/test/worksharing/for/kmp_sch_simd_guided.c as an example.

libomptarget/src/rtl.cpp

198

This can be an assert, I don't think that the compiler should call it with this flag

212

Should it be != here and in other places?

215

Better to emit it in terms of OpenMP pragmas and clauses here and in all other places

[gtbercea]

• Fix messages.

[gtbercea]

• Add test.

[gtbercea]

• Clean test.

[ABataev]

Tests for error messages

[gtbercea]

ping

[grokos]

libomptarget/src/device.h

107

0 -> OMP_REQ_UNDEFINED?

[gtbercea]

libomptarget/src/device.h

107

OMP_REQ_UNDEFINED is not in scope here and not needed. This 0 never gets used. In rtl.cpp this field will be overwritten with OMP_REQ_UNDEFINED if applicable.

[gtbercea]

• Fix test.

[gtbercea]

• Add comment.

[Hahnfeld]

In this patch __tgt_rtl_init_requires does nothing. Do we really want to add the prototype now? If yes, are we sure that we won't need to pass more information in the future? Once we add it, the runtime library must maintain compatibility.

Moreover, I'm not sure if we should require plugins to implement the function if they don't need any addition checks. I'd propose to make this optional, ie look for the symbol via dlsym and disable the call if it's not found.

libomptarget/test/offloading/requires.c

28–33

This will only work if the test compiler already supports requires. What's wrong with calling __tgt_register_requires twice as I suggested offline?

[gtbercea]

• Fix test to make it compatible with previous clang versions.

[gtbercea]

• Make implementation of init requires optional.

[gtbercea]

In D60223#1507576, @Hahnfeld wrote:

In this patch __tgt_rtl_init_requires does nothing. Do we really want to add the prototype now? If yes, are we sure that we won't need to pass more information in the future? Once we add it, the runtime library must maintain compatibility.

Moreover, I'm not sure if we should require plugins to implement the function if they don't need any addition checks. I'd propose to make this optional, ie look for the symbol via dlsym and disable the call if it's not found.

Should be optional now.

[grokos]

I don't have any other comments or suggestions, the patch looks good.

[Hahnfeld]

In D60223#1507576, @Hahnfeld wrote:

In this patch __tgt_rtl_init_requires does nothing. Do we really want to add the prototype now? If yes, are we sure that we won't need to pass more information in the future? Once we add it, the runtime library must maintain compatibility.

Ping. At the moment this patch just stores the RequiresFlags in the plugins and never uses it.

libomptarget/test/offloading/requires.c

26–27

Could you please run clang-format on this file?

[gtbercea]

In D60223#1510396, @Hahnfeld wrote:

In D60223#1507576, @Hahnfeld wrote:

In this patch __tgt_rtl_init_requires does nothing. Do we really want to add the prototype now? If yes, are we sure that we won't need to pass more information in the future? Once we add it, the runtime library must maintain compatibility.

Ping. At the moment this patch just stores the RequiresFlags in the plugins and never uses it.

If you check the declare target link patch it will be used there. I can move this part of the code to that patch.

[gtbercea]

• Format test.

[gtbercea]

• Remove requires init on plugin side.

[gtbercea]

@Hahnfeld just updated the patch. Please let me know if you have any more comments.

[Hahnfeld]

LG in case we don't need __tgt_register_requires to return a code in the near future (like for saying "that's not supported")

[Hahnfeld]

In D60223#1510738, @Hahnfeld wrote:

LG in case we don't need __tgt_register_requires to return a code in the near future (like for saying "that's not supported")

Sorry, mixed this up with __tgt_rtl_init_requires which is now gone. LGTM without restrictions :)

libomptarget/src/rtl.h

128–132

Nit: This shouldn't be a copy of the comment for RegisterLib

libomptarget/include/omptarget.h

enum tgt_map_type {
   OMP_TGT_MAPTYPE_RETURN_PARAM    = 0x040,
   // private variable - not mapped
   OMP_TGT_MAPTYPE_PRIVATE         = 0x080,
   // copy by value - not mapped
   OMP_TGT_MAPTYPE_LITERAL         = 0x100,
   // mapping is implicit
   OMP_TGT_MAPTYPE_IMPLICIT        = 0x200,
   // member of struct, member given by [16 MSBs] - 1
   OMP_TGT_MAPTYPE_MEMBER_OF       = 0xffff000000000000
 };
 
 enum OpenMPOffloadingDeclareTargetFlags {
   /// Mark the entry as having a 'link' attribute.
   OMP_DECLARE_TARGET_LINK = 0x01,
   /// Mark the entry as being a global constructor.
   OMP_DECLARE_TARGET_CTOR = 0x02,
   /// Mark the entry as being a global destructor.
   OMP_DECLARE_TARGET_DTOR = 0x04
 };
 
+enum OpenMPOffloadingRequiresDirFlags {
+  /// flag undefined.
+  OMP_REQ_UNDEFINED               = 0x000,
+  /// no requires directive present.
+  OMP_REQ_NONE                    = 0x001,
+  /// reverse_offload clause.
+  OMP_REQ_REVERSE_OFFLOAD         = 0x002,
+  /// unified_address clause.
+  OMP_REQ_UNIFIED_ADDRESS         = 0x004,
+  /// unified_shared_memory clause.
+  OMP_REQ_UNIFIED_SHARED_MEMORY   = 0x008,
+  /// dynamic_allocators clause.
+  OMP_REQ_DYNAMIC_ALLOCATORS      = 0x010
+};
+
 /// This struct is a record of an entry point or global. For a function
 /// entry point the size is expected to be zero
 struct __tgt_offload_entry {
   void *addr;   // Pointer to the offload entry info (function or global)
   char *name;   // Name of the function or global
   size_t size;  // Size of the entry info (0 if it is a function)
   int32_t flags; // Flags associated with the entry, e.g. 'link'.
   int32_t reserved; // Reserved, to be used by the runtime library.

 int omp_target_memcpy_rect(void *dst, void *src, size_t element_size,
     int num_dims, const size_t *volume, const size_t *dst_offsets,
     const size_t *src_offsets, const size_t *dst_dimensions,
     const size_t *src_dimensions, int dst_device, int src_device);
 int omp_target_associate_ptr(void *host_ptr, void *device_ptr, size_t size,
     size_t device_offset, int device_num);
 int omp_target_disassociate_ptr(void *host_ptr, int device_num);
 
+/// add the clauses of the requires directives in a given file
+void __tgt_register_requires(int64_t flags);
+
 /// adds a target shared library to the target execution image
 void __tgt_register_lib(__tgt_bin_desc *desc);
 
 /// removes a target shared library from the target execution image
 void __tgt_unregister_lib(__tgt_bin_desc *desc);
 
 // creates the host to target data mapping, stores it in the
 // libomptarget.so internal structure (an entry in a stack of data maps) and

libomptarget/include/omptargetplugin.h

 
 // Return an integer different from zero if the provided device image can be
 // supported by the runtime. The functionality is similar to comparing the
 // result of __tgt__rtl__load__binary to NULL. However, this is meant to be a
 // lightweight query to determine if the RTL is suitable for an image without
 // having to load the library, which can be expensive.
 int32_t __tgt_rtl_is_valid_binary(__tgt_device_image *Image);
 
+// Initialize the requires flags for the device.
+void __tgt_rtl_init_requires(int64_t RequiresFlags);
+
 // Initialize the specified device. In case of success return 0; otherwise
 // return an error code.
 int32_t __tgt_rtl_init_device(int32_t ID);
 
 // Pass an executable image section described by image to the specified
 // device and prepare an address table of target entities. In case of error,
 // return NULL. Otherwise, return a pointer to the built address table.
 // Individual entries in the table may also be NULL, when the corresponding

libomptarget/plugins/cuda/src/rtl.cpp

public:
   // OpenMP properties
   std::vector<int> NumTeams;
   std::vector<int> NumThreads;
 
   // OpenMP Environment properties
   int EnvNumTeams;
   int EnvTeamLimit;
 
+  // OpenMP Requires Flags
+  int64_t RequiresFlags;
+
   //static int EnvNumThreads;
   static const int HardTeamLimit = 1<<16; // 64k
   static const int HardThreadLimit = 1024;
   static const int DefaultNumTeams = 128;
   static const int DefaultNumThreads = 128;
 
   // Record entry point associated with device
   void addOffloadEntry(int32_t device_id, __tgt_offload_entry entry) {
#endif // OMPTARGET_DEBUG
     envStr = getenv("OMP_NUM_TEAMS");
     if (envStr) {
       // OMP_NUM_TEAMS has been set
       EnvNumTeams = std::stoi(envStr);
       DP("Parsed OMP_NUM_TEAMS=%d\n", EnvNumTeams);
     } else {
       EnvNumTeams = -1;
     }
+
+    // Default state.
+    RequiresFlags = OMP_REQ_UNDEFINED;
   }
 
   ~RTLDeviceInfoTy() {
     // Close modules
     for (auto &module : Modules)
       if (module) {
         CUresult err = cuModuleUnload(module);
         if (err != CUDA_SUCCESS) {

 #endif
 
 int32_t __tgt_rtl_is_valid_binary(__tgt_device_image *image) {
   return elf_check_machine(image, 190); // EM_CUDA = 190.
 }
 
 int32_t __tgt_rtl_number_of_devices() { return DeviceInfo.NumberOfDevices; }
 
+void __tgt_rtl_init_requires(int64_t RequiresFlags) {
+  DP("Init requires flags to %ld\n", RequiresFlags);
+  DeviceInfo.RequiresFlags = RequiresFlags;
+}
+
 int32_t __tgt_rtl_init_device(int32_t device_id) {
 
   CUdevice cuDevice;
   DP("Getting device %d\n", device_id);
   CUresult err = cuDeviceGet(&cuDevice, device_id);
   if (err != CUDA_SUCCESS) {
     DP("Error when getting CUDA device with id = %d\n", device_id);
     CUDA_ERR_STRING(err);
if (e->size) {
       DP("Entry point " DPxMOD " maps to global %s (" DPxMOD ")\n",
           DPxPTR(e - HostBegin), e->name, DPxPTR(cuptr));
       entry.addr = (void *)cuptr;
 
       DeviceInfo.addOffloadEntry(device_id, entry);
 
       continue;
     }
 
     CUfunction fun;
     err = cuModuleGetFunction(&fun, cumod, e->name);
 
     if (err != CUDA_SUCCESS) {
       DP("Loading '%s' (Failed)\n", e->name);
       CUDA_ERR_STRING(err);
       return NULL;
     }

libomptarget/plugins/exports

 VERS1.0 {
   global:
     __tgt_rtl_is_valid_binary;
     __tgt_rtl_number_of_devices;
+    __tgt_rtl_init_requires;
     __tgt_rtl_init_device;
     __tgt_rtl_load_binary;
     __tgt_rtl_data_alloc;
     __tgt_rtl_data_submit;
     __tgt_rtl_data_retrieve;
     __tgt_rtl_data_delete;
     __tgt_rtl_run_target_team_region;
     __tgt_rtl_run_target_region;
   local:
     *;
 };

libomptarget/plugins/generic-elf-64bit/src/rtl.cpp

 
 /// Class containing all the device information.
 class RTLDeviceInfoTy {
   std::vector<std::list<FuncOrGblEntryTy>> FuncGblEntries;
 
 public:
   std::list<DynLibTy> DynLibs;
 
+  // OpenMP Requires Flags
+  int64_t RequiresFlags;
+
   // Record entry point associated with device.
   void createOffloadTable(int32_t device_id, __tgt_offload_entry *begin,
                           __tgt_offload_entry *end) {
     assert(device_id < (int32_t)FuncGblEntries.size() &&
            "Unexpected device id!");
     FuncGblEntries[device_id].emplace_back();
     FuncOrGblEntryTy &E = FuncGblEntries[device_id].back();
 
public:
   RTLDeviceInfoTy(int32_t num_devices) {
 #ifdef OMPTARGET_DEBUG
     if (char *envStr = getenv("LIBOMPTARGET_DEBUG")) {
       DebugLevel = std::stoi(envStr);
     }
 #endif // OMPTARGET_DEBUG
 
     FuncGblEntries.resize(num_devices);
+
+    // Default state.
+    RequiresFlags = OMP_REQ_UNDEFINED;
   }
 
   ~RTLDeviceInfoTy() {
     // Close dynamic libraries
     for (auto &lib : DynLibs) {
       if (lib.Handle) {
         dlclose(lib.Handle);
         remove(lib.FileName);
#else
   return elf_check_machine(image, TARGET_ELF_ID);
 #endif
 }
 
 int32_t __tgt_rtl_number_of_devices() { return NUMBER_OF_DEVICES; }
 
 int32_t __tgt_rtl_init_device(int32_t device_id) { return OFFLOAD_SUCCESS; }
 
+void __tgt_rtl_init_requires(int64_t RequiresFlags) {
+  DP("Init requires flags to %ld\n", RequiresFlags);
+  DeviceInfo.RequiresFlags = RequiresFlags;
+}
+
 __tgt_target_table *__tgt_rtl_load_binary(int32_t device_id,
                                           __tgt_device_image *image) {
 
   DP("Dev %d: load binary from " DPxMOD " image\n", device_id,
      DPxPTR(image->ImageStart));
 
   assert(device_id >= 0 && device_id < NUMBER_OF_DEVICES && "bad dev id");
 

libomptarget/src/device.h

struct DeviceTy {
   PendingCtorsDtorsPerLibrary PendingCtorsDtors;
 
   ShadowPtrListTy ShadowPtrMap;
 
   std::mutex DataMapMtx, PendingGlobalsMtx, ShadowMtx;
 
   uint64_t loopTripCnt;
 
+  int64_t RTLRequiresFlags;
+
   DeviceTy(RTLInfoTy *RTL)
       : DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false), InitFlag(),
         HasPendingGlobals(false), HostDataToTargetMap(),
         PendingCtorsDtors(), ShadowPtrMap(), DataMapMtx(), PendingGlobalsMtx(),
-        ShadowMtx(), loopTripCnt(0) {}
+        ShadowMtx(), loopTripCnt(0), RTLRequiresFlags(0) {}
 
   // The existence of mutexes makes DeviceTy non-copyable. We need to
   // provide a copy constructor and an assignment operator explicitly.
   DeviceTy(const DeviceTy &d)
       : DeviceID(d.DeviceID), RTL(d.RTL), RTLDeviceID(d.RTLDeviceID),
         IsInit(d.IsInit), InitFlag(), HasPendingGlobals(d.HasPendingGlobals),
         HostDataToTargetMap(d.HostDataToTargetMap),
         PendingCtorsDtors(d.PendingCtorsDtors), ShadowPtrMap(d.ShadowPtrMap),
         DataMapMtx(), PendingGlobalsMtx(),
-        ShadowMtx(), loopTripCnt(d.loopTripCnt) {}
+        ShadowMtx(), loopTripCnt(d.loopTripCnt),
+        RTLRequiresFlags(d.RTLRequiresFlags) {}
 
   DeviceTy& operator=(const DeviceTy &d) {
     DeviceID = d.DeviceID;
     RTL = d.RTL;
     RTLDeviceID = d.RTLDeviceID;
     IsInit = d.IsInit;
     HasPendingGlobals = d.HasPendingGlobals;
     HostDataToTargetMap = d.HostDataToTargetMap;
     PendingCtorsDtors = d.PendingCtorsDtors;
     ShadowPtrMap = d.ShadowPtrMap;
     loopTripCnt = d.loopTripCnt;
+    RTLRequiresFlags = d.RTLRequiresFlags;
 
     return *this;
   }
 
   long getMapEntryRefCnt(void *HstPtrBegin);
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
       bool &IsNew, bool IsImplicit, bool UpdateRefCount = true);

libomptarget/src/device.cpp

DP("WARNING: Pointer is already mapped but section extends beyond mapped "
         "region\n");
   }
 
   return lr;
 }
 
 // Used by target_data_begin
 // Return the target pointer begin (where the data will be moved).
-// Allocate memory if this is the first occurrence if this mapping.
+// Allocate memory if this is the first occurrence of this mapping.
 // Increment the reference counter.
 // If NULL is returned, then either data allocation failed or the user tried
 // to do an illegal mapping.
 void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
     int64_t Size, bool &IsNew, bool IsImplicit, bool UpdateRefCount) {
   void *rc = NULL;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete) {
   }
 
   DataMapMtx.unlock();
   return rc;
 }
 
 /// Init device, should not be called directly.
 void DeviceTy::init() {
+  RTL->init_requires(RTLRequiresFlags);
   int32_t rc = RTL->init_device(RTLDeviceID);
   if (rc == OFFLOAD_SUCCESS) {
     IsInit = true;
   }
 }
 
 /// Thread-safe method to initialize the device only once.
 int32_t DeviceTy::initOnce() {

libomptarget/src/exports

 VERS1.0 {
   global:
+    __tgt_register_requires;
     __tgt_register_lib;
     __tgt_unregister_lib;
     __tgt_target_data_begin;
     __tgt_target_data_end;
     __tgt_target_data_update;
     __tgt_target;
     __tgt_target_teams;
     __tgt_target_data_begin_nowait;

libomptarget/src/interface.cpp

 static void HandleTargetOutcome(bool success) {
   switch (TargetOffloadPolicy) {
     case tgt_disabled:
       if (success) {
         FATAL_MESSAGE0(1, "expected no offloading while offloading is disabled");
       }
       break;
     case tgt_default:
-        FATAL_MESSAGE0(1, "default offloading policy must switched to " 
+        FATAL_MESSAGE0(1, "default offloading policy must switched to "
             "mandatory or disabled");
       break;
     case tgt_mandatory:
       if (!success) {
         FATAL_MESSAGE0(1, "failure of target construct while offloading is mandatory");
       }
       break;
   }
 }
 
 ////////////////////////////////////////////////////////////////////////////////
+/// adds requires flags
+EXTERN void __tgt_register_requires(int64_t flags) {
+  RTLs.RegisterRequires(flags);
+}
+
+////////////////////////////////////////////////////////////////////////////////
 /// adds a target shared library to the target execution image
 EXTERN void __tgt_register_lib(__tgt_bin_desc *desc) {
   RTLs.RegisterLib(desc);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// unloads a target shared library
 EXTERN void __tgt_unregister_lib(__tgt_bin_desc *desc) {

libomptarget/src/rtl.h

struct RTLInfoTy {
   typedef void *(data_alloc_ty)(int32_t, int64_t, void *);
   typedef int32_t(data_submit_ty)(int32_t, void *, void *, int64_t);
   typedef int32_t(data_retrieve_ty)(int32_t, void *, void *, int64_t);
   typedef int32_t(data_delete_ty)(int32_t, void *);
   typedef int32_t(run_region_ty)(int32_t, void *, void **, ptrdiff_t *,
                                  int32_t);
   typedef int32_t(run_team_region_ty)(int32_t, void *, void **, ptrdiff_t *,
                                       int32_t, int32_t, int32_t, uint64_t);
+  typedef int32_t(init_requires_ty)(int64_t);
 
   int32_t Idx;                     // 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;         // Number of devices this RTL deals with.
 
   void *LibraryHandler;
 
 #ifdef OMPTARGET_DEBUG
   std::string RTLName;
 #endif
 
   // Functions implemented in the RTL.
   is_valid_binary_ty *is_valid_binary;
   number_of_devices_ty *number_of_devices;
   init_device_ty *init_device;
   load_binary_ty *load_binary;
   data_alloc_ty *data_alloc;
   data_submit_ty *data_submit;
   data_retrieve_ty *data_retrieve;
   data_delete_ty *data_delete;
   run_region_ty *run_region;
   run_team_region_ty *run_team_region;
+  init_requires_ty *init_requires;
 
   // Are there images associated with this RTL.
   bool isUsed;
 
   // 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;
 
   // The existence of the mutex above makes RTLInfoTy non-copyable.
   // We need to provide a copy constructor explicitly.
   RTLInfoTy()
       : Idx(-1), NumberOfDevices(-1), LibraryHandler(0),
 #ifdef OMPTARGET_DEBUG
         RTLName(),
 #endif
         is_valid_binary(0), number_of_devices(0), init_device(0),
         load_binary(0), data_alloc(0), data_submit(0), data_retrieve(0),
-        data_delete(0), run_region(0), run_team_region(0), isUsed(false),
-        Mtx() {}
+        data_delete(0), run_region(0), run_team_region(0),
+        init_requires(0), isUsed(false), Mtx() {}
 
   RTLInfoTy(const RTLInfoTy &r) : Mtx() {
     Idx = r.Idx;
     NumberOfDevices = r.NumberOfDevices;
     LibraryHandler = r.LibraryHandler;
 #ifdef OMPTARGET_DEBUG
     RTLName = r.RTLName;
 #endif
     is_valid_binary = r.is_valid_binary;
     number_of_devices = r.number_of_devices;
     init_device = r.init_device;
     load_binary = r.load_binary;
     data_alloc = r.data_alloc;
     data_submit = r.data_submit;
     data_retrieve = r.data_retrieve;
     data_delete = r.data_delete;
     run_region = r.run_region;
     run_team_region = r.run_team_region;
+    init_requires = r.init_requires;
     isUsed = r.isUsed;
   }
 };
 
 /// RTLs identified in the system.
 class RTLsTy {
 private:
   // Mutex-like object to guarantee thread-safety and unique initialization
   // (i.e. the library attempts to load the RTLs (plugins) only once).
   std::once_flag initFlag;
   void LoadRTLs(); // not thread-safe
 
 public:
   // List of the detected runtime libraries.
   std::list<RTLInfoTy> AllRTLs;
 
   // Array of pointers to the detected runtime libraries that have compatible
   // binaries.
   std::vector<RTLInfoTy *> UsedRTLs;
 
+  int64_t RequiresFlags;
+
   explicit RTLsTy() {}
 
   // Register a shared library with all (compatible) RTLs.
+  void RegisterRequires(int64_t flags);
+
+  // Register a shared library with all (compatible) RTLs.
   void RegisterLib(__tgt_bin_desc *desc);
 
   // Unregister a shared library from all RTLs.
   void UnregisterLib(__tgt_bin_desc *desc);
 };
 extern RTLsTy RTLs;
 extern std::mutex RTLsMtx;
 
 

libomptarget/src/rtl.cpp

if (!(*((void**) &R.data_delete) = dlsym(
               dynlib_handle, "__tgt_rtl_data_delete")))
       continue;
     if (!(*((void**) &R.run_region) = dlsym(
               dynlib_handle, "__tgt_rtl_run_target_region")))
       continue;
     if (!(*((void**) &R.run_team_region) = dlsym(
               dynlib_handle, "__tgt_rtl_run_target_team_region")))
       continue;
+    if (!(*((void**) &R.init_requires) = dlsym(
+              dynlib_handle, "__tgt_rtl_init_requires")))
+      continue;
 
     // No devices are supported by this RTL?
     if (!(R.NumberOfDevices = R.number_of_devices())) {
       DP("No devices supported in this RTL\n");
       continue;
     }
 
     DP("Registering RTL %s supporting %d devices!\n",
for (__tgt_offload_entry *entry = img->EntriesBegin;
         Device.PendingCtorsDtors[desc].PendingCtors.push_back(entry->addr);
       } else if (entry->flags & OMP_DECLARE_TARGET_DTOR) {
         // Dtors are pushed in reverse order so they are executed from end
         // to beginning when unregistering the library!
         DP("Adding dtor " DPxMOD " to the pending list.\n",
             DPxPTR(entry->addr));
         Device.PendingCtorsDtors[desc].PendingDtors.push_front(entry->addr);
       }
 
       if (entry->flags & OMP_DECLARE_TARGET_LINK) {
         DP("The \"link\" attribute is not yet supported!\n");
       }
     }
     Device.PendingGlobalsMtx.unlock();
   }
 }
 
+void RTLsTy::RegisterRequires(int64_t flags) {
+  // TODO: add more elaborate check.
+  // Minimal check: only set requires flags if previous value
+  // is undefined. This ensures that only the first call to this
+  // function will set the requires flags. All subsequent calls
+  // will be checked for compatibility.
+  assert(flags != OMP_REQ_UNDEFINED &&
+         "illegal undefined flag for requires directive!");
+  if (RequiresFlags == OMP_REQ_UNDEFINED) {
+    RequiresFlags = flags;
+    return;
+  }
+
+  // If multiple compilation units are present enforce
+  // consistency across all of them for require clauses:
+  //  - reverse_offload
+  //  - unified_address
+  //  - unified_shared_memory
+  if ((RequiresFlags & OMP_REQ_REVERSE_OFFLOAD) !=
+      (flags & OMP_REQ_REVERSE_OFFLOAD)) {
+    FATAL_MESSAGE0(1,
+        "'#pragma omp requires reverse_offload' not used consistently!");
+  }
+  if ((RequiresFlags & OMP_REQ_UNIFIED_ADDRESS) !=
+          (flags & OMP_REQ_UNIFIED_ADDRESS)) {
+    FATAL_MESSAGE0(1,
+        "'#pragma omp requires unified_address' not used consistently!");
+  }
+  if ((RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) !=
+          (flags & OMP_REQ_UNIFIED_SHARED_MEMORY)) {
+    FATAL_MESSAGE0(1,
+        "'#pragma omp requires unified_shared_memory' not used consistently!");
+  }
+
+  // TODO: insert any other missing checks
+
+  DP("New requires flags %ld compatible with existing %ld!\n",
+     flags, RequiresFlags);
+}
+
 void RTLsTy::RegisterLib(__tgt_bin_desc *desc) {
   // Attempt to load all plugins available in the system.
   std::call_once(initFlag, &RTLsTy::LoadRTLs, this);
 
   RTLsMtx.lock();
   // Register the images with the RTLs that understand them, if any.
   for (int32_t i = 0; i < desc->NumDeviceImages; ++i) {
     // Obtain the image.
for (auto &R : RTLs.AllRTLs) {
         size_t start = Devices.size();
         Devices.resize(start + R.NumberOfDevices, device);
         for (int32_t device_id = 0; device_id < R.NumberOfDevices;
             device_id++) {
           // global device ID
           Devices[start + device_id].DeviceID = start + device_id;
           // RTL local device ID
           Devices[start + device_id].RTLDeviceID = device_id;
+          // RTL requires flags
+          Devices[start + device_id].RTLRequiresFlags = RequiresFlags;
         }
 
         // Initialize the index of this RTL and save it in the used RTLs.
         R.Idx = (RTLs.UsedRTLs.empty())
                     ? 0
                     : RTLs.UsedRTLs.back()->Idx +
                           RTLs.UsedRTLs.back()->NumberOfDevices;
         assert((size_t) R.Idx == start &&

libomptarget/test/offloading/requires.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
+
+/*
+  Test for the 'requires' clause check.
+  When a target region is used, the requires flags are set in the
+  runtime for the entire compilation unit. If the flags are set again,
+  (for whatever reason) the set must be consistent with previously
+  set values.
+*/
+#include <stdio.h>
+#include <omp.h>
+
+// ---------------------------------------------------------------------------
+// Various definitions copied from OpenMP RTL
+
+extern void __tgt_register_requires(int64_t);
+
+// End of definitions copied from OpenMP RTL.
+// ---------------------------------------------------------------------------
+
+void run_reg_requires()
+{
+  // Before the target region is registered, the requires registers the status
+  // of the requires clauses. Since there are no requires clauses in this file
+  // the flags state can only be OMP_REQ_NONE i.e. 1.
+
+  // This is the 2nd time this function is called so it should print the debug
+  // info belonging to the check.
+  __tgt_register_requires(1);
+  __tgt_register_requires(1);
+  // DEBUG: New requires flags 1 compatible with existing 1!
+}
+
+// ---------------------------------------------------------------------------
+int main()
+{
+  run_reg_requires();
+
+  // This also runs reg requires for the first time.
+  #pragma omp target
+  {}
+
+  return 0;
+}