Emit info message when use_device_address variable does not have a device counterpart.

4 files changed, 43 insertions, 1 deletions

diff --git a/openmp/libomptarget/include/Debug.h b/openmp/libomptarget/include/Debug.h
index 91c12688789a..17aa8746afe7 100644
--- a/openmp/libomptarget/include/Debug.h
+++ b/openmp/libomptarget/include/Debug.h
@@ -55,6 +55,8 @@ enum OpenMPInfoType : uint32_t {
   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,
 };
diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h
index 5f147f9c77e6..7bc4629c6220 100644
--- a/openmp/libomptarget/include/device.h
+++ b/openmp/libomptarget/include/device.h
@@ -312,7 +312,9 @@ struct TargetPointerResultTy {
     /// 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 @@ struct TargetPointerResultTy {
 
   bool isHostPointer() const { return Flags.IsHostPointer; }
 
+  bool isContained() const { return Flags.IsContained; }
+
   /// The corresponding target pointer
   void *TargetPointer = nullptr;
 
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 0459648a4802..3094e899d0de 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -732,6 +732,15 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
           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;
diff --git a/openmp/libomptarget/test/mapping/target_wrong_use_device_addr.c b/openmp/libomptarget/test/mapping/target_wrong_use_device_addr.c
new file mode 100644
index 000000000000..bb5a1d44031a
--- /dev/null
+++ b/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;
+}
+