Message ID | 4bbb4970cd48424873573a50e627786cc9cf3378.1579292772.git.julian@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | Mixed static/dynamic data lifetimes with OpenACC (PR92843) | expand |
Hi! On 2020-01-17T13:18:20-0800, Julian Brown <julian@codesourcery.com> wrote: > This patch prevents "exit data" directives from copying back data that > was mapped with an acc_map_data API call. This matches the behaviour > expected by the pr92843-1.c test > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, > n->refcount--; > > if (copyfrom > + && n->refcount != REFCOUNT_INFINITY > && (kind != GOMP_MAP_FROM || n->refcount == 0)) > gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start, > (void *) (n->tgt->tgt_start + n->tgt_offset Such a change -- re-posted as <http://mid.mail-archive.com/7b2f54faa0a25a7a445ad86bf4726202a1190a4f.1590182783.git.julian@codesourcery.com>, and <http://mid.mail-archive.com/0585b4fda1ba5c76dce2ac5053a55e2ceef06041.1592343756.git.julian@codesourcery.com> -- does fix the 'libgomp.oacc-c-c++-common/pr92843-1.c' test case, but I don't agree that the change is correct. Instead, I have pushed "[OpenACC] Revert always-copyfrom behavior for 'GOMP_MAP_FORCE_FROM' in 'libgomp/oacc-mem.c:goacc_exit_data_internal'", <http://mid.mail-archive.com/87wo3ky5vn.fsf@euler.schwinge.homeip.net>. (This moves us past the point where it used to 'abort' before, but doesn't resolve the XFAIL, yet, which needs changes from "[OpenACC] Adjust dynamic reference count semantics", <http://mid.mail-archive.com/5e9472b80dc475214a4a082ef54ee919d7f9dcff.1592343756.git.julian@codesourcery.com>.) Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 45ab2b169d7..783e7f363fb 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1235,6 +1235,7 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, n->refcount--; if (copyfrom + && n->refcount != REFCOUNT_INFINITY && (kind != GOMP_MAP_FROM || n->refcount == 0)) gomp_copy_dev2host (acc_dev, aq, (void *) cur_node.host_start, (void *) (n->tgt->tgt_start + n->tgt_offset diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c index f16c46a37bf..786a12a8504 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c @@ -1,7 +1,6 @@ /* Verify that 'acc_copyout' etc. is a no-op if there's still a structured reference count. */ -/* { dg-xfail-run-if "TODO PR92843" { *-*-* } } */ /* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ #include <assert.h> @@ -96,6 +95,9 @@ test_acc_map_data () verify_array (h, N, c1); assign_array (h, N, c1); + /* Note that we're not expecting this (nor the copyouts below) to perform + an actual "finalize" or copyout since the data was mapped with + acc_map_data. */ #pragma acc exit data copyout (h[0:N]) finalize assert (acc_is_present (h, N)); verify_array (h, N, c1);