Index: openmp/libomptarget/include/Debug.h =================================================================== --- openmp/libomptarget/include/Debug.h +++ openmp/libomptarget/include/Debug.h @@ -54,6 +54,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 cotunerpart + OMP_INFOTYPE_EMPTY_MAPPING = 0x0040, // Enable every flag. OMP_INFOTYPE_ALL = 0xffffffff, }; 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; 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/src/omptarget.cpp =================================================================== --- openmp/libomptarget/src/omptarget.cpp +++ openmp/libomptarget/src/omptarget.cpp @@ -627,6 +627,15 @@ } else Device.ShadowMtx.unlock(); } + + // 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/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: env LIBOMPTARGET_INFO=64 %libomptarget-run-fail-generic 2>&1 \ +// RUN: | %fcheck-generic + +#include + +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; +}