[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

57 ↗

(On Diff #466169)

cotunerpart -> counterpart

[ronlieb]

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

[RaviNarayanaswamy]

LGTM

[doru1004]

Commit: a91cb9ce39dc42e6a7a2c4fe97580e51eb1c2961

openmp/libomptarget/include/device.h

 struct TargetPointerResultTy {
   struct {
     /// If the map table entry is just created
     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;
-  } Flags = {0, 0, 0};
+    /// If the pointer is contained.
+    unsigned IsContained : 1;
+  } Flags = {0, 0, 0, 0};
 
   bool isPresent() const { return Flags.IsPresent; }
 
   bool isHostPointer() const { return Flags.IsHostPointer; }
 
+  bool isContained() const { return Flags.IsContained; }
+
   /// The corresponding map table entry which is stable.
   HostDataToTargetTy *Entry = nullptr;
 
   /// The corresponding target pointer
   void *TargetPointer = nullptr;
 };
 
 /// Map for shadow pointers

openmp/libomptarget/src/device.cpp

if (Entry && !IsHostPtr && !HasPresentModifier) {
           return {{false /* IsNewEntry */, false /* IsHostPointer */},
                   nullptr /* Entry */,
                   nullptr /* TargetPointer */};
         }
       }
     }
   }
 
-  return {{IsNew, IsHostPtr, IsPresent}, Entry, TargetPointer};
+  return {{IsNew, IsHostPtr, IsPresent, LR.Flags.IsContained},
+          Entry,
+          TargetPointer};
 }
 
 // Used by targetDataBegin, targetDataEnd, targetDataUpdate and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from targetDataEnd.
 TargetPointerResultTy
 DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
                          bool UpdateRefCount, bool UseHoldRefCount,

openmp/libomptarget/src/omptarget.cpp

if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
           return OFFLOAD_FAIL;
         }
         if (PointerTpr.Entry->addEventIfNecessary(Device, AsyncInfo) !=
             OFFLOAD_SUCCESS)
           return OFFLOAD_FAIL;
       } else
         Device.ShadowMtx.unlock();
     }
+
+    // Check if variable can be used on the device:
+    bool IsStructMember = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;
+    if (ArgTypes[I] != 0 && !IsStructMember && !IsImplicit &&
+        !TPR.isPresent() && !TPR.isContained() && !TPR.isHostPointer())
+      MESSAGE("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: %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 message: 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;
+}