[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 517183.

Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D134845

Files:
  openmp/libomptarget/include/Debug.h
  openmp/libomptarget/include/device.h
  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,27 @@
+// 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;
+}
+
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -732,6 +732,15 @@
           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;
Index: openmp/libomptarget/include/device.h
===================================================================
--- openmp/libomptarget/include/device.h
+++ openmp/libomptarget/include/device.h
@@ -312,7 +312,9 @@
     /// 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;
@@ -348,6 +350,8 @@
 
   bool isHostPointer() const { return Flags.IsHostPointer; }
 
+  bool isContained() const { return Flags.IsContained; }
+
   /// The corresponding target pointer
   void *TargetPointer = nullptr;
 
Index: openmp/libomptarget/include/Debug.h
===================================================================
--- openmp/libomptarget/include/Debug.h
+++ openmp/libomptarget/include/Debug.h
@@ -55,6 +55,8 @@
   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,
 };


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D134845.517183.patch
Type: text/x-patch
Size: 3072 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20230426/185e0c8a/attachment.bin>