public inbox for gcc-cvs@sourceware.org help / color / mirror / Atom feed
From: Giuliano Belinassi <giulianob@gcc.gnu.org> To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/autopar_devel] [OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries Date: Sat, 22 Aug 2020 22:02:32 +0000 (GMT) [thread overview] Message-ID: <20200822220232.2A60B3949D80@sourceware.org> (raw) 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; -}
reply other threads:[~2020-08-22 22:02 UTC|newest] Thread overview: [no followups] expand[flat|nested] mbox.gz Atom feed
Reply instructions: You may reply publicly to this message via plain-text email using any one of the following methods: * Save the following mbox file, import it into your mail client, and reply-to-all from there: mbox Avoid top-posting and favor interleaved quoting: https://en.wikipedia.org/wiki/Posting_style#Interleaved_style * Reply using the --to, --cc, and --in-reply-to switches of git-send-email(1): git send-email \ --in-reply-to=20200822220232.2A60B3949D80@sourceware.org \ --to=giulianob@gcc.gnu.org \ --cc=gcc-cvs@gcc.gnu.org \ /path/to/YOUR_REPLY https://kernel.org/pub/software/scm/git/docs/git-send-email.html * If your mail client supports setting the In-Reply-To header via mailto: links, try the mailto: linkBe sure your reply has a Subject: header at the top and a blank line before the message body.
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).