openacc: Remove unnecessary detach finalization

The call to gomp_detach_pointer in gomp_unmap_vars_internal does not
need to force finalization, and doing so may mask mismatched pointer
attachments/detachments. This patch removes the forcing.

2020-07-16  Julian Brown  <julian@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

libgomp/
	* target.c (gomp_unmap_vars_internal): Remove unnecessary forcing of
	finalization for detach operation.
	* testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c:
	New test.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>

diff --git a/libgomp/target.c b/libgomp/target.c
index d6b3572..00c75fb 100644 (file)
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1437,7 +1437,7 @@ gomp_unmap_vars_internal (struct target_mem_desc *tgt, bool do_copyfrom,
       if (k != NULL && tgt->list[i].do_detach)
        gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
                                             + tgt->list[i].offset,
-                            k->refcount == 1, NULL);
+                            false, NULL);
     }
 
   for (i = 0; i < tgt->list_count; i++)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
new file mode 100644 (file)
index 0000000..fc1f59e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/structured-detach-underflow.c
@@ -0,0 +1,28 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <stdio.h>
+
+int main ()
+{
+  struct {
+    int *arr;
+  } mystr;
+  int localarr[16];
+  mystr.arr = localarr;
+
+  #pragma acc enter data copyin(mystr, localarr[0:16])
+
+  #pragma acc data attach(mystr.arr)
+  {
+    #pragma acc exit data detach(mystr.arr)
+    fprintf (stderr, "CheCKpOInT1\n");
+    /* { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" } */
+  }
+  /* { dg-shouldfail "" }
+     { dg-output "(\n|\r\n|\r)libgomp: attach count underflow(\n|\r\n|\r)$" } */
+  fprintf (stderr, "CheCKpOInT2\n");
+
+  #pragma acc exit data copyout(mystr, localarr[0:16])
+
+  return 0;
+}