diff mbox series

Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c')

Message ID 87k10msnk2.fsf@euler.schwinge.homeip.net
State New
Headers show
Series Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c') | expand

Commit Message

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

On 2019-12-09T12:52:02+0100, I wrote:
> See attached "[PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c'",
> committed to trunk in r279120, "to document the status quo", which does
> match my understanding of the OpenACC 2.6 semantics.

I've now pushed "Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more
[PR92854]" to master branch in commit
af8fd1a99d9a21f8088ebb11250cd06a3f275052, and releases/gcc-10 branch in
commit 364f46de9f02dc00e8ff51cc9e2662ae37520389, see attached.


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
diff mbox series

Patch

From 364f46de9f02dc00e8ff51cc9e2662ae37520389 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 29 May 2020 14:11:27 +0200
Subject: [PATCH] Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more
 [PR92854]

	libgomp/
	PR libgomp/92854
	* testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: Extend some
	more.

(cherry picked from commit af8fd1a99d9a21f8088ebb11250cd06a3f275052)
---
 .../libgomp.oacc-c-c++-common/pr92854-1.c     | 64 ++++++++++++++-----
 1 file changed, 47 insertions(+), 17 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
index 6ba96b6bf8f9..79cebf65c348 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
@@ -1,31 +1,61 @@ 
-/* Verify that 'acc_unmap_data' unmaps even in presence of dynamic reference
-   counts.  */
+/* Verify that 'acc_unmap_data' unmaps even in presence of structured and
+   dynamic reference counts, but the device memory remains allocated.  */
 
 /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
 
 #include <assert.h>
 #include <stdlib.h>
+#include <string.h>
 #include <openacc.h>
 
 int
 main ()
 {
   const int N = 180;
-
-  char *h = (char *) malloc (N);
-  char *d = (char *) acc_malloc (N);
-  if (!d)
-    abort ();
-  acc_map_data (h, d, N);
-
-  char *d_ = (char *) acc_create (h + 3, N - 77);
-  assert (d_ == d + 3);
-
-  d_ = (char *) acc_create (h, N);
-  assert (d_ == d);
-
-  acc_unmap_data (h);
-  assert (!acc_is_present (h, N));
+  const int N_i = 537;
+  const int C = 37;
+
+  unsigned char *h = (unsigned char *) malloc (N);
+  assert (h);
+  unsigned char *d = (unsigned char *) acc_malloc (N);
+  assert (d);
+
+  for (int i = 0; i < N_i; ++i)
+    {
+      acc_map_data (h, d, N);
+      assert (acc_is_present (h, N));
+#pragma acc parallel present(h[0:N])
+      {
+	if (i == 0)
+	  memset (h, C, N);
+      }
+
+      unsigned char *d_ = (unsigned char *) acc_create (h + 3, N - 77);
+      assert (d_ == d + 3);
+
+#pragma acc data create(h[6:N - 44])
+      {
+	d_ = (unsigned char *) acc_create (h, N);
+	assert (d_ == d);
+
+#pragma acc enter data create(h[0:N])
+
+	assert (acc_is_present (h, N));
+	acc_unmap_data (h);
+	assert (!acc_is_present (h, N));
+      }
+
+      /* We can however still access the device memory.  */
+#pragma acc parallel loop deviceptr(d)
+      for (int j = 0; j < N; ++j)
+	d[j] += i * j;
+    }
+
+  acc_memcpy_from_device(h, d, N);
+  for (int j = 0; j < N; ++j)
+    assert (h[j] == ((C + N_i * (N_i - 1) / 2 * j) % 256));
+
+  acc_free (d);
 
   return 0;
 }
-- 
2.26.2