Message ID | 87pp0aaksc.fsf@kepler.schwinge.homeip.net |
---|---|
State | New |
Headers | show |
On Mon, Oct 19, 2015 at 18:24:35 +0200, Thomas Schwinge wrote: > Chung-Lin, would you please have a look at the following (on > gomp-4_0-branch)? Also, anyone else got any ideas off-hand? > > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors) > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors) > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test Maybe it was caused by this change in gomp_unmap_vars? https://gcc.gnu.org/ml/gcc-patches/2015-06/msg01376.html Looking at the code, I don't see any difference in async_refcount handling, but I was unable to test it without having hardware :( -- Ilya
On Mon, Oct 19, 2015 at 07:43:59PM +0300, Ilya Verbin wrote: > On Mon, Oct 19, 2015 at 18:24:35 +0200, Thomas Schwinge wrote: > > Chung-Lin, would you please have a look at the following (on > > gomp-4_0-branch)? Also, anyone else got any ideas off-hand? > > > > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors) > > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test > > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 (test for excess errors) > > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/data-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 execution test > > Maybe it was caused by this change in gomp_unmap_vars? > https://gcc.gnu.org/ml/gcc-patches/2015-06/msg01376.html > > Looking at the code, I don't see any difference in async_refcount handling, but > I was unable to test it without having hardware :( I think that is the only patch that could have affected it. The copy_from change is from the old behavior, where basically all concurrent mappings ored into the copy_from flag and when refcount went to 0, if there were any mappings with from or tofrom, it copied back, the OpenMP 4.5 behavior is that whether data is copied from the device is determined solely by the mapping kind of the mapping that performs the refcount decrease to 0. Plus there is the always flag which requests the data copying operation always, no matter what the refcount is (either on the mapping/refcount increase side, or unmapping/refcount decrease size). Jakub
Hi! On Mon, 19 Oct 2015 18:24:35 +0200, I wrote: > Chung-Lin, would you please have a look at the following (on > gomp-4_0-branch)? Also, anyone else got any ideas off-hand? Ilya, Jakub, thanks for your comments! > On Tue, 23 Jun 2015 13:51:39 +0200, Jakub Jelinek <jakub@redhat.com> wrote: > > On Tue, Jun 23, 2015 at 02:40:43PM +0300, Ilya Verbin wrote: > > > On Sat, Jun 20, 2015 at 00:35:14 +0300, Ilya Verbin wrote: > > > > Given that a mapped variable in 4.1 can have different kinds across nested data > > > > regions, we need to store map-type not only for each var, but also for each > > > > structured mapping. Here is my WIP patch, is it sane? :) > > > > Attached testcase works OK on the device with non-shared memory. > > > > > > A bit updated version with a fix for GOMP_MAP_TO_PSET. > > > make check-target-libgomp passed. > > > > Ok, thanks. > > > > > include/gcc/ > > > * gomp-constants.h (GOMP_MAP_ALWAYS_TO_P, > > > GOMP_MAP_ALWAYS_FROM_P): Define. > > > libgomp/ > > > * libgomp.h (struct target_var_desc): New. > > > (struct target_mem_desc): Replace array of splay_tree_key with array of > > > target_var_desc. > > > (struct splay_tree_key_s): Move copy_from to target_var_desc. > > > * oacc-mem.c (gomp_acc_remove_pointer): Use copy_from from > > > target_var_desc. > > > * oacc-parallel.c (GOACC_parallel): Use copy_from from target_var_desc. > > > * target.c (gomp_map_vars_existing): Copy data to device if map-type is > > > 'always to' or 'always tofrom'. > > > (gomp_map_vars): Use key from target_var_desc. Set copy_from and > > > always_copy_from. > > > (gomp_copy_from_async): Use key and copy_from from target_var_desc. > > > (gomp_unmap_vars): Copy data from device if always_copy_from is set. > > > (gomp_offload_image_to_device): Do not use copy_from. > > > * testsuite/libgomp.c/target-11.c: New test. > > (That's gomp-4_1-branch r224838. The attached > gomp-4_1-branch-r224838.patch is a variant that applies on top of > gomp-4_0-branch r228972.) This change introduces regressions in OpenACC > async clause handling. > Testing on gomp-4_0-branch r228972 plus the attached > gomp-4_1-branch-r224838.patch: > > PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/asyncwait-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none (test for excess errors) > [-PASS:-]{+FAIL:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/asyncwait-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none execution test > > Same for C++. With an XFAIL added (Chung-Lin, please remove that one once you come up with a fix), and merge conflicts resolved as follows, I have now merged gomp-4_1-branch r224838 in gomp-4_0-branch r229178: commit cbef8ef8e3b6bf7ea3705b1fae5462be9e619a56 Merge: 3596aeb a568354 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Thu Oct 22 17:50:08 2015 +0000 svn merge -r 224607:224838 svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_1-branch git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@229178 138bc75d-0d04-0410-961f-82ee72b054a4 include/ChangeLog.gomp41 | 5 + include/gomp-constants.h | 6 ++ libgomp/ChangeLog.gomp41 | 18 ++++ libgomp/libgomp.h | 15 ++- libgomp/oacc-mem.c | 2 +- libgomp/oacc-parallel.c | 6 +- libgomp/target.c | 106 +++++++++++++-------- libgomp/testsuite/libgomp.c/target-11.c | 51 ++++++++++ .../libgomp.oacc-c-c++-common/asyncwait-1.c | 2 + 9 files changed, 162 insertions(+), 49 deletions(-) diff --cc libgomp/oacc-mem.c index 7fcf199,c0fcb07..a90c912 --- libgomp/oacc-mem.c +++ libgomp/oacc-mem.c @@@ -685,7 -650,8 +685,7 @@@ gomp_acc_remove_pointer (void *h, bool } } - t->list[0]->copy_from = force_copyfrom ? 1 : 0; - if (force_copyfrom) - t->list[0].copy_from = 1; ++ t->list[0].copy_from = force_copyfrom ? 1 : 0; gomp_mutex_unlock (&acc_dev->lock); diff --cc libgomp/oacc-parallel.c index 2b90c9f,8ea3dd1..e4ecc87 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@@ -261,16 -135,12 +261,16 @@@ GOACC_parallel_keyed (int device, void devaddrs = gomp_alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) - devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start - + tgt->list[i].key->tgt_offset); + { - if (tgt->list[i] != NULL) - devaddrs[i] = (void *) (tgt->list[i]->tgt->tgt_start - + tgt->list[i]->tgt_offset); ++ if (tgt->list[i].key != NULL) ++ devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start ++ + tgt->list[i].key->tgt_offset); + else + devaddrs[i] = NULL; + } - acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, sizes, kinds, - num_gangs, num_workers, vector_length, async, - tgt); + acc_dev->openacc.exec_func (tgt_fn, mapnum, hostaddrs, devaddrs, + async, dims, tgt); /* If running synchronously, unmap immediately. */ if (async < acc_async_noval) diff --cc libgomp/target.c index 4587361,05c9b71..c2e1996 --- libgomp/target.c +++ libgomp/target.c @@@ -714,9 -721,6 +740,8 @@@ gomp_load_image_to_device (struct gomp_ k->tgt_offset = target_table[i].start; k->refcount = 1; k->async_refcount = 0; - k->copy_from = false; - tgt->list[i] = k; ++ tgt->list[i].key = k; + tgt->refcount++; array->left = NULL; array->right = NULL; splay_tree_insert (&devicep->mem_map, array); @@@ -742,9 -746,6 +767,8 @@@ k->tgt_offset = target_var->start; k->refcount = 1; k->async_refcount = 0; - k->copy_from = false; - tgt->list[i] = k; ++ tgt->list[i].key = k; + tgt->refcount++; array->left = NULL; array->right = NULL; splay_tree_insert (&devicep->mem_map, array); diff --cc libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c index d478ce2,22cef6d..f3b490a --- libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @@@ -1,4 -1,4 +1,6 @@@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ ++/* <http://news.gmane.org/find-root.php?message_id=%3C87pp0aaksc.fsf%40kepler.schwinge.homeip.net%3E>. ++ { dg-xfail-run-if "TODO" { *-*-* } } */ /* { dg-additional-options "-lcuda" } */ #include <openacc.h> Grüße Thomas
On 22/10/15 20:27, Thomas Schwinge wrote: > diff --cc libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c > index d478ce2,22cef6d..f3b490a > --- libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c > @@@ -1,4 -1,4 +1,6 @@@ > /* { dg-do run { target openacc_nvidia_accel_selected } } */ > ++/*<http://news.gmane.org/find-root.php?message_id=%3C87pp0aaksc.fsf%40kepler.schwinge.homeip.net%3E>. > ++ { dg-xfail-run-if "TODO" { *-*-* } } */ > /* { dg-additional-options "-lcuda" } */ > > #include <openacc.h> This failure shows up on trunk. Should it also be xfailed there? Thanks, - Tom
Hi! On Wed, 18 Nov 2015 16:17:39 +0100, Tom de Vries <Tom_deVries@mentor.com> wrote: > On 22/10/15 20:27, Thomas Schwinge wrote: > > diff --cc libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c > > index d478ce2,22cef6d..f3b490a > > --- libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c > > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c > > @@@ -1,4 -1,4 +1,6 @@@ > > /* { dg-do run { target openacc_nvidia_accel_selected } } */ > > ++/*<http://news.gmane.org/find-root.php?message_id=%3C87pp0aaksc.fsf%40kepler.schwinge.homeip.net%3E>. > > ++ { dg-xfail-run-if "TODO" { *-*-* } } */ > > /* { dg-additional-options "-lcuda" } */ > > > > #include <openacc.h> > > This failure shows up on trunk. Should it also be xfailed there? I added the XFAIL as part of my recent r234575 "Update OpenACC test cases" commit, <http://news.gmane.org/find-root.php?message_id=%3C878u109ew4.fsf%40hertz.schwinge.homeip.net%3E>. Chung-Lin, for avoidance of doubt, please remove that XFAIL once you get to commit your fix for this issue (currently waiting for Jakub's approval). Grüße Thomas
diff --git include/gomp-constants.h include/gomp-constants.h index b55f68b..540a31e 100644 --- include/gomp-constants.h +++ include/gomp-constants.h @@ -111,6 +111,12 @@ enum gomp_map_kind #define GOMP_MAP_POINTER_P(X) \ ((X) == GOMP_MAP_POINTER) +#define GOMP_MAP_ALWAYS_TO_P(X) \ + (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM)) + +#define GOMP_MAP_ALWAYS_FROM_P(X) \ + (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM)) + /* Asynchronous behavior. Keep in sync with libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */ diff --git libgomp/libgomp.h libgomp/libgomp.h index d86da7d..8fd7d08 100644 --- libgomp/libgomp.h +++ libgomp/libgomp.h @@ -641,6 +641,15 @@ typedef struct splay_tree_node_s *splay_tree_node; typedef struct splay_tree_s *splay_tree; typedef struct splay_tree_key_s *splay_tree_key; +struct target_var_desc { + /* Splay key. */ + splay_tree_key key; + /* True if data should be copied from device to host at the end. */ + bool copy_from; + /* True if data always should be copied from device to host at the end. */ + bool always_copy_from; +}; + struct target_mem_desc { /* Reference count. */ uintptr_t refcount; @@ -660,9 +669,9 @@ struct target_mem_desc { /* Corresponding target device descriptor. */ struct gomp_device_descr *device_descr; - /* List of splay keys to remove (or decrease refcount) + /* List of target items to remove (or decrease refcount) at the end of region. */ - splay_tree_key list[]; + struct target_var_desc list[]; }; struct splay_tree_key_s { @@ -678,8 +687,6 @@ struct splay_tree_key_s { uintptr_t refcount; /* Asynchronous reference count. */ uintptr_t async_refcount; - /* True if data should be copied from device to host at the end. */ - bool copy_from; }; #include "splay-tree.h" diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c index 7fcf199..a90c912 100644 --- libgomp/oacc-mem.c +++ libgomp/oacc-mem.c @@ -685,7 +685,7 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum) } } - t->list[0]->copy_from = force_copyfrom ? 1 : 0; + t->list[0].copy_from = force_copyfrom ? 1 : 0; gomp_mutex_unlock (&acc_dev->lock); diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index 2b90c9f..e4ecc87 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -262,9 +262,9 @@ GOACC_parallel_keyed (int device, void (*fn) (void *), devaddrs = gomp_alloca (sizeof (void *) * mapnum); for (i = 0; i < mapnum; i++) { - if (tgt->list[i] != NULL) - devaddrs[i] = (void *) (tgt->list[i]->tgt->tgt_start - + tgt->list[i]->tgt_offset); + if (tgt->list[i].key != NULL) + devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start + + tgt->list[i].key->tgt_offset); else devaddrs[i] = NULL; } diff --git libgomp/target.c libgomp/target.c index 4587361..c2e1996 100644 --- libgomp/target.c +++ libgomp/target.c @@ -168,6 +168,12 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, splay_tree_key oldn, (void *) newn->host_start, (void *) newn->host_end, (void *) oldn->host_start, (void *) oldn->host_end); } + + if (GOMP_MAP_ALWAYS_TO_P (kind)) + devicep->host2dev_func (devicep->target_id, + (void *) (oldn->tgt->tgt_start + oldn->tgt_offset), + (void *) newn->host_start, + newn->host_end - newn->host_start); oldn->refcount++; } @@ -267,7 +273,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, int kind = get_kind (short_mapkind, kinds, i); if (hostaddrs[i] == NULL) { - tgt->list[i] = NULL; + tgt->list[i].key = NULL; continue; } cur_node.host_start = (uintptr_t) hostaddrs[i]; @@ -278,12 +284,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); if (n) { - tgt->list[i] = n; + tgt->list[i].key = n; + tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask); + tgt->list[i].always_copy_from + = GOMP_MAP_ALWAYS_FROM_P (kind & typemask); gomp_map_vars_existing (devicep, n, &cur_node, kind & typemask); } else { - tgt->list[i] = NULL; + tgt->list[i].key = NULL; size_t align = (size_t) 1 << (kind >> rshift); not_found_cnt++; @@ -304,7 +313,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, break; else { - tgt->list[j] = NULL; + tgt->list[j].key = NULL; i++; } } @@ -352,7 +361,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, size_t j; for (i = 0; i < mapnum; i++) - if (tgt->list[i] == NULL) + if (tgt->list[i].key == NULL) { int kind = get_kind (short_mapkind, kinds, i); if (hostaddrs[i] == NULL) @@ -366,18 +375,23 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, splay_tree_key n = splay_tree_lookup (mem_map, k); if (n) { - tgt->list[i] = n; + tgt->list[i].key = n; + tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask); + tgt->list[i].always_copy_from + = GOMP_MAP_ALWAYS_FROM_P (kind & typemask); gomp_map_vars_existing (devicep, n, k, kind & typemask); } else { size_t align = (size_t) 1 << (kind >> rshift); - tgt->list[i] = k; + tgt->list[i].key = k; tgt_size = (tgt_size + align - 1) & ~(align - 1); k->tgt = tgt; k->tgt_offset = tgt_size; tgt_size += k->host_end - k->host_start; - k->copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask); + tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask); + tgt->list[i].always_copy_from + = GOMP_MAP_ALWAYS_FROM_P (kind & typemask); k->refcount = 1; k->async_refcount = 0; tgt->refcount++; @@ -395,6 +409,8 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, case GOMP_MAP_TOFROM: case GOMP_MAP_FORCE_TO: case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_TOFROM: /* FIXME: Perhaps add some smarts, like if copying several adjacent fields from host to target, use some host buffer to avoid sending each var individually. */ @@ -427,7 +443,9 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, break; else { - tgt->list[j] = k; + tgt->list[j].key = k; + tgt->list[j].copy_from = false; + tgt->list[j].always_copy_from = false; k->refcount++; gomp_map_pointer (tgt, (uintptr_t) *(void **) hostaddrs[j], @@ -479,11 +497,11 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, { for (i = 0; i < mapnum; i++) { - if (tgt->list[i] == NULL) + if (tgt->list[i].key == NULL) cur_node.tgt_offset = (uintptr_t) NULL; else - cur_node.tgt_offset = tgt->list[i]->tgt->tgt_start - + tgt->list[i]->tgt_offset; + cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start + + tgt->list[i].key->tgt_offset; /* FIXME: see above FIXME comment. */ devicep->host2dev_func (devicep->target_id, (void *) (tgt->tgt_start @@ -523,17 +541,17 @@ gomp_copy_from_async (struct target_mem_desc *tgt) gomp_mutex_lock (&devicep->lock); for (i = 0; i < tgt->list_count; i++) - if (tgt->list[i] == NULL) + if (tgt->list[i].key == NULL) ; - else if (tgt->list[i]->refcount > 1) + else if (tgt->list[i].key->refcount > 1) { - tgt->list[i]->refcount--; - tgt->list[i]->async_refcount++; + tgt->list[i].key->refcount--; + tgt->list[i].key->async_refcount++; } else { - splay_tree_key k = tgt->list[i]; - if (k->copy_from) + splay_tree_key k = tgt->list[i].key; + if (tgt->list[i].copy_from) devicep->dev2host_func (devicep->target_id, (void *) k->host_start, (void *) (k->tgt->tgt_start + k->tgt_offset), k->host_end - k->host_start); @@ -561,25 +579,33 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom) size_t i; for (i = 0; i < tgt->list_count; i++) - if (tgt->list[i] == NULL) - ; - else if (tgt->list[i]->refcount > 1) - tgt->list[i]->refcount--; - else if (tgt->list[i]->async_refcount > 0) - tgt->list[i]->async_refcount--; - else - { - splay_tree_key k = tgt->list[i]; - if (k->copy_from && do_copyfrom) - devicep->dev2host_func (devicep->target_id, (void *) k->host_start, - (void *) (k->tgt->tgt_start + k->tgt_offset), - k->host_end - k->host_start); - splay_tree_remove (&devicep->mem_map, k); - if (k->tgt->refcount > 1) - k->tgt->refcount--; - else - gomp_unmap_tgt (k->tgt); - } + { + splay_tree_key k = tgt->list[i].key; + if (k == NULL) + continue; + + bool do_unmap = false; + if (k->refcount > 1) + k->refcount--; + else if (k->async_refcount > 0) + k->async_refcount--; + else + do_unmap = true; + + if ((do_unmap && do_copyfrom && tgt->list[i].copy_from) + || tgt->list[i].always_copy_from) + devicep->dev2host_func (devicep->target_id, (void *) k->host_start, + (void *) (k->tgt->tgt_start + k->tgt_offset), + k->host_end - k->host_start); + if (do_unmap) + { + splay_tree_remove (&devicep->mem_map, k); + if (k->tgt->refcount > 1) + k->tgt->refcount--; + else + gomp_unmap_tgt (k->tgt); + } + } if (tgt->refcount > 1) tgt->refcount--; @@ -714,8 +740,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt_offset = target_table[i].start; k->refcount = 1; k->async_refcount = 0; - k->copy_from = false; - tgt->list[i] = k; + tgt->list[i].key = k; tgt->refcount++; array->left = NULL; array->right = NULL; @@ -742,8 +767,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version, k->tgt_offset = target_var->start; k->refcount = 1; k->async_refcount = 0; - k->copy_from = false; - tgt->list[i] = k; + tgt->list[i].key = k; tgt->refcount++; array->left = NULL; array->right = NULL; diff --git libgomp/testsuite/libgomp.c/target-11.c libgomp/testsuite/libgomp.c/target-11.c new file mode 100644 index 0000000..4562d88 --- /dev/null +++ libgomp/testsuite/libgomp.c/target-11.c @@ -0,0 +1,51 @@ +/* { dg-require-effective-target offload_device } */ + +#include <assert.h> + +int main () +{ + int aa = 0, bb = 0, cc = 0, dd = 0; + + #pragma omp target data map(tofrom: aa) map(to: bb) map(from: cc, dd) + { + int ok; + aa = bb = cc = 1; + + /* Set dd on target to 0 for the further check. */ + #pragma omp target map(always to: dd) + { dd; } + + dd = 1; + #pragma omp target map(tofrom: aa) map(always to: bb) \ + map(always from: cc) map(to: dd) map(from: ok) + { + /* bb is always to, aa and dd are not. */ + ok = (aa == 0) && (bb == 1) && (dd == 0); + aa = bb = cc = dd = 2; + } + + assert (ok); + assert (aa == 1); + assert (bb == 1); + assert (cc == 2); /* cc is always from. */ + assert (dd == 1); + + dd = 3; + #pragma omp target map(from: cc) map(always to: dd) map(from: ok) + { + ok = (dd == 3); /* dd is always to. */ + cc = dd = 4; + } + + assert (ok); + assert (cc == 2); + assert (dd == 3); + } + + assert (aa == 2); + assert (bb == 1); + assert (cc == 4); + assert (dd == 4); + + return 0; +}