[Openmp-commits] [PATCH] D134845: [OpenMP][libomptarget] Emit runtime message when variable is incorrectly mapped to device

[Gheorghe-Teodor Bercea via Phabricator via Openmp-commits]

doru1004 created this revision.
doru1004 added reviewers: mjklemm, trws, dreachem, carlo.bertolli, ronlieb.
doru1004 added a project: OpenMP.
Herald added subscribers: guansong, yaxunl.
Herald added a project: All.
doru1004 requested review of this revision.
Herald added a reviewer: jdoerfert.
Herald added subscribers: openmp-commits, sstefan1.

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.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D134845

Files:
  openmp/libomptarget/include/device.h
  openmp/libomptarget/src/device.cpp
  openmp/libomptarget/src/omptarget.cpp
  openmp/libomptarget/test/mapping/target_wrong_use_device_addr.c


Index: openmp/libomptarget/test/mapping/target_wrong_use_device_addr.c
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/mapping/target_wrong_use_device_addr.c
@@ -0,0 +1,26 @@
+// 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;
+}
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -627,6 +627,14 @@
       } 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;
Index: openmp/libomptarget/src/device.cpp
===================================================================
--- openmp/libomptarget/src/device.cpp
+++ openmp/libomptarget/src/device.cpp
@@ -356,7 +356,9 @@
     }
   }
 
-  return {{IsNew, IsHostPtr, IsPresent}, Entry, TargetPointer};
+  return {{IsNew, IsHostPtr, IsPresent, LR.Flags.IsContained},
+          Entry,
+          TargetPointer};
 }
 
 // Used by targetDataBegin, targetDataEnd, targetDataUpdate and target.
Index: openmp/libomptarget/include/device.h
===================================================================
--- openmp/libomptarget/include/device.h
+++ openmp/libomptarget/include/device.h
@@ -283,12 +283,16 @@
     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;
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D134845.463725.patch
Type: text/x-patch
Size: 2825 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20220929/acd46833/attachment.bin>