From a02f1adbfe619ab19cf142438e0a02950d3594da Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Thu, 4 Jun 2020 16:01:07 +0200 Subject: [PATCH] [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) : 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. --- libgomp/oacc-mem.c | 10 ++ .../libgomp.oacc-c-c++-common/struct-1.c | 146 ++++++++++++++++++ .../struct-refcount-1.c | 47 ------ 3 files changed, 156 insertions(+), 47 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c delete 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 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 +#include +#include + +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 -#include -#include - -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; -} -- 2.30.2