[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 updated this revision to Diff 463729.

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D134845/new/

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.463729.patch
Type: text/x-patch
Size: 2825 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20220929/cf4b62f5/attachment.bin>