diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h --- a/openmp/libomptarget/include/device.h +++ b/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; diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp --- a/openmp/libomptarget/src/device.cpp +++ b/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. diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp --- a/openmp/libomptarget/src/omptarget.cpp +++ b/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; 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 100755 --- /dev/null +++ b/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 + +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; +}