public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/autopar_devel] [OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries
@ 2020-08-22 22:02 Giuliano Belinassi
0 siblings, 0 replies; only message in thread
From: Giuliano Belinassi @ 2020-08-22 22:02 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:f2572f30a3deef67cd58b4961e06c0de637dfced
commit f2572f30a3deef67cd58b4961e06c0de637dfced
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu Jun 4 16:01:07 2020 +0200
[OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries
Currently, we don't at all evaluate 'finalize' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.
libgomp/
* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
Evaluate 'finalize' individually for each entry.
* testsuite/libgomp.oacc-c-c++-common/struct-1.c: New file.
* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: Remove
file.
Diff:
---
libgomp/oacc-mem.c | 10 ++
.../testsuite/libgomp.oacc-c-c++-common/struct-1.c | 146 +++++++++++++++++++++
.../libgomp.oacc-c-c++-common/struct-refcount-1.c | 47 -------
3 files changed, 156 insertions(+), 47 deletions(-)
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index b7c85cf5976..a34f4cf0e91 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1184,6 +1184,16 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
int elems = sizes[i];
for (int j = 1; j <= elems; j++)
{
+ assert (i + j < mapnum);
+
+ kind = kinds[i + j] & 0xff;
+
+ finalize = false;
+ if (kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FORCE_DETACH)
+ finalize = true;
+
struct splay_tree_key_s k;
k.host_start = (uintptr_t) hostaddrs[i + j];
k.host_end = k.host_start + sizes[i + j];
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
new file mode 100644
index 00000000000..285be84f244
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
@@ -0,0 +1,146 @@
+/* Test dynamic refcount of separate structure members. */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+ signed char a;
+ float b;
+};
+
+static void test(unsigned variant)
+{
+ 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 (variant & 4)
+ {
+ if (variant & 8)
+ {
+#pragma acc enter data create(s.b)
+ }
+ else
+ acc_create(&s.b, sizeof s.b);
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+
+ if (variant & 16)
+ {
+#pragma acc enter data create(s.a)
+ }
+ else
+ acc_create(&s.a, sizeof s.a);
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+
+ if (variant & 32)
+ {
+#pragma acc enter data create(s.a)
+ acc_create(&s.b, sizeof s.b);
+#pragma acc enter data create(s.b)
+#pragma acc enter data create(s.b)
+ acc_create(&s.a, sizeof s.a);
+ acc_create(&s.a, sizeof s.a);
+ acc_create(&s.a, sizeof s.a);
+ }
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+ }
+
+#pragma acc parallel \
+ copy(s.a, s.b)
+ {
+ }
+
+ if (variant & 32)
+ {
+ if (variant & 1)
+ {
+#pragma acc exit data delete(s.a) finalize
+ }
+ else
+ acc_delete_finalize(&s.a, sizeof s.a);
+ }
+ else
+ {
+ if (variant & 1)
+ {
+#pragma acc exit data delete(s.a)
+ }
+ else
+ acc_delete(&s.a, sizeof s.a);
+ if (variant & 4)
+ {
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+ if (variant & 1)
+ {
+#pragma acc exit data delete(s.a)
+ }
+ else
+ acc_delete(&s.a, sizeof s.a);
+ }
+ }
+#if ACC_MEM_SHARED
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+#else
+ assert(!acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+#endif
+
+ if (variant & 32)
+ {
+ if (variant & 2)
+ {
+#pragma acc exit data delete(s.b) finalize
+ }
+ else
+ acc_delete_finalize(&s.b, sizeof s.b);
+ }
+ else
+ {
+ if (variant & 2)
+ {
+#pragma acc exit data delete(s.b)
+ }
+ else
+ acc_delete(&s.b, sizeof s.b);
+ if (variant & 4)
+ {
+#if ACC_MEM_SHARED
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+#else
+ assert(!acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+#endif
+ if (variant & 2)
+ {
+#pragma acc exit data delete(s.b)
+ }
+ else
+ acc_delete(&s.b, sizeof s.b);
+ }
+ }
+#if ACC_MEM_SHARED
+ assert(acc_is_present(&s.a, sizeof s.a));
+ assert(acc_is_present(&s.b, sizeof s.b));
+#else
+ assert(!acc_is_present(&s.a, sizeof s.a));
+ assert(!acc_is_present(&s.b, sizeof s.b));
+#endif
+}
+
+int main()
+{
+ for (unsigned variant = 0; variant < 64; ++variant)
+ test(variant);
+
+ return 0;
+}
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
deleted file mode 100644
index bde5890d667..00000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
+++ /dev/null
@@ -1,47 +0,0 @@
-/* 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;
-}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2020-08-22 22:02 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-08-22 22:02 [gcc/devel/autopar_devel] [OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries Giuliano Belinassi
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).