From: Julian Brown Date: Fri, 31 Jul 2020 16:02:56 +0000 (-0700) Subject: Merge branch 'releases/gcc-10' into devel/omp/gcc-10 X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=4f0f8bf894d4a8d14f3f1f49942116bf315036a8;p=thirdparty%2Fgcc.git Merge branch 'releases/gcc-10' into devel/omp/gcc-10 --- 4f0f8bf894d4a8d14f3f1f49942116bf315036a8 diff --cc gcc/gimplify.c index 37ea9e0673cb,dfee6b7e7aac..9484117a8bce --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@@ -8839,27 -8760,14 +8839,20 @@@ gimplify_scan_omp_clauses (tree *list_p case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: - case OACC_HOST_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - /* For target {,enter ,exit }data only the array slice is - mapped, but not the pointer to it. */ - remove = true; - break; case OACC_ENTER_DATA: case OACC_EXIT_DATA: - if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + case OACC_HOST_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) ++ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) + && !(prev_list_p + && OMP_CLAUSE_CODE (*prev_list_p) == OMP_CLAUSE_MAP + && ((OMP_CLAUSE_MAP_KIND (*prev_list_p) + == GOMP_MAP_DECLARE_ALLOCATE) + || (OMP_CLAUSE_MAP_KIND (*prev_list_p) + == GOMP_MAP_DECLARE_DEALLOCATE)))) + /* For target {,enter ,exit }data only the array slice is + mapped, but not the pointer to it. */ remove = true; break; default: diff --cc gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 8eb7451eaf63,ca642156e9fd..266ead351926 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@@ -20,8 -20,8 +20,8 @@@ ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5)) -! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=.\\)\\) parm\\.0\\.data - \\(integer\\(kind=.\\)\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } - ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } + ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } @@@ -32,6 -32,6 +32,6 @@@ ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE -! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=.\\)\\) parm\\.1\\.data - \\(integer\\(kind=.\\)\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } } - ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } + ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } END SUBROUTINE f diff --cc libgomp/oacc-mem.c index 884a81b7a050,65757ab2ffca..df4fc0d05133 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@@ -1119,19 -1076,144 +1127,151 @@@ goacc_enter_data_internal (struct gomp_ void **hostaddrs, size_t *sizes, unsigned short *kinds, goacc_aq aq) { + gomp_mutex_lock (&acc_dev->lock); + for (size_t i = 0; i < mapnum; i++) { - int group_last = find_group_last (i, mapnum, sizes, kinds); + splay_tree_key n; + size_t group_last = find_group_last (i, mapnum, sizes, kinds); + bool struct_p = false; + size_t size, groupnum = (group_last - i) + 1; + + switch (kinds[i] & 0xff) + { + case GOMP_MAP_STRUCT: + { + size = (uintptr_t) hostaddrs[group_last] + sizes[group_last] + - (uintptr_t) hostaddrs[i]; + struct_p = true; + } + break; + + case GOMP_MAP_ATTACH: + size = sizeof (void *); + break; + + default: + size = sizes[i]; + } + + n = lookup_host (acc_dev, hostaddrs[i], size); - if (n && struct_p) + if ((kinds[i] & 0xff) == GOMP_MAP_DECLARE_ALLOCATE) - gomp_acc_declare_allocate (true, group_last > i, &hostaddrs[i], - &sizes[i], &kinds[i]); - else - gomp_map_vars_async (acc_dev, aq, - (group_last - i) + 1, - &hostaddrs[i], NULL, - &sizes[i], &kinds[i], true, - GOMP_MAP_VARS_OPENACC_ENTER_DATA); ++ { ++ gomp_mutex_unlock (&acc_dev->lock); ++ gomp_acc_declare_allocate (true, group_last > i, &hostaddrs[i], ++ &sizes[i], &kinds[i]); ++ gomp_mutex_lock (&acc_dev->lock); ++ } ++ else if (n && struct_p) + { + for (size_t j = i + 1; j <= group_last; j++) + { + struct splay_tree_key_s cur_node; + cur_node.host_start = (uintptr_t) hostaddrs[j]; + cur_node.host_end = cur_node.host_start + sizes[j]; + splay_tree_key n2 + = splay_tree_lookup (&acc_dev->mem_map, &cur_node); + if (!n2 + || n2->tgt != n->tgt + || n2->host_start - n->host_start + != n2->tgt_offset - n->tgt_offset) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("Trying to map into device [%p..%p) structure " + "element when other mapped elements from the " + "same structure weren't mapped together with " + "it", (void *) cur_node.host_start, + (void *) cur_node.host_end); + } + } + /* This is a special case because we must increment the refcount by + the number of mapped struct elements, rather than by one. */ + if (n->refcount != REFCOUNT_INFINITY) + n->refcount += groupnum - 1; + n->dynamic_refcount += groupnum - 1; + } + else if (n && groupnum == 1) + { + void *h = hostaddrs[i]; + size_t s = sizes[i]; + + if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH) + { + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, + (uintptr_t) h, s, NULL); + /* OpenACC 'attach'/'detach' doesn't affect structured/dynamic + reference counts ('n->refcount', 'n->dynamic_refcount'). */ + } + else + goacc_map_var_existing (acc_dev, h, s, n); + } + else if (n && groupnum > 1) + { + assert (n->refcount != REFCOUNT_INFINITY + && n->refcount != REFCOUNT_LINK); + + for (size_t j = i + 1; j <= group_last; j++) + if ((kinds[j] & 0xff) == GOMP_MAP_ATTACH) + { + splay_tree_key m + = lookup_host (acc_dev, hostaddrs[j], sizeof (void *)); + gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m, + (uintptr_t) hostaddrs[j], sizes[j], NULL); + } + + bool processed = false; + + struct target_mem_desc *tgt = n->tgt; + for (size_t j = 0; j < tgt->list_count; j++) + if (tgt->list[j].key == n) + { + /* We are processing a group of mappings (e.g. + [GOMP_MAP_TO, GOMP_MAP_TO_PSET, GOMP_MAP_POINTER]). + Find the right group in the target_mem_desc's variable + list, and increment the refcounts for each item in that + group. */ + for (size_t k = 0; k < groupnum; k++) + if (j + k < tgt->list_count + && tgt->list[j + k].key + && !tgt->list[j + k].is_attach) + { + tgt->list[j + k].key->refcount++; + tgt->list[j + k].key->dynamic_refcount++; + } + processed = true; + break; + } + + if (!processed) + { + gomp_mutex_unlock (&acc_dev->lock); + gomp_fatal ("dynamic refcount incrementing failed for " + "pointer/pset"); + } + } + else if (hostaddrs[i]) + { + /* The data is not mapped already. Map it now, unless the first + member in the group has a NULL pointer (e.g. a non-present + optional parameter). */ + gomp_mutex_unlock (&acc_dev->lock); + + struct target_mem_desc *tgt + = gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL, + &sizes[i], &kinds[i], true, + GOMP_MAP_VARS_ENTER_DATA); + assert (tgt); + + gomp_mutex_lock (&acc_dev->lock); + + for (size_t j = 0; j < tgt->list_count; j++) + { + n = tgt->list[j].key; + if (n && !tgt->list[j].is_attach) + n->dynamic_refcount++; + } + } i = group_last; } diff --cc libgomp/target.c index 35a65cc27937,3e292eb8c627..a1558a2711a8 --- a/libgomp/target.c +++ b/libgomp/target.c @@@ -673,11 -666,9 +673,10 @@@ gomp_map_vars_internal (struct gomp_dev struct splay_tree_s *mem_map = &devicep->mem_map; struct splay_tree_key_s cur_node; struct target_mem_desc *tgt - = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); - tgt->list_count = mapnum; + = gomp_malloc (sizeof (*tgt) + + sizeof (tgt->list[0]) * (mapnum + nca_data_row_num)); + tgt->list_count = mapnum + nca_data_row_num; - tgt->refcount = (pragma_kind == GOMP_MAP_VARS_ENTER_DATA - || pragma_kind == GOMP_MAP_VARS_OPENACC_ENTER_DATA) ? 0 : 1; + tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1; tgt->device_descr = devicep; tgt->prev = NULL; struct gomp_coalesce_buf cbuf, *cbufp = NULL; @@@ -825,28 -815,6 +824,28 @@@ has_firstprivate = true; continue; } + else if (GOMP_MAP_NONCONTIG_ARRAY_P (kind & typemask)) + { + /* Ignore non-contiguous arrays for now, we process them together + later. */ + tgt->list[i].key = NULL; + tgt->list[i].offset = 0; + not_found_cnt++; + + /* The map for the non-contiguous array itself is never copied from + during unmapping, its the data rows that count. Set copy-from + flags to false here. */ + tgt->list[i].copy_from = false; + tgt->list[i].always_copy_from = false; - tgt->list[i].do_detach = false; ++ tgt->list[i].is_attach = false; + + size_t align = (size_t) 1 << (kind >> rshift); + if (tgt_align < align) + tgt_align = align; + + continue; + } + cur_node.host_start = (uintptr_t) hostaddrs[i]; if (!GOMP_MAP_POINTER_P (kind & typemask)) cur_node.host_end = cur_node.host_start + sizes[i]; @@@ -1344,100 -1264,6 +1344,100 @@@ array++; } } + + /* Processing of non-contiguous array rows. */ + if (nca_info) + { + struct target_var_desc *next_var_desc = &tgt->list[mapnum]; + for (i = 0; i < nca_info->num_ncarray; i++) + { + struct goacc_ncarray *nca = &nca_info->ncarray[i]; + int kind = get_kind (short_mapkind, kinds, nca->map_index); + size_t align = (size_t) 1 << (kind >> rshift); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + + assert (nca->ptr == hostaddrs[nca->map_index]); + + /* For the map of the non-contiguous array itself, adjust so that + the passed device address points to the beginning of the + ptrblock. Remember to adjust the first-dimension's bias here. */ + tgt->list[nca->map_index].key->tgt_offset + = tgt_size - nca->descr->dims[0].base; + + void *target_ptrblock = (void*) tgt->tgt_start + tgt_size; + tgt_size += nca->ptrblock_size; + + /* Add splay key for each data row in current non-contiguous + array. */ + for (size_t j = 0; j < nca->data_row_num; j++) + { + struct target_var_desc *row_desc = next_var_desc++; + void *row = nca->data_rows[j]; + cur_node.host_start = (uintptr_t) row; + cur_node.host_end = cur_node.host_start + nca->data_row_size; + splay_tree_key k = splay_tree_lookup (mem_map, &cur_node); + if (k) + { + assert (k->refcount != REFCOUNT_LINK); + gomp_map_vars_existing (devicep, aq, k, &cur_node, row_desc, + kind & typemask, + cbufp); + } + else + { + tgt->refcount++; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + + k = &array->key; + k->host_start = (uintptr_t) row; + k->host_end = k->host_start + nca->data_row_size; + + k->tgt = tgt; + k->refcount = 1; - k->virtual_refcount = 0; ++ k->dynamic_refcount = 0; + k->aux = NULL; + k->tgt_offset = tgt_size; + + tgt_size += nca->data_row_size; + + row_desc->key = k; + row_desc->copy_from + = GOMP_MAP_COPY_FROM_P (kind & typemask); + row_desc->always_copy_from + = GOMP_MAP_COPY_FROM_P (kind & typemask); - row_desc->do_detach = false; ++ row_desc->is_attach = false; + row_desc->offset = 0; + row_desc->length = nca->data_row_size; + + array->left = NULL; + array->right = NULL; + splay_tree_insert (mem_map, array); + + if (GOMP_MAP_COPY_TO_P (kind & typemask)) + gomp_copy_host2dev (devicep, aq, + (void *) tgt->tgt_start + k->tgt_offset, + (void *) k->host_start, + nca->data_row_size, + true, cbufp); + array++; + } + nca->tgt_data_rows[j] + = (void *) (k->tgt->tgt_start + k->tgt_offset); + } + + /* Now we have the target memory allocated, and target offsets of all + row blocks assigned and calculated, we can construct the + accelerator side ptrblock and copy it in. */ + if (nca->ptrblock_size) + { + void *ptrblock = goacc_noncontig_array_create_ptrblock + (nca, target_ptrblock); + gomp_copy_host2dev (devicep, aq, target_ptrblock, ptrblock, + nca->ptrblock_size, true, cbufp); + free (ptrblock); + } + } + } } if (pragma_kind == GOMP_MAP_VARS_TARGET)