Skip to content

Commit

Permalink
[libomptarget] Make omp_target_is_present checks storage instead of z…
Browse files Browse the repository at this point in the history
…ero length array.

Consider checking whether a pointer has been mapped can be achieved via omp_get_mapped_ptr.
omp_target_is_present is more needed to check whether the storage being pointed is mapped.
This restore the old behavior of omp_target_is_present before D123093
Fixes llvm/llvm-project#54899

Reviewed By: jdenny

Differential Revision: https://reviews.llvm.org/D123891
  • Loading branch information
ye-luo authored and memfrob committed Oct 4, 2022
1 parent 08812f3 commit 20f13eb
Show file tree
Hide file tree
Showing 2 changed files with 19 additions and 3 deletions.
6 changes: 5 additions & 1 deletion openmp/libomptarget/src/api.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -108,8 +108,12 @@ EXTERN int omp_target_is_present(const void *ptr, int device_num) {
DeviceTy &Device = *PM->Devices[device_num];
bool IsLast; // not used
bool IsHostPtr;
// omp_target_is_present tests whether a host pointer refers to storage that
// is mapped to a given device. However, due to the lack of the storage size,
// only check 1 byte. Cannot set size 0 which checks whether the pointer (zero
// lengh array) is mapped instead of the referred storage.
TargetPointerResultTy TPR =
Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
Device.getTgtPtrBegin(const_cast<void *>(ptr), 1, IsLast,
/*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr);
int rc = (TPR.TargetPointer != NULL);
Expand Down
16 changes: 14 additions & 2 deletions openmp/libomptarget/test/mapping/target_implicit_partial_map.c
Original file line number Diff line number Diff line change
Expand Up @@ -13,15 +13,27 @@ int main() {

#pragma omp target data map(alloc: arr[50:2]) // partially mapped
{
// CHECK: arr[50] must present: 1
fprintf(stderr, "arr[50] must present: %d\n",
omp_target_is_present(&arr[50], omp_get_default_device()));

// CHECK: arr[0] should not present: 0
fprintf(stderr, "arr[0] should not present: %d\n",
omp_target_is_present(&arr[0], omp_get_default_device()));

// CHECK: arr[49] should not present: 0
fprintf(stderr, "arr[49] should not present: %d\n",
omp_target_is_present(&arr[49], omp_get_default_device()));

#pragma omp target // would implicitly map with full size but already present
{
arr[50] = 5;
arr[51] = 6;
} // must treat as present (dec ref count) even though full size not present
} // wouldn't delete if previous ref count dec didn't happen

// CHECK: still present: 0
fprintf(stderr, "still present: %d\n",
// CHECK: arr[50] still present: 0
fprintf(stderr, "arr[50] still present: %d\n",
omp_target_is_present(&arr[50], omp_get_default_device()));

return 0;
Expand Down

0 comments on commit 20f13eb

Please sign in to comment.