From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 63997 invoked by alias); 31 Jul 2015 16:16:19 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 63987 invoked by uid 89); 31 Jul 2015 16:16:19 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL,BAYES_00,KAM_LAZY_DOMAIN_SECURITY,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Fri, 31 Jul 2015 16:16:16 +0000 Received: from int-mx14.intmail.prod.int.phx2.redhat.com (int-mx14.intmail.prod.int.phx2.redhat.com [10.5.11.27]) by mx1.redhat.com (Postfix) with ESMTPS id DC996A10DA; Fri, 31 Jul 2015 16:16:14 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-30.ams2.redhat.com [10.36.116.30]) by int-mx14.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t6VGGDeB002598 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Fri, 31 Jul 2015 12:16:14 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t6VGGBcC030679; Fri, 31 Jul 2015 18:16:11 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t6VGGAOT030678; Fri, 31 Jul 2015 18:16:10 +0200 Date: Fri, 31 Jul 2015 16:28:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: gcc-patches@gcc.gnu.org Subject: [gomp4.1] Start of structure element mapping support Message-ID: <20150731161610.GF1780@tucnak.redhat.com> Reply-To: Jakub Jelinek MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-07/txt/msg02685.txt.bz2 Hi! This patch is the start of implementation of struct element mapping. I'm not handling structure element based array sections (neither array based, nor pointer/reference based) yet, nor C++. If the whole struct is already mapped, then that mapping is used, otherwise we require that either all the fields are already mapped, or none of them (otherwise runtime error). If none of them, then we allocate enough room for the first to last mapped field, and place all the individual allocations into the allocated space. 2015-07-31 Jakub Jelinek * gimplify.c (gimplify_scan_omp_clauses): Handle map clauses with COMPONENT_REF. * omp-low.c (lower_omp_target): Handle GOMP_MAP_STRUCT. Handle GOMP_MAP_RELEASE for zero-length array sections. * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_STRUCT. gcc/c/ * c-parser.c (c_parser_omp_variable_list): Parse struct elements in map/to/from clauses. * c-typeck.c (handle_omp_array_sections): Handle GOMP_MAP_RELEASE for zero-length array sections. (c_finish_omp_clauses): Handle struct elements in map/to/from OpenMP clauses. gcc/cp/ * semantics.c (handle_omp_array_sections): Handle GOMP_MAP_RELEASE for zero-length array sections. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_STRUCT. libgomp/ * target.c (gomp_map_fields_existing): New function. (gomp_map_vars): Handle GOMP_MAP_STRUCT. * testsuite/libgomp.c/target-21.c: New test. --- gcc/gimplify.c.jj 2015-07-31 16:55:01.482411392 +0200 +++ gcc/gimplify.c 2015-07-31 16:57:22.307320290 +0200 @@ -6202,6 +6202,7 @@ gimplify_scan_omp_clauses (tree *list_p, { struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; + hash_map *struct_map_to_clause = NULL; ctx = new_omp_context (region_type); outer_ctx = ctx->outer_context; @@ -6442,6 +6443,11 @@ gimplify_scan_omp_clauses (tree *list_p, } if (!DECL_P (decl)) { + if (TREE_CODE (decl) == COMPONENT_REF) + { + while (TREE_CODE (decl) == COMPONENT_REF) + decl = TREE_OPERAND (decl, 0); + } if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) @@ -6449,6 +6455,128 @@ gimplify_scan_omp_clauses (tree *list_p, remove = true; break; } + if (DECL_P (decl)) + { + if (error_operand_p (decl)) + { + remove = true; + break; + } + + if (TYPE_SIZE_UNIT (TREE_TYPE (decl)) == NULL + || (TREE_CODE (TYPE_SIZE_UNIT (TREE_TYPE (decl))) + != INTEGER_CST)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "mapping field %qE of variable length " + "structure", OMP_CLAUSE_DECL (c)); + remove = true; + break; + } + + tree offset; + HOST_WIDE_INT bitsize, bitpos; + machine_mode mode; + int unsignedp, volatilep = 0; + tree base + = get_inner_reference (OMP_CLAUSE_DECL (c), &bitsize, + &bitpos, &offset, &mode, &unsignedp, + &volatilep, false); + gcc_assert (base == decl + && (offset == NULL_TREE + || TREE_CODE (offset) == INTEGER_CST)); + + splay_tree_node n + = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); + if (n == NULL || (n->value & GOVD_MAP) == 0) + { + *list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT); + OMP_CLAUSE_DECL (*list_p) = decl; + OMP_CLAUSE_SIZE (*list_p) = size_int (1); + OMP_CLAUSE_CHAIN (*list_p) = c; + if (struct_map_to_clause == NULL) + struct_map_to_clause = new hash_map; + struct_map_to_clause->put (decl, *list_p); + list_p = &OMP_CLAUSE_CHAIN (*list_p); + flags = GOVD_MAP | GOVD_EXPLICIT; + if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) + flags |= GOVD_SEEN; + goto do_add_decl; + } + else + { + tree *osc = struct_map_to_clause->get (decl), *sc; + if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) + n->value |= GOVD_SEEN; + offset_int o1, o2; + if (offset) + o1 = wi::to_offset (offset); + else + o1 = 0; + if (bitpos) + o1 = o1 + bitpos / BITS_PER_UNIT; + for (sc = &OMP_CLAUSE_CHAIN (*osc); *sc != c; + sc = &OMP_CLAUSE_CHAIN (*sc)) + if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF) + break; + else + { + tree offset2; + HOST_WIDE_INT bitsize2, bitpos2; + base = get_inner_reference (OMP_CLAUSE_DECL (*sc), + &bitsize2, &bitpos2, + &offset2, &mode, + &unsignedp, &volatilep, + false); + if (base != decl) + break; + gcc_assert (offset == NULL_TREE + || TREE_CODE (offset) == INTEGER_CST); + tree d1 = OMP_CLAUSE_DECL (*sc); + tree d2 = OMP_CLAUSE_DECL (c); + while (TREE_CODE (d1) == COMPONENT_REF) + if (TREE_CODE (d2) == COMPONENT_REF + && TREE_OPERAND (d1, 1) + == TREE_OPERAND (d2, 1)) + { + d1 = TREE_OPERAND (d1, 0); + d2 = TREE_OPERAND (d2, 0); + } + else + break; + if (d1 == d2) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE appears more than once in map " + "clauses", OMP_CLAUSE_DECL (c)); + remove = true; + break; + } + if (offset2) + o2 = wi::to_offset (offset2); + else + o2 = 0; + if (bitpos2) + o2 = o2 + bitpos2 / BITS_PER_UNIT; + if (wi::ltu_p (o1, o2) + || (wi::eq_p (o1, o2) && bitpos < bitpos2)) + break; + } + if (!remove) + OMP_CLAUSE_SIZE (*osc) + = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), + size_one_node); + if (!remove && *sc != c) + { + *list_p = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = *sc; + *sc = c; + continue; + } + } + } break; } flags = GOVD_MAP | GOVD_EXPLICIT; @@ -6790,6 +6918,8 @@ gimplify_scan_omp_clauses (tree *list_p, } gimplify_omp_ctxp = ctx; + if (struct_map_to_clause) + delete struct_map_to_clause; } struct gimplify_adjust_omp_clauses_data --- gcc/omp-low.c.jj 2015-07-31 16:55:01.272414510 +0200 +++ gcc/omp-low.c 2015-07-31 16:57:22.317320141 +0200 @@ -12954,6 +12954,7 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_STRUCT: break; case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: @@ -13303,6 +13304,7 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION; break; default: --- gcc/tree-pretty-print.c.jj 2015-07-31 16:55:01.484411362 +0200 +++ gcc/tree-pretty-print.c 2015-07-31 16:57:22.320320097 +0200 @@ -643,6 +643,9 @@ dump_omp_clause (pretty_printer *pp, tre case GOMP_MAP_FIRSTPRIVATE_POINTER: pp_string (pp, "firstprivate"); break; + case GOMP_MAP_STRUCT: + pp_string (pp, "struct"); + break; default: gcc_unreachable (); } --- gcc/c/c-parser.c.jj 2015-07-31 16:55:01.481411407 +0200 +++ gcc/c/c-parser.c 2015-07-31 16:57:22.313320201 +0200 @@ -10190,10 +10190,25 @@ c_parser_omp_variable_list (c_parser *pa t = error_mark_node; break; } - /* FALL THROUGH. */ + /* FALLTHROUGH */ case OMP_CLAUSE_MAP: case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: + while (c_parser_next_token_is (parser, CPP_DOT)) + { + location_t op_loc = c_parser_peek_token (parser)->location; + c_parser_consume_token (parser); + if (!c_parser_next_token_is (parser, CPP_NAME)) + { + c_parser_error (parser, "expected identifier"); + t = error_mark_node; + break; + } + tree ident = c_parser_peek_token (parser)->value; + c_parser_consume_token (parser); + t = build_component_ref (op_loc, t, ident); + } + /* FALLTHROUGH */ case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_REDUCTION: while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE)) --- gcc/c/c-typeck.c.jj 2015-07-31 16:55:01.482411392 +0200 +++ gcc/c/c-typeck.c 2015-07-31 16:58:09.246623290 +0200 @@ -12040,6 +12040,7 @@ handle_omp_array_sections (tree c, bool case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; break; default: @@ -12117,7 +12118,7 @@ tree c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, map_head; + bitmap_head aligned_head, map_head, map_field_head; tree c, t, type, *pc; tree simdlen = NULL_TREE, safelen = NULL_TREE; bool branch_seen = false; @@ -12130,6 +12131,7 @@ c_finish_omp_clauses (tree clauses, bool bitmap_initialize (&lastprivate_head, &bitmap_default_obstack); bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); + bitmap_initialize (&map_field_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -12574,8 +12576,49 @@ c_finish_omp_clauses (tree clauses, bool break; } if (t == error_mark_node) - remove = true; - else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) + { + remove = true; + break; + } + if (TREE_CODE (t) == COMPONENT_REF + && is_omp + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) + { + if (DECL_BIT_FIELD (TREE_OPERAND (t, 1))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "bit-field %qE in %qs clause", + t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + else if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE does not have a mappable type in %qs clause", + t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]); + remove = true; + } + while (TREE_CODE (t) == COMPONENT_REF) + { + if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) + == UNION_TYPE) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qE is a member of a union", t); + remove = true; + break; + } + t = TREE_OPERAND (t, 0); + } + if (remove) + break; + if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) + { + if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + break; + } + } + if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) { error_at (OMP_CLAUSE_LOCATION (c), "%qE is not a variable in %qs clause", t, @@ -12597,6 +12640,7 @@ c_finish_omp_clauses (tree clauses, bool == GOMP_MAP_FIRSTPRIVATE_POINTER) || (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR))) + && t == OMP_CLAUSE_DECL (c) && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) { error_at (OMP_CLAUSE_LOCATION (c), @@ -12613,7 +12657,12 @@ c_finish_omp_clauses (tree clauses, bool remove = true; } else - bitmap_set_bit (&map_head, DECL_UID (t)); + { + bitmap_set_bit (&map_head, DECL_UID (t)); + if (t != OMP_CLAUSE_DECL (c) + && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF) + bitmap_set_bit (&map_field_head, DECL_UID (t)); + } break; case OMP_CLAUSE_TO_DECLARE: --- gcc/cp/semantics.c.jj 2015-07-31 16:55:01.485411348 +0200 +++ gcc/cp/semantics.c 2015-07-31 16:57:22.303320349 +0200 @@ -4836,6 +4836,7 @@ handle_omp_array_sections (tree c, bool case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; break; default: --- include/gomp-constants.h.jj 2015-07-31 16:55:01.604409581 +0200 +++ include/gomp-constants.h 2015-07-31 16:55:38.711858574 +0200 @@ -102,6 +102,14 @@ enum gomp_map_kind /* If not already present, allocate. And unconditionally copy to and from device. */ GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_TOFROM), + /* Map a sparse struct; the address is the base of the structure, alignment + it's required alignment, and size is the number of adjacent entries + that belong to the struct. The adjacent entries should be sorted by + increasing address, so it is easy to determine lowest needed address + (address of the first adjacent entry) and highest needed address + (address of the last adjacent entry plus its size). */ + GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_ALWAYS + | GOMP_MAP_FLAG_SPECIAL | 0), /* OpenMP 4.1 alias for forced deallocation. */ GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC, /* Decrement usage count and deallocate if zero. */ --- libgomp/target.c.jj 2015-07-31 16:55:01.981403983 +0200 +++ libgomp/target.c 2015-07-31 16:55:38.710858589 +0200 @@ -245,6 +245,66 @@ gomp_map_pointer (struct target_mem_desc sizeof (void *)); } +static void +gomp_map_fields_existing (struct target_mem_desc *tgt, splay_tree_key n, + size_t first, size_t i, void **hostaddrs, + size_t *sizes, void *kinds) +{ + struct gomp_device_descr *devicep = tgt->device_descr; + struct splay_tree_s *mem_map = &devicep->mem_map; + struct splay_tree_key_s cur_node; + int kind; + const bool short_mapkind = true; + const int typemask = short_mapkind ? 0xff : 0x7; + + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizes[i]; + splay_tree_key n2 = splay_tree_lookup (mem_map, &cur_node); + kind = get_kind (short_mapkind, kinds, i); + if (n2 + && n2->tgt == n->tgt + && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) + { + gomp_map_vars_existing (devicep, n2, &cur_node, + &tgt->list[i], kind & typemask); + return; + } + if (sizes[i] == 0) + { + if (cur_node.host_start > (uintptr_t) hostaddrs[first - 1]) + { + cur_node.host_start--; + n2 = splay_tree_lookup (mem_map, &cur_node); + cur_node.host_start++; + if (n2 + && n2->tgt == n->tgt + && n2->host_start - n->host_start + == n2->tgt_offset - n->tgt_offset) + { + gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i], + kind & typemask); + return; + } + } + cur_node.host_end++; + n2 = splay_tree_lookup (mem_map, &cur_node); + cur_node.host_end--; + if (n2 + && n2->tgt == n->tgt + && n2->host_start - n->host_start == n2->tgt_offset - n->tgt_offset) + { + gomp_map_vars_existing (devicep, n2, &cur_node, &tgt->list[i], + kind & typemask); + return; + } + } + gomp_mutex_unlock (&devicep->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); +} + attribute_hidden struct target_mem_desc * gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds, @@ -304,6 +364,37 @@ gomp_map_vars (struct gomp_device_descr tgt->list[i].offset = ~(uintptr_t) 0; continue; } + else if ((kind & typemask) == GOMP_MAP_STRUCT) + { + size_t first = i + 1; + size_t last = i + sizes[i]; + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = (uintptr_t) hostaddrs[last] + + sizes[last]; + tgt->list[i].key = NULL; + tgt->list[i].offset = ~(uintptr_t) 2; + splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + if (n == NULL) + { + size_t align = (size_t) 1 << (kind >> rshift); + if (tgt_align < align) + tgt_align = align; + tgt_size -= (uintptr_t) hostaddrs[first] + - (uintptr_t) hostaddrs[i]; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += cur_node.host_end - (uintptr_t) hostaddrs[i]; + not_found_cnt += last - i; + for (i = first; i <= last; i++) + tgt->list[i].key = NULL; + i--; + continue; + } + for (i = first; i <= last; i++) + gomp_map_fields_existing (tgt, n, first, i, hostaddrs, + sizes, kinds); + i--; + 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]; @@ -406,7 +497,8 @@ gomp_map_vars (struct gomp_device_descr if (not_found_cnt) tgt->array = gomp_malloc (not_found_cnt * sizeof (*tgt->array)); splay_tree_node array = tgt->array; - size_t j; + size_t j, field_tgt_offset = 0, field_tgt_clear = ~(size_t) 0; + uintptr_t field_tgt_base = 0; for (i = 0; i < mapnum; i++) if (tgt->list[i].key == NULL) @@ -414,24 +506,53 @@ gomp_map_vars (struct gomp_device_descr int kind = get_kind (short_mapkind, kinds, i); if (hostaddrs[i] == NULL) continue; - if ((kind & typemask) == GOMP_MAP_FIRSTPRIVATE) + switch (kind & typemask) { - size_t align = (size_t) 1 << (kind >> rshift); + size_t align, len, first, last; + splay_tree_key n; + case GOMP_MAP_FIRSTPRIVATE: + align = (size_t) 1 << (kind >> rshift); tgt_size = (tgt_size + align - 1) & ~(align - 1); tgt->list[i].offset = tgt_size; - size_t len = sizes[i]; + len = sizes[i]; devicep->host2dev_func (devicep->target_id, (void *) (tgt->tgt_start + tgt_size), (void *) hostaddrs[i], len); tgt_size += len; continue; - } - switch (kind & typemask) - { case GOMP_MAP_FIRSTPRIVATE_INT: case GOMP_MAP_USE_DEVICE_PTR: case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: continue; + case GOMP_MAP_STRUCT: + first = i + 1; + last = i + sizes[i]; + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = (uintptr_t) hostaddrs[last] + + sizes[last]; + if (tgt->list[first].key != NULL) + continue; + n = splay_tree_lookup (mem_map, &cur_node); + if (n == NULL) + { + size_t align = (size_t) 1 << (kind >> rshift); + tgt_size -= (uintptr_t) hostaddrs[first] + - (uintptr_t) hostaddrs[i]; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += (uintptr_t) hostaddrs[first] + - (uintptr_t) hostaddrs[i]; + field_tgt_base = (uintptr_t) hostaddrs[first]; + field_tgt_offset = tgt_size; + field_tgt_clear = last; + tgt_size += cur_node.host_end + - (uintptr_t) hostaddrs[first]; + continue; + } + for (i = first; i <= last; i++) + gomp_map_fields_existing (tgt, n, first, i, hostaddrs, + sizes, kinds); + i--; + continue; default: break; } @@ -449,10 +570,20 @@ gomp_map_vars (struct gomp_device_descr { size_t align = (size_t) 1 << (kind >> rshift); 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; + if (field_tgt_clear != ~(size_t) 0) + { + k->tgt_offset = k->host_start - field_tgt_base + + field_tgt_offset; + if (i == field_tgt_clear) + field_tgt_clear = ~(size_t) 0; + } + else + { + tgt_size = (tgt_size + align - 1) & ~(align - 1); + k->tgt_offset = tgt_size; + tgt_size += k->host_end - k->host_start; + } 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); @@ -571,6 +702,12 @@ gomp_map_vars (struct gomp_device_descr cur_node.tgt_offset = (uintptr_t) hostaddrs[i]; else if (tgt->list[i].offset == ~(uintptr_t) 1) cur_node.tgt_offset = 0; + else if (tgt->list[i].offset == ~(uintptr_t) 2) + cur_node.tgt_offset = tgt->list[i + 1].key->tgt->tgt_start + + tgt->list[i + 1].key->tgt_offset + + tgt->list[i + 1].offset + + (uintptr_t) hostaddrs[i] + - (uintptr_t) hostaddrs[i + 1]; else cur_node.tgt_offset = tgt->tgt_start + tgt->list[i].offset; --- libgomp/testsuite/libgomp.c/target-21.c.jj 2015-07-31 17:00:30.415527080 +0200 +++ libgomp/testsuite/libgomp.c/target-21.c 2015-07-31 17:32:56.098638516 +0200 @@ -0,0 +1,55 @@ +extern void abort (void); +union U { int x; long long y; }; +struct T { int a; union U b; int c; }; +struct S { int s; int u; struct T v; union U w; }; + +int +main () +{ + struct S s; + s.s = 0; + s.u = 1; + s.v.a = 2; + s.v.b.y = 3LL; + s.v.c = 19; + s.w.x = 4; + int err = 0; + #pragma omp target map (to:s.v.b, s.u) map (from: s.w, err) + { + err = 0; + if (s.u != 1 || s.v.b.y != 3LL) + err = 1; + s.w.x = 6; + } + if (err || s.w.x != 6) + abort (); + s.u++; + s.v.a++; + s.v.b.y++; + s.w.x++; + #pragma omp target data map (tofrom: s) + #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b) + { + err = 0; + if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7) + err = 1; + s.w.x = 8; + } + if (err || s.w.x != 8) + abort (); + s.u++; + s.v.a++; + s.v.b.y++; + s.w.x++; + #pragma omp target data map (from: s.w) map (to: s.v.b, s.u) + #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b) + { + err = 0; + if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9) + err = 1; + s.w.x = 11; + } + if (err || s.w.x != 11) + abort (); + return 0; +} Jakub