[OpenACC] Repair/restore 'is_tgt_unmapped' checking

	libgomp/
	* oacc-mem.c (goacc_exit_datum): Repair 'is_tgt_unmapped'
	checking.
	(acc_unmap_data, goacc_exit_data_internal): Restore
	'is_tgt_unmapped' checking.
	* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: New
	file.
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Adjust.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Likewise.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
(cherry picked from commit 06ec61726d192659cd446e59a91e78745037f0fd)

libgomp/oacc-mem.c

@@ -485,7 +485,8 @@ acc_unmap_data (void *h)
   tgt->tgt_end = 0;
   tgt->to_free = NULL;
 
-  gomp_remove_var (acc_dev, n);
+  bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+  assert (is_tgt_unmapped);
 
   gomp_mutex_unlock (&acc_dev->lock);
 
@@ -727,8 +728,16 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
 	gomp_remove_var_async (acc_dev, n, aq);
       else
 	{
+	  size_t num_mappings = 0;
+	  /* If the target_mem_desc represents a single data mapping, we can
+	     check that it is freed when this splay tree key's refcount reaches
+	     zero.  Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping with
+	     multiple members), fall back to skipping the test.  */
+	  for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
+	    if (n->tgt->list[l_i].key)
+	      ++num_mappings;
 	  bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
-	  assert (is_tgt_unmapped);
+	  assert (is_tgt_unmapped || num_mappings > 1);
 	}
     }
 
@@ -1145,7 +1154,28 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 				  cur_node.host_end - cur_node.host_start);
 
 	    if (n->refcount == 0)
-	      gomp_remove_var_async (acc_dev, n, aq);
+	      {
+		if (aq)
+		  /* TODO We can't do the 'is_tgt_unmapped' checking -- see the
+		     'gomp_unref_tgt' comment in
+		     <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
+		     PR92881.  */
+		  gomp_remove_var_async (acc_dev, n, aq);
+		else
+		  {
+		    size_t num_mappings = 0;
+		    /* If the target_mem_desc represents a single data mapping,
+		       we can check that it is freed when this splay tree key's
+		       refcount reaches zero.  Otherwise (e.g. for a
+		       'GOMP_MAP_STRUCT' mapping with multiple members), fall
+		       back to skipping the test.  */
+		    for (size_t l_i = 0; l_i < n->tgt->list_count; ++l_i)
+		      if (n->tgt->list[l_i].key)
+			++num_mappings;
+		    bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
+		    assert (is_tgt_unmapped || num_mappings > 1);
+		  }
+	      }
 	  }
 	  break;
 
@@ -1177,7 +1207,29 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 			     && str->refcount != REFCOUNT_INFINITY)
 		      str->refcount--;
 		    if (str->refcount == 0)
-		      gomp_remove_var_async (acc_dev, str, aq);
+		      {
+			if (aq)
+			  /* TODO We can't do the 'is_tgt_unmapped' checking --
+			     see the 'gomp_unref_tgt' comment in
+			     <http://mid.mail-archive.com/878snl36eu.fsf@euler.schwinge.homeip.net>;
+			     PR92881.  */
+			  gomp_remove_var_async (acc_dev, str, aq);
+			else
+			  {
+			    size_t num_mappings = 0;
+			    /* If the target_mem_desc represents a single data
+			       mapping, we can check that it is freed when this
+			       splay tree key's refcount reaches zero.
+			       Otherwise (e.g. for a 'GOMP_MAP_STRUCT' mapping
+			       with multiple members), fall back to skipping
+			       the test.  */
+			    for (size_t l_i = 0; l_i < str->tgt->list_count; ++l_i)
+			      if (str->tgt->list[l_i].key)
+				++num_mappings;
+			    bool is_tgt_unmapped = gomp_remove_var (acc_dev, str);
+			    assert (is_tgt_unmapped || num_mappings > 1);
+			  }
+		      }
 		  }
 	      }
 	    i += elems;

libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c

+/* Test dynamic unmapping of separate structure members.  */
+
+/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+  char a;
+  float b;
+};
+
+void test (bool use_directives)
+{
+  struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+  assert (acc_is_present (&s.a, sizeof s.a));
+  assert (acc_is_present (&s.b, sizeof s.b));
+
+  if (use_directives)
+    {
+#pragma acc exit data delete(s.a)
+    }
+  else
+    acc_delete (&s.a, sizeof s.a);
+  assert (!acc_is_present (&s.a, sizeof s.a));
+  assert (acc_is_present (&s.b, sizeof s.b));
+  if (use_directives)
+    {
+#pragma acc exit data delete(s.b)
+    }
+  else
+    acc_delete (&s.b, sizeof s.b);
+  assert (!acc_is_present (&s.a, sizeof s.a));
+  assert (!acc_is_present (&s.b, sizeof s.b));
+}
+
+int main ()
+{
+  test (true);
+  test (false);
+
+  return 0;
+}

libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90

@@ -40,19 +40,20 @@ program dtype
   if (.not. acc_is_present(var%a(5:n - 5))) stop 11
   if (.not. acc_is_present(var%b(5:n - 5))) stop 12
   if (.not. acc_is_present(var)) stop 13
-!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
-  if (acc_get_device_type() .ne. acc_device_host) then
-     if (acc_is_present(var%a(5:n - 5))) stop 21
-     if (acc_is_present(var%b(5:n - 5))) stop 22
-  end if
   print *, "CheCKpOInT1"
   ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-  if (.not. acc_is_present(var)) stop 23
-  !TODO { dg-output "STOP 23(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
   !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
   !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
   print *, "CheCKpOInT2"
   ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_get_device_type() .ne. acc_device_host) then
+     if (acc_is_present(var%a(5:n - 5))) stop 21
+     if (acc_is_present(var%b(5:n - 5))) stop 22
+  end if
+  if (.not. acc_is_present(var)) stop 23
 
 !$acc end data
 

libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90

@@ -21,16 +21,17 @@ program main
   if (.not. acc_is_present(var%a)) stop 1
   if (.not. acc_is_present(var)) stop 2
 
-  !$acc exit data delete(var%a) finalize
-  if (acc_is_present(var%a)) stop 3
   print *, "CheCKpOInT1"
   ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-  if (.not. acc_is_present(var)) stop 4
-  !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !$acc exit data delete(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
   !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
   !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
   print *, "CheCKpOInT2"
   ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
 
   !$acc end data
   if (acc_is_present(var%a)) stop 5

libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90

@@ -23,16 +23,17 @@ program main
   if (.not. acc_is_present(var%a)) stop 1
   if (.not. acc_is_present(var)) stop 2
 
-  !$acc exit data delete(var%a) finalize
-  if (acc_is_present(var%a)) stop 3
   print *, "CheCKpOInT1"
   ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-  if (.not. acc_is_present(var)) stop 4
-  !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !$acc exit data delete(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
   !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
   !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
   print *, "CheCKpOInT2"
   ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
 
   !$acc end data
   if (acc_is_present(var%a)) stop 5

libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90

@@ -23,16 +23,17 @@ program main
   if (.not. acc_is_present(var%a)) stop 1
   if (.not. acc_is_present(var)) stop 2
 
-  !$acc exit data delete(var%a)
-  if (acc_is_present(var%a)) stop 3
   print *, "CheCKpOInT1"
   ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-  if (.not. acc_is_present(var)) stop 4
-  !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !$acc exit data delete(var%a)
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
   !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
   !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
   print *, "CheCKpOInT2"
   ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
 
   !$acc end data
   if (acc_is_present(var%a)) stop 5

libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90

@@ -24,16 +24,17 @@ program main
   if (.not. acc_is_present(var)) stop 2
 
   !$acc exit data detach(var%a)
-  !$acc exit data delete(var%a) finalize
-  if (acc_is_present(var%a)) stop 3
   print *, "CheCKpOInT1"
   ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-  if (.not. acc_is_present(var)) stop 4
-  !TODO { dg-output "STOP 4(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !$acc exit data delete(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
   !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
   !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
   print *, "CheCKpOInT2"
   ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  if (acc_is_present(var%a)) stop 3
+  if (.not. acc_is_present(var)) stop 4
 
   !$acc end data
   if (acc_is_present(var%a)) stop 5

libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90

@@ -23,15 +23,16 @@ program main
   if (.not. acc_is_present(var%a)) stop 1
   if (.not. acc_is_present(var)) stop 2
 
-  !$acc exit data detach(var%a) finalize
   print *, "CheCKpOInT1"
   ! { dg-output ".*CheCKpOInT1(\n|\r\n|\r)" }
-  !$acc exit data delete(var%a)
-  !TODO { dg-output "(\n|\r\n|\r)libgomp: struct not mapped for detach operation(\n|\r\n|\r)$" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+  !$acc exit data detach(var%a) finalize
+  !TODO     goacc_exit_data_internal: Assertion `is_tgt_unmapped || num_mappings > 1' failed.
+  !TODO { dg-output ".*\[Aa\]ssert.*is_tgt_unmapped" { target { ! openacc_host_selected } } } ! Scan for what we expect in the "XFAILed" case (without actually XFAILing).
   !TODO { dg-shouldfail "XFAILed" { ! openacc_host_selected } } ! ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
   !TODO { dg-final { if { [dg-process-target { xfail { ! openacc_host_selected } }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } } ! ... so that we still get an XFAIL visible in the log.
   print *, "CheCKpOInT2"
   ! { dg-output ".CheCKpOInT2(\n|\r\n|\r)" { target { openacc_host_selected } } }
+  !$acc exit data delete(var%a)
   if (acc_is_present(var%a)) stop 3
   if (.not. acc_is_present(var)) stop 4