From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 2100) id 2A60B3949D80; Sat, 22 Aug 2020 22:02:32 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 2A60B3949D80 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1598133752; bh=eCEmpa+dhqV7Q2E8UBoGzOjkvama1kBHL/1wnd46Gjk=; h=From:To:Subject:Date:From; b=vZ2ML1Rr1u0D3AOIRF44MsIYMkuGTenXNlJXDQhDqzhSMANoGYtijQMO5+P0/WtIT PtC59c6myRllipJEjrnMdr/EDvdmAZV/7RvxYGH068qhEpztOBCJS/MGFMQBWw4cT4 mFxAskyFH1MOwcRRLIK1neiPeNfXFGfC2mfi6xOk= Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit From: Giuliano Belinassi To: gcc-cvs@gcc.gnu.org Subject: [gcc/devel/autopar_devel] [OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/devel/autopar_devel X-Git-Oldrev: 6364a238ebb02c2f820a5d4611231f81ce7955fc X-Git-Newrev: f2572f30a3deef67cd58b4961e06c0de637dfced Message-Id: <20200822220232.2A60B3949D80@sourceware.org> Date: Sat, 22 Aug 2020 22:02:32 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Sat, 22 Aug 2020 22:02:32 -0000 https://gcc.gnu.org/g:f2572f30a3deef67cd58b4961e06c0de637dfced commit f2572f30a3deef67cd58b4961e06c0de637dfced Author: Thomas Schwinge 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) : 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 +#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; -}