* [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c'
@ 2019-12-09 11:52 Thomas Schwinge
2019-12-11 17:04 ` [PR92854] Add 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c' Thomas Schwinge
` (2 more replies)
0 siblings, 3 replies; 4+ messages in thread
From: Thomas Schwinge @ 2019-12-09 11:52 UTC (permalink / raw)
To: gcc-patches; +Cc: Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 231 bytes --]
Hi!
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.
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-PR92854-Add-libgomp.oacc-c-c-common-pr92854-1..trunk.patch --]
[-- Type: text/x-diff, Size: 2089 bytes --]
From e14bd9d202bc4140d825a396ddaf64a5930ee3d1 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 9 Dec 2019 11:40:17 +0000
Subject: [PATCH] [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c'
... to document the status quo.
libgomp/
PR libgomp/92854
* testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279120 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 3 ++
.../libgomp.oacc-c-c++-common/pr92854-1.c | 31 +++++++++++++++++++
2 files changed, 34 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 6ef2f24e4d5..aac3b1887b0 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,5 +1,8 @@
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
+ PR libgomp/92854
+ * testsuite/libgomp.oacc-c-c++-common/pr92854-1.c: New file.
+
* testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New file.
* target.c (gomp_exit_data): Use 'gomp_remove_var'.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
new file mode 100644
index 00000000000..6ba96b6bf8f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92854-1.c
@@ -0,0 +1,31 @@
+/* Verify that 'acc_unmap_data' unmaps even in presence of dynamic reference
+ counts. */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdlib.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));
+
+ return 0;
+}
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 4+ messages in thread
* [PR92854] Add 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c'
2019-12-09 11:52 [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c' Thomas Schwinge
@ 2019-12-11 17:04 ` Thomas Schwinge
2020-06-04 18:17 ` Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c') Thomas Schwinge
2020-06-04 18:22 ` [OpenACC] Remove 'tgt' reference counting from 'acc_unmap_data' [PR92854] Thomas Schwinge
2 siblings, 0 replies; 4+ messages in thread
From: Thomas Schwinge @ 2019-12-11 17:04 UTC (permalink / raw)
To: gcc-patches; +Cc: Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 540 bytes --]
Hi!
See attached "[PR92854] Add
'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c',
'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c'", committed to
trunk in r279231, "to document the status quo", which does match my
understanding of the OpenACC 2.6 semantics.
The TODO in 'libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c'
is being tracked in PR92888 "[OpenACC] Failure to resolve back via
'acc_hostptr' an 'acc_deviceptr' retrieved for a '#pragma acc declare'd
variable".
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-PR92854-Add-libgomp.oacc-c-c-common-acc_map_da.trunk.patch --]
[-- Type: text/x-diff, Size: 10027 bytes --]
From ebcbd5ae0e1451cc97914cb825fc1017edb98e26 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 11 Dec 2019 16:48:59 +0000
Subject: [PATCH] [PR92854] Add
'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c',
'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c'
... to document the status quo.
libgomp/
PR libgomp/92854
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c:
New file.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c:
Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279231 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 16 +++++++++
.../acc_map_data-device_already-1.c | 36 +++++++++++++++++++
.../acc_map_data-device_already-2.c | 35 ++++++++++++++++++
.../acc_map_data-device_already-3.c | 31 ++++++++++++++++
.../acc_map_data-host_already-1.c | 33 +++++++++++++++++
.../acc_map_data-host_already-2.c | 32 +++++++++++++++++
.../acc_map_data-host_already-3.c | 27 ++++++++++++++
7 files changed, 210 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 6635ed7b44b..404722e20e3 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,19 @@
+2019-12-11 Thomas Schwinge <thomas@codesourcery.com>
+
+ PR libgomp/92854
+ * testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c:
+ New file.
+ * testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c:
+ Likewise.
+
2019-12-11 Thomas Schwinge <thomas@codesourcery.com>
Julian Brown <julian@codesourcery.com>
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c
new file mode 100644
index 00000000000..b48a1adbbb6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-1.c
@@ -0,0 +1,36 @@
+/* Verify that we refuse 'acc_map_data' when the "device address [...] is
+ already mapped". */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 131;
+
+ char *h1 = (char *) malloc (N);
+ assert (h1);
+ void *d = acc_malloc (N);
+ assert (d);
+ acc_map_data (h1, d, N);
+
+ char *h2 = (char *) malloc (N);
+ assert (h2);
+ /* Try to arrange a setting such that a later 'acc_unmap_data' would find the
+ device memory object still referenced elsewhere. This is not possible,
+ given the semantics of 'acc_map_data'. */
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_map_data (h2, d, N);
+
+ return 0;
+}
+
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+131\\\] is already mapped" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c
new file mode 100644
index 00000000000..4fe0662cabb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-2.c
@@ -0,0 +1,35 @@
+/* Verify that we refuse 'acc_map_data' when the "device address [...] is
+ already mapped". */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 132;
+
+ char *h1 = (char *) malloc (N);
+ assert (h1);
+ void *d = acc_create (h1, N);
+ assert (d);
+
+ char *h2 = (char *) malloc (N);
+ assert (h2);
+ /* Try to arrange a setting such that a later 'acc_unmap_data' would find the
+ device memory object still referenced elsewhere. This is not possible,
+ given the semantics of 'acc_map_data'. */
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_map_data (h2, d, N);
+
+ return 0;
+}
+
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+132\\\] is already mapped" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c
new file mode 100644
index 00000000000..44ebaa0eb3f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c
@@ -0,0 +1,31 @@
+/* Verify that we refuse 'acc_map_data' when the "device address [...] is
+ already mapped". */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdio.h>
+#include <openacc.h>
+
+double global_var;
+#pragma acc declare create (global_var)
+
+int
+main ()
+{
+ double var;
+ void *d = acc_deviceptr (&global_var);
+ assert (d);
+ /* Try to arrange a setting such that a later 'acc_unmap_data' would find the
+ device memory object still referenced elsewhere. This is not possible,
+ given the semantics of 'acc_map_data'. */
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_map_data (&var, d, sizeof var);
+
+ return 0;
+}
+
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+8\\\] is already mapped" { xfail *-*-* } } TODO PR92888 */
+/* { dg-shouldfail "TODO PR92888" { this-really-should-fail } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c
new file mode 100644
index 00000000000..1fff806613c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-1.c
@@ -0,0 +1,33 @@
+/* Verify that we refuse 'acc_map_data' when the "host address [...] is already
+ mapped". */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 101;
+
+ char *h = (char *) malloc (N);
+ assert (h);
+ void *d1 = acc_malloc (N);
+ assert (d1);
+ acc_map_data (h, d1, N);
+
+ void *d2 = acc_malloc (N);
+ assert (d2);
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_map_data (h, d2, N);
+
+ return 0;
+}
+
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "host address \\\[\[0-9a-fA-FxX\]+, \\\+101\\\] is already mapped" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c
new file mode 100644
index 00000000000..fc804692d1b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-2.c
@@ -0,0 +1,32 @@
+/* Verify that we refuse 'acc_map_data' when the "host address [...] is already
+ mapped". */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ const int N = 102;
+
+ char *h = (char *) malloc (N);
+ assert (h);
+ void *d1 = acc_create (h, N);
+ assert (d1);
+
+ void *d2 = acc_malloc (N);
+ assert (d2);
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_map_data (h, d2, N);
+
+ return 0;
+}
+
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "host address \\\[\[0-9a-fA-FxX\]+, \\\+102\\\] is already mapped" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c
new file mode 100644
index 00000000000..6a80ebfef46
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-host_already-3.c
@@ -0,0 +1,27 @@
+/* Verify that we refuse 'acc_map_data' when the "host address [...] is already
+ mapped". */
+
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <assert.h>
+#include <stdio.h>
+#include <openacc.h>
+
+float global_var;
+#pragma acc declare create (global_var)
+
+int
+main ()
+{
+ void *d = acc_malloc (sizeof global_var);
+ assert (d);
+ fprintf (stderr, "CheCKpOInT\n");
+ acc_map_data (&global_var, d, sizeof global_var);
+
+ return 0;
+}
+
+
+/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
+/* { dg-output "host address \\\[\[0-9a-fA-FxX\]+, \\\+4\\\] is already mapped" } */
+/* { dg-shouldfail "" } */
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 4+ messages in thread
* Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c')
2019-12-09 11:52 [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c' Thomas Schwinge
2019-12-11 17:04 ` [PR92854] Add 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c' Thomas Schwinge
@ 2020-06-04 18:17 ` Thomas Schwinge
2020-06-04 18:22 ` [OpenACC] Remove 'tgt' reference counting from 'acc_unmap_data' [PR92854] Thomas Schwinge
2 siblings, 0 replies; 4+ messages in thread
From: Thomas Schwinge @ 2020-06-04 18:17 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 719 bytes --]
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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-Extend-libgomp.oacc-c-c-common-pr92854-1.c-some-more.patch --]
[-- Type: text/x-diff, Size: 2636 bytes --]
From af8fd1a99d9a21f8088ebb11250cd06a3f275052 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.
---
.../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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-Extend-libgomp.oacc-c-c-common-pr92854-1.c-some-.g10.patch --]
[-- Type: text/x-diff, Size: 2706 bytes --]
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
^ permalink raw reply [flat|nested] 4+ messages in thread
* [OpenACC] Remove 'tgt' reference counting from 'acc_unmap_data' [PR92854]
2019-12-09 11:52 [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c' Thomas Schwinge
2019-12-11 17:04 ` [PR92854] Add 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c' Thomas Schwinge
2020-06-04 18:17 ` Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c') Thomas Schwinge
@ 2020-06-04 18:22 ` Thomas Schwinge
2 siblings, 0 replies; 4+ messages in thread
From: Thomas Schwinge @ 2020-06-04 18:22 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 524 bytes --]
Hi!
Per the discussion/conclusion in PR92854, I've now pushed "[OpenACC]
Remove 'tgt' reference counting from 'acc_unmap_data' [PR92854]" to
master branch in commit 4662f7fe7863b19fcc20ba58c22880f8d6661f3a, and
releases/gcc-10 branch in commit
0c59837c89bd62e2addf4b34704a1ebe7e3bffab, 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
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0001-OpenACC-Remove-tgt-reference-counting-from-acc_unmap.patch --]
[-- Type: text/x-diff, Size: 1070 bytes --]
From 4662f7fe7863b19fcc20ba58c22880f8d6661f3a Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 29 May 2020 14:12:16 +0200
Subject: [PATCH] [OpenACC] Remove 'tgt' reference counting from
'acc_unmap_data' [PR92854]
libgomp/
PR libgomp/92854
* oacc-mem.c (acc_unmap_data): Remove 'tgt' reference counting.
---
libgomp/oacc-mem.c | 14 +++++++-------
1 file changed, 7 insertions(+), 7 deletions(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index e2fb651a2334..6314f5d8b686 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -477,13 +477,13 @@ acc_unmap_data (void *h)
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("cannot unmap target block");
}
- else if (tgt->refcount > 1)
- tgt->refcount--;
- else
- {
- free (tgt->array);
- free (tgt);
- }
+
+ /* Above, we've verified that the mapping must have been set up by
+ 'acc_map_data'. */
+ assert (tgt->refcount == 1);
+
+ free (tgt->array);
+ free (tgt);
gomp_mutex_unlock (&acc_dev->lock);
--
2.26.2
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #3: 0001-OpenACC-Remove-tgt-reference-counting-from-acc_u.g10.patch --]
[-- Type: text/x-diff, Size: 1140 bytes --]
From 0c59837c89bd62e2addf4b34704a1ebe7e3bffab Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 29 May 2020 14:12:16 +0200
Subject: [PATCH] [OpenACC] Remove 'tgt' reference counting from
'acc_unmap_data' [PR92854]
libgomp/
PR libgomp/92854
* oacc-mem.c (acc_unmap_data): Remove 'tgt' reference counting.
(cherry picked from commit 4662f7fe7863b19fcc20ba58c22880f8d6661f3a)
---
libgomp/oacc-mem.c | 14 +++++++-------
1 file changed, 7 insertions(+), 7 deletions(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index e2fb651a2334..6314f5d8b686 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -477,13 +477,13 @@ acc_unmap_data (void *h)
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("cannot unmap target block");
}
- else if (tgt->refcount > 1)
- tgt->refcount--;
- else
- {
- free (tgt->array);
- free (tgt);
- }
+
+ /* Above, we've verified that the mapping must have been set up by
+ 'acc_map_data'. */
+ assert (tgt->refcount == 1);
+
+ free (tgt->array);
+ free (tgt);
gomp_mutex_unlock (&acc_dev->lock);
--
2.26.2
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2020-06-04 18:22 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2019-12-09 11:52 [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c' Thomas Schwinge
2019-12-11 17:04 ` [PR92854] Add 'libgomp.oacc-c-c++-common/acc_map_data-device_already-*.c', 'libgomp.oacc-c-c++-common/acc_map_data-host_already-*.c' Thomas Schwinge
2020-06-04 18:17 ` Extend 'libgomp.oacc-c-c++-common/pr92854-1.c' some more [PR92854] (was: [PR92854] Add 'libgomp.oacc-c-c++-common/pr92854-1.c') Thomas Schwinge
2020-06-04 18:22 ` [OpenACC] Remove 'tgt' reference counting from 'acc_unmap_data' [PR92854] Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).