[OpenMP] Fix ref count dec for implicit map of partial data

[jdenny]

Summary

D85342 broke this case. The new test case presents an example.

[jdenny]

Don't examine LIBOMPTARGET_DEBUG=1 output in the test. It doesn't work in release builds.

[grokos]

Nice catch! Looks good.

[jdenny]

Thanks!

openmp/libomptarget/src/omptarget.cpp

if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&
            "\n",
            Padding, DPxPTR(HstPtrBegin));
         HstPtrBegin = (char *)HstPtrBegin - Padding;
         DataSize += Padding;
       }
     }
 
     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;
     bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
     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"));
       if (HasPresentModifier) {
         // This should be an error upon entering an "omp target exit data".  It
         // should not be an error upon exiting an "omp target data" or "omp
         // target".  For "omp target data", Clang thus doesn't include present
         // modifiers for end calls.  For "omp target", we have not found a valid

openmp/libomptarget/test/mapping/target_implicit_partial_map.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu
+// RUN: env LIBOMPTARGET_DEBUG=1 \
+// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
+// RUN: | %fcheck-aarch64-unknown-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu
+// RUN: env LIBOMPTARGET_DEBUG=1 \
+// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu
+// RUN: env LIBOMPTARGET_DEBUG=1 \
+// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
+// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
+
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu
+// RUN: env LIBOMPTARGET_DEBUG=1 \
+// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
+// RUN: | %fcheck-x86_64-pc-linux-gnu
+//
+// END.
+
+#include <stdio.h>
+
+int main() {
+  int arr[100];
+
+  // CHECK: size=[[#%u,SIZE:]]
+  fprintf(stderr, "size=%ld\n", 2 * sizeof arr[0]);
+
+#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: Deleting tgt data 0x[[#%x,]] of size [[#SIZE]]
+
+  return 0;
+}