[gcc(refs/users/giulianob/heads/autopar_rebase2)] [OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries
Giuliano Belinassi
giulianob@gcc.gnu.org
Mon Aug 17 23:12:19 GMT 2020
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;
-}
More information about the Gcc-cvs
mailing list