[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

443 ↗

(On Diff #193558)

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

443 ↗

(On Diff #193558)

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 ↗

(On Diff #194191)

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

190

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 ↗

(On Diff #195399)

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 ↗

(On Diff #195399)

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 ↗

(On Diff #195800)

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

149 ↗

(On Diff #196892)

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

195

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

209

Should it be != here and in other places?

212

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–129

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/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);

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

 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 the clauses of the requires directive.
+  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

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;
+}
 No newline at end of file