[OpenMP][libomptarget] Emit runtime message when variable is incorrectly mapped to device

[doru1004]

Summary

Enable the runtime to emit a warning message when a mapped variable does not have an appropriate device counterpart that can be used. The current patch uses the use_device_addr as an example but other cases may be possible. The patch has been tested with the warning as an error to make sure that it does not disturb any of the existing tests. The warning currently uses MESSAGE but can be changed to REPORT if reviewers think it would be more appropriate.

[RaviNarayanaswamy]

I am not sure if we should generate this message. The spec says "Each list item must have a corresponding list item in the device data environment or be accessible on the target device."
The pointer may be in the unified_shared_memory.

[trws]

The message seems useful, accidentally putting x there rather than x[:0] like the user probably intended is likely to be a relatively common mistake unfortunately. Would it be feasible to, when the message would be printed for a pointer-type variable, check if it's target is accessible and suggest the appropriate fix, or at least say "maybe you meant to use an array section, the array section target is valid"? That would make it immensely easier for an lower experience user to take action on.

[mjklemm]

In D134845#3827023, @trws wrote:

The message seems useful, accidentally putting x there rather than x[:0] like the user probably intended is likely to be a relatively common mistake unfortunately. Would it be feasible to, when the message would be printed for a pointer-type variable, check if it's target is accessible and suggest the appropriate fix, or at least say "maybe you meant to use an array section, the array section target is valid"? That would make it immensely easier for an lower experience user to take action on.

I would agree to that sentiment. If the runtime can easily detect the issue and report, it should do so. Just letting the code die with a SEGFAULT is certainly fun for folks who love to debug, but is not really productive.

[RaviNarayanaswamy]

This message could be printed when LIBOMPTARGET_INFO="something". I don't think it should be printed by default as it could be false positive.

[doru1004]

In D134845#3827802, @RaviNarayanaswamy wrote:

This message could be printed when LIBOMPTARGET_INFO="something". I don't think it should be printed by default as it could be false positive.

Happy to enable it as an INFO.

In D134845#3825276, @RaviNarayanaswamy wrote:

I am not sure if we should generate this message. The spec says "Each list item must have a corresponding list item in the device data environment or be accessible on the target device."
The pointer may be in the unified_shared_memory.

We don't emit the message when IsHostPtr is true (i.e. when USM is enabled and not under close modifier).

[ronlieb]

@RaviNarayanaswamy can this one move forward to an LOTM Looks Ok To Me ? or better ?

[RaviNarayanaswamy]

Minor spelling otherwise LGTM

openmp/libomptarget/include/Debug.h

58

cotunerpart -> counterpart

[ronlieb]

@RaviNarayanaswamy LGTM, please formally approve if you are satsified , thx

[RaviNarayanaswamy]

LGTM

[doru1004]

Commit: a91cb9ce39dc42e6a7a2c4fe97580e51eb1c2961

openmp/libomptarget/include/Debug.h

enum OpenMPInfoType : uint32_t {
   // Dump the contents of the device pointer map at kernel exit or failure.
   OMP_INFOTYPE_DUMP_TABLE = 0x0004,
   // Indicate when an address is added to the device mapping table.
   OMP_INFOTYPE_MAPPING_CHANGED = 0x0008,
   // Print kernel information from target device plugins.
   OMP_INFOTYPE_PLUGIN_KERNEL = 0x0010,
   // Print whenever data is transferred to the device
   OMP_INFOTYPE_DATA_TRANSFER = 0x0020,
+  // Print whenever data does not have a viable device counterpart.
+  OMP_INFOTYPE_EMPTY_MAPPING = 0x0040,
   // Enable every flag.
   OMP_INFOTYPE_ALL = 0xffffffff,
 };
 
 inline std::atomic<uint32_t> &getInfoLevelInternal() {
   static std::atomic<uint32_t> InfoLevel;
   static std::once_flag Flag{};
   std::call_once(Flag, []() {

openmp/libomptarget/include/device.h

struct FlagTy {
     unsigned IsNewEntry : 1;
     /// If the pointer is actually a host pointer (when unified memory enabled)
     unsigned IsHostPointer : 1;
     /// If the pointer is present in the mapping table.
     unsigned IsPresent : 1;
     /// Flag indicating that this was the last user of the entry and the ref
     /// count is now 0.
     unsigned IsLast : 1;
-  } Flags = {0, 0, 0, 0};
+    /// If the pointer is contained.
+    unsigned IsContained : 1;
+  } Flags = {0, 0, 0, 0, 0};
 
   TargetPointerResultTy(const TargetPointerResultTy &) = delete;
   TargetPointerResultTy &operator=(const TargetPointerResultTy &TPR) = delete;
   TargetPointerResultTy() {}
 
   TargetPointerResultTy(FlagTy Flags, HostDataToTargetTy *Entry,
                         void *TargetPointer)
       : Flags(Flags), TargetPointer(TargetPointer), Entry(Entry) {
~TargetPointerResultTy() {
     if (Entry)
       Entry->unlock();
   }
 
   bool isPresent() const { return Flags.IsPresent; }
 
   bool isHostPointer() const { return Flags.IsHostPointer; }
 
+  bool isContained() const { return Flags.IsContained; }
+
   /// The corresponding target pointer
   void *TargetPointer = nullptr;
 
   HostDataToTargetTy *getEntry() const { return Entry; }
   void setEntry(HostDataToTargetTy *HDTTT,
                 HostDataToTargetTy *OwnedTPR = nullptr) {
     if (Entry)
       Entry->unlock();

openmp/libomptarget/src/omptarget.cpp

if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
           REPORT("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
         if (PointerTpr.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=
             OFFLOAD_SUCCESS)
           return OFFLOAD_FAIL;
       }
     }
+
+    // Check if variable can be used on the device:
+    bool IsStructMember = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
+    if (getInfoLevel() & OMP_INFOTYPE_EMPTY_MAPPING && ArgTypes[I] != 0 &&
+        !IsStructMember && !IsImplicit && !TPR.isPresent() &&
+        !TPR.isContained() && !TPR.isHostPointer())
+      INFO(OMP_INFOTYPE_EMPTY_MAPPING, Device.DeviceID,
+           "variable %s does not have a valid device counterpart\n",
+           (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 namespace {
 /// This structure contains information to deallocate a target pointer, aka.
 /// used to fix up the shadow map and potentially delete the entry from the

openmp/libomptarget/test/mapping/target_wrong_use_device_addr.c

+// RUN: %libomptarget-compile-generic -fopenmp-version=51 -g
+// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+#include <stdio.h>
+
+int main() {
+  float arr[10];
+  float *x = &arr[0];
+
+  // CHECK: host addr=0x[[#%x,HOST_ADDR:]]
+  fprintf(stderr, "host addr=%p\n", x);
+
+#pragma omp target data map(to : x [0:10])
+  {
+// CHECK: Libomptarget device 0 info: variable x does not have a valid device
+// counterpart
+#pragma omp target data use_device_addr(x)
+    {
+      // CHECK-NOT: device addr=0x[[#%x,HOST_ADDR:]]
+      fprintf(stderr, "device addr=%p\n", x);
+    }
+  }
+
+  return 0;
+}
+