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 @@ -497,6 +497,7 @@ } bool IsLast, IsHostPtr; + bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT; bool UpdateRef = !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) || (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ); bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE; @@ -504,9 +505,8 @@ bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT; // If PTR_AND_OBJ, HstPtrBegin is address of pointee - void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast, - UpdateRef, IsHostPtr, - /*MustContain=*/true); + void *TgtPtrBegin = Device.getTgtPtrBegin( + HstPtrBegin, DataSize, IsLast, UpdateRef, IsHostPtr, !IsImplicit); if (!TgtPtrBegin && (DataSize || HasPresentModifier)) { DP("Mapping does not exist (%s)\n", (HasPresentModifier ? "'present' map type modifier" : "ignored")); diff --git a/openmp/libomptarget/test/mapping/target_implicit_partial_map.c b/openmp/libomptarget/test/mapping/target_implicit_partial_map.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/mapping/target_implicit_partial_map.c @@ -0,0 +1,39 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu +// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \ +// RUN: | %fcheck-aarch64-unknown-linux-gnu + +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu +// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64-ibm-linux-gnu + +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu +// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \ +// RUN: | %fcheck-powerpc64le-ibm-linux-gnu + +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu +// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu +// +// END. + +#include +#include + +int main() { + int arr[100]; + +#pragma omp target data map(alloc: arr[50:2]) // partially mapped + { +#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", + omp_target_is_present(&arr[50], omp_get_default_device())); + + return 0; +}