diff --git a/openmp/libomptarget/include/Debug.h b/openmp/libomptarget/include/Debug.h --- a/openmp/libomptarget/include/Debug.h +++ b/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, }; 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 @@ -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; 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 @@ -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; 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 --- /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 + +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; +} +