[OpenACC] Repair/restore 'is_tgt_unmapped' checking (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts)

Message ID 87bllysmqi.fsf@euler.schwinge.homeip.net
State New
Headers show
Series
  • [OpenACC] Repair/restore 'is_tgt_unmapped' checking (was: [PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts)
Related show

Commit Message

Thomas Schwinge June 4, 2020, 6:35 p.m.
Hi!

On 2020-05-20T20:11:00+0100, Julian Brown <julian@codesourcery.com> wrote:
> On Wed, 20 May 2020 16:52:02 +0200

> Thomas Schwinge <thomas@codesourcery.com> wrote:

>> On 2019-12-17T22:03:47-0800, Julian Brown <julian@codesourcery.com>

>> wrote:

>> > --- a/libgomp/oacc-mem.c

>> > +++ b/libgomp/oacc-mem.c

>>

>> >  static int

>> > -find_group_last (int pos, size_t mapnum, unsigned short *kinds)

>> > +find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds) {

>> >    unsigned char kind0 = kinds[pos] & 0xff;

>> > -  int first_pos = pos, last_pos = pos;

>> > +  int first_pos = pos;

>> >

>> > -  if (kind0 == GOMP_MAP_TO_PSET)

>> > +  switch (kind0)

>> >      {

>> > +    case GOMP_MAP_TO_PSET:

>> >        while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)

>> > -  last_pos = ++pos;

>> > +  pos++;

>> >        /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET.  */

>> > -      assert (last_pos > first_pos);

>> > -    }

>> > -  else

>> > -    {

>> > +      assert (pos > first_pos);

>> > +      break;

>> > +

>> > +    case GOMP_MAP_STRUCT:

>> > +      pos += sizes[pos];

>> > +      break;

>> > +

>> > +    case GOMP_MAP_POINTER:

>> > +    case GOMP_MAP_ALWAYS_POINTER:

>> > +      /* These mappings are only expected after some other mapping.  If we

>> > +   see one by itself, something has gone wrong.  */

>> > +      gomp_fatal ("unexpected mapping");

>> > +      break;

>> > +

>> > +    default:

>> >        /* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other mapping.  */

>> > -      if (pos + 1 < mapnum

>> > -    && (kinds[pos + 1] & 0xff) == GOMP_MAP_ALWAYS_POINTER)

>> > -  return pos + 1;

>> > +      if (pos + 1 < mapnum)

>> > +  {

>> > +    unsigned char kind1 = kinds[pos + 1] & 0xff;

>> > +    if (kind1 == GOMP_MAP_ALWAYS_POINTER)

>> > +      return pos + 1;

>> > +  }

>> >

>> > -      /* We can have one or several GOMP_MAP_POINTER mappings after a to/from

>> > +      /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from (etc.) mapping.  */

>> >        while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)

>> > -  last_pos = ++pos;

>> > +  pos++;

>> >      }

>> >

>> > -  return last_pos;

>> > +  return pos;

>> >  }

>>

>> So this now causes grouped (!) mapping of all of 'GOMP_MAP_STRUCT',

>> that is, all its "members" at once.

>>

>> This, I suppose, mandated the removal of (some of) the

>> 'is_tgt_unmapped' checking (unfortunately committed not here, but as

>> part of r279621 "OpenACC reference count overhaul"), where we had

>> unmapping code (conceptually) similar to:

>>

>>     bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);

>>     assert (is_tgt_unmapped);

>>

>> I'd introduced this a little bit earlier, finding this a simple yet

>> effective run-time, low-overhead consistency checking of (certain

>> aspects of) reference counting -- so just noting here that it's

>> somewhat bad that we can't have this anymore "just" because of

>> 'GOMP_MAP_STRUCT'.  (Maybe there is a way to get it back; that's for

>> later?)

>

> I'm actually looking at this now as part of revisiting the refcounting

> work. I'm seeing what I can come up with in terms of being able to keep

> the runtime test (and fixing the other part you mentioned).


Good idea to skip the checking if 'num_mappings > 1', thanks!  You've
included these changes as part of a different, bigger patch; I've split
them out, and pushed "[OpenACC] Repair/restore 'is_tgt_unmapped'
checking" to master branch in commit
06ec61726d192659cd446e59a91e78745037f0fd, and releases/gcc-10 branch in
commit 125621f569cfac9f4caa6afc1976d42b3d21359e, see attached.

Also note how this checking triggers for OpenACC 'finalize' clause usage
in 'libgomp.oacc-fortran/deep-copy-6.f90' (which I had relatedly XFAILed
before); so good to see that it's useful for something!  (..., and did
your reference count self-checking not catch this case?)


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter

Patch

From 125621f569cfac9f4caa6afc1976d42b3d21359e Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 29 May 2020 15:22:42 +0200
Subject: [PATCH] [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                            | 60 +++++++++++++++++--
 .../struct-refcount-1.c                       | 47 +++++++++++++++
 .../libgomp.oacc-fortran/deep-copy-6.f90      | 15 ++---
 .../mdc-refcount-1-1-1.f90                    |  9 +--
 .../mdc-refcount-1-2-1.f90                    |  9 +--
 .../mdc-refcount-1-2-2.f90                    |  9 +--
 .../mdc-refcount-1-3-1.f90                    |  9 +--
 .../mdc-refcount-1-4-1.f90                    |  7 ++-
 8 files changed, 135 insertions(+), 30 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 8e8c7c3093d5..b7c85cf5976f 100644
--- a/libgomp/oacc-mem.c
+++ b/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;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
new file mode 100644
index 000000000000..bde5890d6676
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
@@ -0,0 +1,47 @@ 
+/* 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;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
index a7943d93d542..5837a4039106 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
+++ b/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
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
index 449a6cf6894e..445cbabb8ca4 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90
+++ b/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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
index a7e649d30415..8554534b2f27 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90
+++ b/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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
index 3402fafc7e28..8e696cc70e80 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90
+++ b/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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
index 7504969d9a56..070a6f8e1493 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90
+++ b/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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
index fedae0db054b..b22e411567ff 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90
+++ b/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
 
-- 
2.26.2