Hi! This patch updates the mapping of vars to what has been ratified in OpenMP 4.5, or where left unspecified, hopefully follows the discussed intent.
In particular: 1) for C++ references we map what they refer to, and on target construct privatize the references themselves (but not what they point to, because that is mapped) 2) same var may not be present in both data sharing and mapping clauses 3) structure element based array sections (or C++ references) don't have the structure elements privatized, but mapped with an always pointer store at the start of the region (except exit data; and update doesn't touch the structure elements) 4) omp_target_is_present on one past the last element really is about what mapping starts at that point, so essentially it is checking if the first byte at the specified address is mapped 5) zero length array sections pointing to one past the last element really are about what mapping starts at that point >From the above, 3) is really not specified in the standard and just based on the discussions we had, hopefully OpenMP 5.0 will clarify, and 4)/5) are fuzzy in the standard and also based on the discussions. 2015-11-05 Jakub Jelinek <ja...@redhat.com> gcc/ * gimplify.c (omp_notice_variable): For references check whether what it refers to has mappable type, rather than the reference itself. (gimplify_scan_omp_clauses): Add support for GOMP_MAP_FIRSTPRIVATE_REFERENCE and GOMP_MAP_ALWAYS_POINTER, remove old handling of structure element based array sections. (gimplify_adjust_omp_clauses_1): For implicit references to variables with reference type and when not ref to scalar or ref to pointer, map what they refer to using tofrom and use GOMP_MAP_FIRSTPRIVATE_REFERENCE for the reference. (gimplify_adjust_omp_clauses): Remove GOMP_MAP_ALWAYS_POINTER from target exit data. Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE. Drop OMP_CLAUSE_MAP_PRIVATE support. * omp-low.c (scan_sharing_clauses): Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE, drop OMP_CLAUSE_MAP_PRIVATE support. (lower_omp_target): Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE and GOMP_MAP_ALWAYS_POINTER. Drop OMP_CLAUSE_MAP_PRIVATE support. * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE and GOMP_MAP_ALWAYS_POINTER. Simplify. * tree-vect-stmts.c (vectorizable_simd_clone_call): Add SIMD_CLONE_ARG_TYPE_LINEAR_{REF,VAL,UVAL}_VARIABLE_STEP cases. gcc/c/ * c-parser.c (c_parser_omp_target_data, c_parser_omp_target_enter_data, c_parser_omp_target_exit_data, c_parser_omp_target): Allow GOMP_MAP_ALWAYS_POINTER. * c-typeck.c (handle_omp_array_sections): For structure element based array sections use GOMP_MAP_ALWAYS_POINTER instead of GOMP_MAP_FIRSTPRIVATE_POINTER. (c_finish_omp_clauses): Drop generic_field_head, structure elements are now always mapped even as array section bases, diagnose same var in data sharing and mapping clauses. gcc/cp/ * parser.c (cp_parser_omp_target_data, cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data, cp_parser_omp_target): Allow GOMP_MAP_ALWAYS_POINTER and GOMP_MAP_FIRSTPRIVATE_REFERENCE. * semantics.c (handle_omp_array_sections): For structure element based array sections use GOMP_MAP_ALWAYS_POINTER instead of GOMP_MAP_FIRSTPRIVATE_POINTER. (finish_omp_clauses): Drop generic_field_head, structure elements are now always mapped even as array section bases, diagnose same var in data sharing and mapping clauses. For references map what they refer to using GOMP_MAP_ALWAYS_POINTER for structure elements and GOMP_MAP_FIRSTPRIVATE_REFERENCE otherwise. gcc/testsuite/ * c-c++-common/gomp/clauses-2.c (foo): Adjust for diagnostics of variables in both data sharing and mapping clauses and for structure element based array sections being mapped rather than privatized. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_ALWAYS_POINTER and GOMP_MAP_FIRSTPRIVATE_REFERENCE. libgomp/ * target.c (gomp_map_0len_lookup, gomp_map_val): New inline functions. (gomp_map_vars): Handle GOMP_MAP_ALWAYS_POINTER. For GOMP_MAP_ZERO_LEN_ARRAY_SECTION use gomp_map_0len_lookup. Use gomp_map_val function. (gomp_exit_data): For GOMP_MAP_*ZERO_LEN* use gomp_map_0len_lookup instead of gomp_map_lookup. (omp_target_is_present): Use gomp_map_0len_lookup instead of gomp_map_lookup. * testsuite/libgomp.c/target-12.c (main): Adjust for omp_target_is_present change for one-past-last element. * testsuite/libgomp.c/target-17.c (foo): Drop tests where the same var is both mapped and privatized. * testsuite/libgomp.c/target-19.c (foo): Adjust for different handling of zero-length array sections. * testsuite/libgomp.c/target-29.c: New test. * testsuite/libgomp.c/target-30.c: New test. * testsuite/libgomp.c++/target-14.C: New test. * testsuite/libgomp.c++/target-15.C: New test. * testsuite/libgomp.c++/target-16.C: New test. * testsuite/libgomp.c++/target-17.C: New test. * testsuite/libgomp.c++/target-18.C: New test. * testsuite/libgomp.c++/target-19.C: New test. --- gcc/gimplify.c.jj 2015-11-03 09:21:08.773059315 +0100 +++ gcc/gimplify.c 2015-11-05 10:42:35.772592563 +0100 @@ -5970,8 +5970,13 @@ omp_notice_variable (struct gimplify_omp else if (is_scalar) nflags |= GOVD_FIRSTPRIVATE; } + tree type = TREE_TYPE (decl); if (nflags == flags - && !lang_hooks.types.omp_mappable_type (TREE_TYPE (decl))) + && gimplify_omp_ctxp->target_firstprivatize_array_bases + && lang_hooks.decls.omp_privatize_by_reference (decl)) + type = TREE_TYPE (type); + if (nflags == flags + && !lang_hooks.types.omp_mappable_type (type)) { error ("%qD referenced in target region does not have " "a mappable type", decl); @@ -6226,7 +6231,7 @@ gimplify_scan_omp_clauses (tree *list_p, struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; hash_map<tree, tree> *struct_map_to_clause = NULL; - tree *orig_list_p = list_p; + tree *prev_list_p = NULL; ctx = new_omp_context (region_type); outer_ctx = ctx->outer_context; @@ -6506,7 +6511,9 @@ gimplify_scan_omp_clauses (tree *list_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + 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; @@ -6525,7 +6532,9 @@ gimplify_scan_omp_clauses (tree *list_p, remove = true; break; } - else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) { OMP_CLAUSE_SIZE (c) @@ -6584,6 +6593,25 @@ gimplify_scan_omp_clauses (tree *list_p, break; } + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) + { + /* Error recovery. */ + if (prev_list_p == NULL) + { + remove = true; + break; + } + if (OMP_CLAUSE_CHAIN (*prev_list_p) != c) + { + tree ch = OMP_CLAUSE_CHAIN (*prev_list_p); + if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c) + { + remove = true; + break; + } + } + } + tree offset; HOST_WIDE_INT bitsize, bitpos; machine_mode mode; @@ -6603,56 +6631,64 @@ gimplify_scan_omp_clauses (tree *list_p, splay_tree_node n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); bool ptr = (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_POINTER); - if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE - : GOVD_MAP)) == 0) + == GOMP_MAP_ALWAYS_POINTER); + if (n == NULL || (n->value & GOVD_MAP) == 0) { + tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT); + OMP_CLAUSE_DECL (l) = decl; + OMP_CLAUSE_SIZE (l) = size_int (1); + if (struct_map_to_clause == NULL) + struct_map_to_clause = new hash_map<tree, tree>; + struct_map_to_clause->put (decl, l); if (ptr) { + enum gomp_map_kind mkind + = code == OMP_TARGET_EXIT_DATA + ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_PRIVATE); - OMP_CLAUSE_DECL (c2) = decl; - OMP_CLAUSE_CHAIN (c2) = *orig_list_p; - *orig_list_p = c2; - if (struct_map_to_clause == NULL) - struct_map_to_clause = new hash_map<tree, tree>; - tree *osc; - if (n == NULL || (n->value & GOVD_MAP) == 0) - osc = NULL; - else - osc = struct_map_to_clause->get (decl); - if (osc == NULL) - struct_map_to_clause->put (decl, - tree_cons (NULL_TREE, - c, - NULL_TREE)); - else - *osc = tree_cons (*osc, c, NULL_TREE); - flags = GOVD_PRIVATE | GOVD_EXPLICIT; - goto do_add_decl; + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) + = unshare_expr (OMP_CLAUSE_DECL (c)); + OMP_CLAUSE_CHAIN (c2) = *prev_list_p; + OMP_CLAUSE_SIZE (c2) + = TYPE_SIZE_UNIT (ptr_type_node); + OMP_CLAUSE_CHAIN (l) = c2; + if (OMP_CLAUSE_CHAIN (*prev_list_p) != c) + { + tree c4 = OMP_CLAUSE_CHAIN (*prev_list_p); + tree c3 + = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, mkind); + OMP_CLAUSE_DECL (c3) + = unshare_expr (OMP_CLAUSE_DECL (c4)); + OMP_CLAUSE_SIZE (c3) + = TYPE_SIZE_UNIT (ptr_type_node); + OMP_CLAUSE_CHAIN (c3) = *prev_list_p; + OMP_CLAUSE_CHAIN (c2) = c3; + } + *prev_list_p = l; + prev_list_p = NULL; + } + else + { + OMP_CLAUSE_CHAIN (l) = c; + *list_p = l; + list_p = &OMP_CLAUSE_CHAIN (l); } - *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<tree, tree>; - struct_map_to_clause->put (decl, *list_p); - list_p = &OMP_CLAUSE_CHAIN (*list_p); flags = GOVD_MAP | GOVD_EXPLICIT; - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))) + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr) flags |= GOVD_SEEN; goto do_add_decl; } else { tree *osc = struct_map_to_clause->get (decl); - tree *sc = NULL, *pt = NULL; - if (!ptr && TREE_CODE (*osc) == TREE_LIST) - osc = &TREE_PURPOSE (*osc); - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))) + tree *sc = NULL, *scp = NULL; + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr) n->value |= GOVD_SEEN; offset_int o1, o2; if (offset) @@ -6661,18 +6697,16 @@ gimplify_scan_omp_clauses (tree *list_p, o1 = 0; if (bitpos) o1 = o1 + bitpos / BITS_PER_UNIT; - if (ptr) - pt = osc; - else - sc = &OMP_CLAUSE_CHAIN (*osc); - for (; ptr ? (*pt && (sc = &TREE_VALUE (*pt))) - : *sc != c; - ptr ? (pt = &TREE_CHAIN (*pt)) - : (sc = &OMP_CLAUSE_CHAIN (*sc))) - if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF - && (TREE_CODE (OMP_CLAUSE_DECL (*sc)) - != INDIRECT_REF) - && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_REF) + for (sc = &OMP_CLAUSE_CHAIN (*osc); + *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) + if (ptr && sc == prev_list_p) + break; + else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) + != COMPONENT_REF + && (TREE_CODE (OMP_CLAUSE_DECL (*sc)) + != INDIRECT_REF) + && (TREE_CODE (OMP_CLAUSE_DECL (*sc)) + != ARRAY_REF)) break; else { @@ -6701,6 +6735,8 @@ gimplify_scan_omp_clauses (tree *list_p, &volatilep, false); if (base != decl) break; + if (scp) + continue; gcc_assert (offset == NULL_TREE || TREE_CODE (offset) == INTEGER_CST); tree d1 = OMP_CLAUSE_DECL (*sc); @@ -6739,19 +6775,68 @@ gimplify_scan_omp_clauses (tree *list_p, o2 = o2 + bitpos2 / BITS_PER_UNIT; if (wi::ltu_p (o1, o2) || (wi::eq_p (o1, o2) && bitpos < bitpos2)) - break; + { + if (ptr) + scp = sc; + else + break; + } } + if (remove) + break; + OMP_CLAUSE_SIZE (*osc) + = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), + size_one_node); if (ptr) { - if (!remove) - *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt); - break; + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + tree cl = NULL_TREE; + enum gomp_map_kind mkind + = code == OMP_TARGET_EXIT_DATA + ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; + OMP_CLAUSE_SET_MAP_KIND (c2, mkind); + OMP_CLAUSE_DECL (c2) + = unshare_expr (OMP_CLAUSE_DECL (c)); + OMP_CLAUSE_CHAIN (c2) = scp ? *scp : *prev_list_p; + OMP_CLAUSE_SIZE (c2) + = TYPE_SIZE_UNIT (ptr_type_node); + cl = scp ? *prev_list_p : c2; + if (OMP_CLAUSE_CHAIN (*prev_list_p) != c) + { + tree c4 = OMP_CLAUSE_CHAIN (*prev_list_p); + tree c3 + = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (c3, mkind); + OMP_CLAUSE_DECL (c3) + = unshare_expr (OMP_CLAUSE_DECL (c4)); + OMP_CLAUSE_SIZE (c3) + = TYPE_SIZE_UNIT (ptr_type_node); + OMP_CLAUSE_CHAIN (c3) = *prev_list_p; + if (!scp) + OMP_CLAUSE_CHAIN (c2) = c3; + else + cl = c3; + } + if (scp) + *scp = c2; + if (sc == prev_list_p) + { + *sc = cl; + prev_list_p = NULL; + } + else + { + *prev_list_p = OMP_CLAUSE_CHAIN (c); + list_p = prev_list_p; + prev_list_p = NULL; + OMP_CLAUSE_CHAIN (c) = *sc; + *sc = cl; + continue; + } } - if (!remove) - OMP_CLAUSE_SIZE (*osc) - = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), - size_one_node); - if (!remove && *sc != c) + else if (*sc != c) { *list_p = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = *sc; @@ -6760,6 +6845,13 @@ gimplify_scan_omp_clauses (tree *list_p, } } } + if (!remove + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER + && OMP_CLAUSE_CHAIN (c) + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP + && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_ALWAYS_POINTER)) + prev_list_p = list_p; break; } flags = GOVD_MAP | GOVD_EXPLICIT; @@ -7248,6 +7340,25 @@ gimplify_adjust_omp_clauses_1 (splay_tre OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); OMP_CLAUSE_CHAIN (clause) = nc; } + else if (gimplify_omp_ctxp->target_firstprivatize_array_bases + && lang_hooks.decls.omp_privatize_by_reference (decl)) + { + OMP_CLAUSE_DECL (clause) = build_simple_mem_ref (decl); + OMP_CLAUSE_SIZE (clause) + = unshare_expr (TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl)))); + struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; + gimplify_omp_ctxp = ctx->outer_context; + gimplify_expr (&OMP_CLAUSE_SIZE (clause), + pre_p, NULL, is_gimple_val, fb_rvalue); + gimplify_omp_ctxp = ctx; + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = decl; + OMP_CLAUSE_SIZE (nc) = size_zero_node; + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_REFERENCE); + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); + OMP_CLAUSE_CHAIN (clause) = nc; + } else OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl); } @@ -7375,6 +7486,12 @@ gimplify_adjust_omp_clauses (gimple_seq break; case OMP_CLAUSE_MAP: + if (code == OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) + { + remove = true; + break; + } decl = OMP_CLAUSE_DECL (c); if (!DECL_P (decl)) { @@ -7425,7 +7542,9 @@ gimplify_adjust_omp_clauses (gimple_seq else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) { /* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because for these, TREE_CODE (DECL_SIZE (decl)) will always be @@ -7468,9 +7587,9 @@ gimplify_adjust_omp_clauses (gimple_seq { if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); - if ((n->value & GOVD_SEEN) - && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))) - OMP_CLAUSE_MAP_PRIVATE (c) = 1; + gcc_assert ((n->value & GOVD_SEEN) == 0 + || ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) + == 0)); } break; --- gcc/omp-low.c.jj 2015-11-03 09:21:08.802058898 +0100 +++ gcc/omp-low.c 2015-11-05 10:44:00.003384618 +0100 @@ -2083,7 +2083,9 @@ scan_sharing_clauses (tree clauses, omp_ directly. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && DECL_P (decl) - && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) @@ -2099,7 +2101,9 @@ scan_sharing_clauses (tree clauses, omp_ break; } if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE))) { if (TREE_CODE (decl) == COMPONENT_REF || (TREE_CODE (decl) == INDIRECT_REF @@ -2128,11 +2132,7 @@ scan_sharing_clauses (tree clauses, omp_ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); gcc_assert (DECL_P (decl2)); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_PRIVATE (c)) - install_var_field (decl2, true, 11, ctx); - else - install_var_field (decl2, true, 3, ctx); + install_var_field (decl2, true, 3, ctx); install_var_local (decl2, ctx); install_var_local (decl, ctx); } @@ -2143,9 +2143,6 @@ scan_sharing_clauses (tree clauses, omp_ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_field (decl, true, 7, ctx); - else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_PRIVATE (c)) - install_var_field (decl, true, 11, ctx); else install_var_field (decl, true, 3, ctx); if (is_gimple_omp_offloaded (ctx->stmt)) @@ -2309,7 +2306,9 @@ scan_sharing_clauses (tree clauses, omp_ break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) - && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable) @@ -14363,7 +14362,9 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: case GOMP_MAP_STRUCT: + case GOMP_MAP_ALWAYS_POINTER: break; case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: @@ -14402,7 +14403,8 @@ lower_omp_target (gimple_stmt_iterator * } if (offloaded - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) { if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) { @@ -14421,12 +14423,6 @@ lower_omp_target (gimple_stmt_iterator * continue; } - if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c)) - { - map_cnt++; - continue; - } - if (!maybe_lookup_field (var, ctx)) continue; @@ -14579,7 +14575,9 @@ lower_omp_target (gimple_stmt_iterator * nc = c; ovar = OMP_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE))) break; if (!DECL_P (ovar)) { @@ -14611,14 +14609,7 @@ lower_omp_target (gimple_stmt_iterator * gcc_assert (DECL_P (ovar2)); ovar = ovar2; } - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_PRIVATE (c)) - { - if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar), - ctx)) - continue; - } - else if (!maybe_lookup_field (ovar, ctx)) + if (!maybe_lookup_field (ovar, ctx)) continue; } @@ -14628,12 +14619,7 @@ lower_omp_target (gimple_stmt_iterator * if (nc) { var = lookup_decl_in_outer_ctx (ovar, ctx); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_PRIVATE (c)) - x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), - ctx); - else - x = build_sender_ref (ovar, ctx); + x = build_sender_ref (ovar, ctx); if (maybe_lookup_oacc_reduction (var, ctx)) { gcc_checking_assert (offloaded @@ -15117,7 +15103,7 @@ lower_omp_target (gimple_stmt_iterator * } break; } - /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass, + /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass, so that firstprivate vars holding OMP_CLAUSE_SIZE if needed are already handled. */ for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) @@ -15127,7 +15113,8 @@ lower_omp_target (gimple_stmt_iterator * default: break; case OMP_CLAUSE_MAP: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) { location_t clause_loc = OMP_CLAUSE_LOCATION (c); HOST_WIDE_INT offset = 0; @@ -15181,6 +15168,8 @@ lower_omp_target (gimple_stmt_iterator * } else is_ref = is_reference (var); + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) + is_ref = false; bool ref_to_array = false; if (is_ref) { @@ -15232,8 +15221,10 @@ lower_omp_target (gimple_stmt_iterator * else if (OMP_CLAUSE_CHAIN (c) && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_FIRSTPRIVATE_POINTER) + && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_FIRSTPRIVATE_POINTER + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_FIRSTPRIVATE_REFERENCE))) prev = c; break; case OMP_CLAUSE_PRIVATE: --- gcc/tree-pretty-print.c.jj 2015-11-03 09:21:08.799058941 +0100 +++ gcc/tree-pretty-print.c 2015-11-03 11:58:13.867502798 +0100 @@ -660,9 +660,15 @@ dump_omp_clause (pretty_printer *pp, tre case GOMP_MAP_FIRSTPRIVATE_POINTER: pp_string (pp, "firstprivate"); break; + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + pp_string (pp, "firstprivate ref"); + break; case GOMP_MAP_STRUCT: pp_string (pp, "struct"); break; + case GOMP_MAP_ALWAYS_POINTER: + pp_string (pp, "always_pointer"); + break; default: gcc_unreachable (); } @@ -672,16 +678,22 @@ dump_omp_clause (pretty_printer *pp, tre print_clause_size: if (OMP_CLAUSE_SIZE (clause)) { - if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (clause) - == GOMP_MAP_FIRSTPRIVATE_POINTER)) - pp_string (pp, " [pointer assign, bias: "); - else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET) - pp_string (pp, " [pointer set, len: "); - else - pp_string (pp, " [len: "); + switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP + ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO) + { + case GOMP_MAP_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: + pp_string (pp, " [pointer assign, bias: "); + break; + case GOMP_MAP_TO_PSET: + pp_string (pp, " [pointer set, len: "); + break; + default: + pp_string (pp, " [len: "); + break; + } dump_generic_node (pp, OMP_CLAUSE_SIZE (clause), spc, flags, false); pp_right_bracket (pp); --- gcc/tree-vect-stmts.c.jj 2015-10-14 10:25:50.000000000 +0200 +++ gcc/tree-vect-stmts.c 2015-11-05 10:48:18.025684349 +0100 @@ -2902,6 +2902,9 @@ vectorizable_simd_clone_call (gimple *st case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP: case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP: case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP: + case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP: + case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP: + case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP: /* FORNOW */ i = -1; break; @@ -3174,6 +3177,9 @@ vectorizable_simd_clone_call (gimple *st } break; case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP: + case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP: + case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP: + case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP: default: gcc_unreachable (); } --- gcc/c/c-parser.c.jj 2015-11-03 09:21:09.000000000 +0100 +++ gcc/c/c-parser.c 2015-11-04 14:51:56.710012024 +0100 @@ -14860,6 +14860,7 @@ c_parser_omp_target_data (location_t loc map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_ALWAYS_POINTER: break; default: map_seen |= 1; @@ -14993,6 +14994,7 @@ c_parser_omp_target_enter_data (location map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_ALWAYS_POINTER: break; default: map_seen |= 1; @@ -15079,6 +15081,7 @@ c_parser_omp_target_exit_data (location_ map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_ALWAYS_POINTER: break; default: map_seen |= 1; @@ -15298,6 +15301,7 @@ check_clauses: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_ALWAYS_POINTER: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), --- gcc/c/c-typeck.c.jj 2015-11-03 09:21:08.000000000 +0100 +++ gcc/c/c-typeck.c 2015-11-04 15:17:53.109890507 +0100 @@ -12168,10 +12168,14 @@ handle_omp_array_sections (tree c, bool break; } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, is_omp - ? GOMP_MAP_FIRSTPRIVATE_POINTER - : GOMP_MAP_POINTER); - if (!is_omp && !c_mark_addressable (t)) + if (!is_omp) + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); + else if (TREE_CODE (t) == COMPONENT_REF) + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + else + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER + && !c_mark_addressable (t)) return false; OMP_CLAUSE_DECL (c2) = t; t = build_fold_addr_expr (first); @@ -12239,7 +12243,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, map_field_head, generic_field_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; @@ -12256,7 +12260,6 @@ c_finish_omp_clauses (tree clauses, bool bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); - bitmap_initialize (&generic_field_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -12583,6 +12586,12 @@ c_finish_omp_clauses (tree clauses, bool "%qE appears more than once in data clauses", t); remove = true; } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE + && bitmap_bit_p (&map_head, DECL_UID (t))) + { + error ("%qD appears both in data and map clauses", t); + remove = true; + } else bitmap_set_bit (&generic_head, DECL_UID (t)); break; @@ -12604,6 +12613,11 @@ c_finish_omp_clauses (tree clauses, bool "%qE appears more than once in data clauses", t); remove = true; } + else if (bitmap_bit_p (&map_head, DECL_UID (t))) + { + error ("%qD appears both in data and map clauses", t); + remove = true; + } else bitmap_set_bit (&firstprivate_head, DECL_UID (t)); break; @@ -12795,14 +12809,7 @@ c_finish_omp_clauses (tree clauses, bool break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_POINTER)) - { - if (bitmap_bit_p (&generic_field_head, DECL_UID (t))) - break; - } - else if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + if (bitmap_bit_p (&map_field_head, DECL_UID (t))) break; } } @@ -12845,13 +12852,13 @@ c_finish_omp_clauses (tree clauses, bool error ("%qD appears more than once in data clauses", t); remove = true; } - else + else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - bitmap_set_bit (&generic_head, DECL_UID (t)); - if (t != OMP_CLAUSE_DECL (c) - && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF) - bitmap_set_bit (&generic_field_head, DECL_UID (t)); + error ("%qD appears both in data and map clauses", t); + remove = true; } + else + bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { @@ -12861,6 +12868,12 @@ c_finish_omp_clauses (tree clauses, bool error ("%qD appears more than once in map clauses", t); remove = true; } + else if (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + { + error ("%qD appears both in data and map clauses", t); + remove = true; + } else { bitmap_set_bit (&map_head, DECL_UID (t)); --- gcc/cp/parser.c.jj 2015-11-03 09:21:09.205053109 +0100 +++ gcc/cp/parser.c 2015-11-03 13:31:32.449694248 +0100 @@ -33797,6 +33797,8 @@ cp_parser_omp_target_data (cp_parser *pa map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: break; default: map_seen |= 1; @@ -33888,6 +33890,8 @@ cp_parser_omp_target_enter_data (cp_pars map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: break; default: map_seen |= 1; @@ -33975,6 +33979,8 @@ cp_parser_omp_target_exit_data (cp_parse map_seen = 3; break; case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: break; default: map_seen |= 1; @@ -34238,6 +34244,8 @@ check_clauses: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: case GOMP_MAP_FIRSTPRIVATE_POINTER: + case GOMP_MAP_FIRSTPRIVATE_REFERENCE: + case GOMP_MAP_ALWAYS_POINTER: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), --- gcc/cp/semantics.c.jj 2015-11-03 09:21:08.787059114 +0100 +++ gcc/cp/semantics.c 2015-11-03 16:28:29.133531779 +0100 @@ -4907,9 +4907,20 @@ handle_omp_array_sections (tree c, bool } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER - : GOMP_MAP_POINTER); - if (!is_omp && !cxx_mark_addressable (t)) + if (!is_omp) + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); + else if (TREE_CODE (t) == COMPONENT_REF) + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + else if (REFERENCE_REF_P (t) + && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) + { + t = TREE_OPERAND (t, 0); + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + } + else + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); + if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER + && !cxx_mark_addressable (t)) return false; OMP_CLAUSE_DECL (c2) = t; t = build_fold_addr_expr (first); @@ -4927,15 +4938,18 @@ handle_omp_array_sections (tree c, bool OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; ptr = OMP_CLAUSE_DECL (c2); - if (!is_omp + if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER && TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr)))) { tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_POINTER); + OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2)); OMP_CLAUSE_DECL (c3) = ptr; - OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr); + if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER) + OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr); + else + OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr); OMP_CLAUSE_SIZE (c3) = size_zero_node; OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2); OMP_CLAUSE_CHAIN (c2) = c3; @@ -5659,7 +5673,7 @@ tree finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) { bitmap_head generic_head, firstprivate_head, lastprivate_head; - bitmap_head aligned_head, map_head, map_field_head, generic_field_head; + bitmap_head aligned_head, map_head, map_field_head; tree c, t, *pc; tree safelen = NULL_TREE; bool branch_seen = false; @@ -5673,7 +5687,6 @@ finish_omp_clauses (tree clauses, bool a bitmap_initialize (&aligned_head, &bitmap_default_obstack); bitmap_initialize (&map_head, &bitmap_default_obstack); bitmap_initialize (&map_field_head, &bitmap_default_obstack); - bitmap_initialize (&generic_field_head, &bitmap_default_obstack); for (pc = &clauses, c = clauses; c ; c = *pc) { @@ -5890,6 +5903,12 @@ finish_omp_clauses (tree clauses, bool a error ("%qD appears more than once in data clauses", t); remove = true; } + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE + && bitmap_bit_p (&map_head, DECL_UID (t))) + { + error ("%qD appears both in data and map clauses", t); + remove = true; + } else bitmap_set_bit (&generic_head, DECL_UID (t)); if (!field_ok) @@ -5937,6 +5956,11 @@ finish_omp_clauses (tree clauses, bool a error ("%qD appears more than once in data clauses", t); remove = true; } + else if (bitmap_bit_p (&map_head, DECL_UID (t))) + { + error ("%qD appears both in data and map clauses", t); + remove = true; + } else bitmap_set_bit (&firstprivate_head, DECL_UID (t)); goto handle_field_decl; @@ -6422,7 +6446,10 @@ finish_omp_clauses (tree clauses, bool a } if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) - t = TREE_OPERAND (t, 0); + { + t = TREE_OPERAND (t, 0); + OMP_CLAUSE_DECL (c) = t; + } if (TREE_CODE (t) == COMPONENT_REF && allow_fields && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) @@ -6459,15 +6486,8 @@ finish_omp_clauses (tree clauses, bool a break; if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_POINTER)) - { - if (bitmap_bit_p (&generic_field_head, DECL_UID (t))) - break; - } - else if (bitmap_bit_p (&map_field_head, DECL_UID (t))) - break; + if (bitmap_bit_p (&map_field_head, DECL_UID (t))) + goto handle_map_references; } } if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL) @@ -6475,7 +6495,8 @@ finish_omp_clauses (tree clauses, bool a if (processing_template_decl) break; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)) break; if (DECL_P (t)) error ("%qD is not a variable in %qs clause", t, @@ -6527,17 +6548,13 @@ finish_omp_clauses (tree clauses, bool a error ("%qD appears more than once in data clauses", t); remove = true; } - else + else if (bitmap_bit_p (&map_head, DECL_UID (t))) { - bitmap_set_bit (&generic_head, DECL_UID (t)); - if (t != OMP_CLAUSE_DECL (c) - && (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF - || (REFERENCE_REF_P (OMP_CLAUSE_DECL (c)) - && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), - 0)) - == COMPONENT_REF)))) - bitmap_set_bit (&generic_field_head, DECL_UID (t)); + error ("%qD appears both in data and map clauses", t); + remove = true; } + else + bitmap_set_bit (&generic_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t))) { @@ -6547,6 +6564,12 @@ finish_omp_clauses (tree clauses, bool a error ("%qD appears more than once in map clauses", t); remove = true; } + else if (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + { + error ("%qD appears both in data and map clauses", t); + remove = true; + } else { bitmap_set_bit (&map_head, DECL_UID (t)); @@ -6554,6 +6577,45 @@ finish_omp_clauses (tree clauses, bool a && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF) bitmap_set_bit (&map_field_head, DECL_UID (t)); } + handle_map_references: + if (!remove + && !processing_template_decl + && allow_fields + && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) == REFERENCE_TYPE) + { + t = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + { + OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t); + if (OMP_CLAUSE_SIZE (c) == NULL_TREE) + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t))); + } + else if (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_FIRSTPRIVATE_POINTER + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_FIRSTPRIVATE_REFERENCE) + && (OMP_CLAUSE_MAP_KIND (c) + != GOMP_MAP_ALWAYS_POINTER)) + { + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + if (TREE_CODE (t) == COMPONENT_REF) + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER); + else + OMP_CLAUSE_SET_MAP_KIND (c2, + GOMP_MAP_FIRSTPRIVATE_REFERENCE); + OMP_CLAUSE_DECL (c2) = t; + OMP_CLAUSE_SIZE (c2) = size_zero_node; + OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = c2; + OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t); + if (OMP_CLAUSE_SIZE (c) == NULL_TREE) + OMP_CLAUSE_SIZE (c) + = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t))); + c = c2; + } + } break; case OMP_CLAUSE_TO_DECLARE: --- gcc/testsuite/c-c++-common/gomp/clauses-2.c.jj 2015-11-03 09:21:08.726059990 +0100 +++ gcc/testsuite/c-c++-common/gomp/clauses-2.c 2015-11-04 16:52:53.405837507 +0100 @@ -4,15 +4,15 @@ void bar (int *); void foo (int *p, int q, struct S t, int i, int j, int k, int l) { - #pragma omp target map (q), firstprivate (q) + #pragma omp target map (q), firstprivate (q) /* { dg-error "appears both in data and map clauses" } */ bar (&q); #pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more than once in data clauses" } */ bar (p); #pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */ bar (p); - #pragma omp target map (p[0]) map (p) + #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */ bar (p); - #pragma omp target map (p) , map (p[0]) + #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */ bar (p); #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */ bar (&q); @@ -24,17 +24,17 @@ foo (int *p, int q, struct S t, int i, i bar (&t.r); #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */ bar (&t.r); - #pragma omp target firstprivate (t), map (t.r) + #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */ bar (&t.r); - #pragma omp target map (t.r) firstprivate (t) + #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (&t.r); - #pragma omp target map (t.s[0]) map (t) + #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */ bar (t.s); - #pragma omp target map (t) map(t.s[0]) + #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */ bar (t.s); - #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in data clauses" } */ + #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); - #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in data clauses" } */ + #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ bar (t.s); #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */ bar (t.s); @@ -46,8 +46,8 @@ foo (int *p, int q, struct S t, int i, i bar (t.s); #pragma omp target map (t.r) ,map (t.s[0]) bar (t.s); - #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears more than once in map clauses" } */ - bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 49 } */ - #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in map clauses" } */ - bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 51 } */ + #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */ + bar (t.s); + #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */ + bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } 51 } */ } --- include/gomp-constants.h.jj 2015-10-26 15:38:20.000000000 +0100 +++ include/gomp-constants.h 2015-11-03 10:13:00.621573428 +0100 @@ -111,6 +111,11 @@ enum gomp_map_kind (address of the last adjacent entry plus its size). */ GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_FLAG_SPECIAL | 0), + /* On a location of a pointer/reference that is assumed to be already mapped + earlier, store the translated address of the preceeding mapping. + No refcount is bumped by this, and the store is done unconditionally. */ + GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2 + | GOMP_MAP_FLAG_SPECIAL | 1), /* Forced deallocation of zero length array section. */ GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL_2 @@ -123,7 +128,9 @@ enum gomp_map_kind /* Internal to GCC, not used in libgomp. */ /* Do not map, but pointer assign a pointer instead. */ - GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1) + GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), + /* Do not map, but pointer assign a reference instead. */ + GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2) }; #define GOMP_MAP_COPY_TO_P(X) \ --- libgomp/target.c.jj 2015-11-02 10:44:09.000000000 +0100 +++ libgomp/target.c 2015-11-04 18:46:11.049937173 +0100 @@ -162,7 +162,20 @@ gomp_map_lookup (splay_tree mem_map, spl return splay_tree_lookup (mem_map, key); } -/* Handle the case where gomp_map_lookup found oldn for newn. +static inline splay_tree_key +gomp_map_0len_lookup (splay_tree mem_map, splay_tree_key key) +{ + if (key->host_start != key->host_end) + return splay_tree_lookup (mem_map, key); + + key->host_end++; + splay_tree_key n = splay_tree_lookup (mem_map, key); + key->host_end--; + return n; +} + +/* Handle the case where gomp_map_lookup, splay_tree_lookup or + gomp_map_0len_lookup found oldn for newn. Helper function of gomp_map_vars. */ static inline void @@ -306,6 +319,26 @@ gomp_map_fields_existing (struct target_ (void *) cur_node.host_end); } +static inline uintptr_t +gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) +{ + if (tgt->list[i].key != NULL) + return tgt->list[i].key->tgt->tgt_start + + tgt->list[i].key->tgt_offset + + tgt->list[i].offset; + if (tgt->list[i].offset == ~(uintptr_t) 0) + return (uintptr_t) hostaddrs[i]; + if (tgt->list[i].offset == ~(uintptr_t) 1) + return 0; + if (tgt->list[i].offset == ~(uintptr_t) 2) + return 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]; + return tgt->tgt_start + tgt->list[i].offset; +} + 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, @@ -396,6 +429,13 @@ gomp_map_vars (struct gomp_device_descr i--; continue; } + else if ((kind & typemask) == GOMP_MAP_ALWAYS_POINTER) + { + tgt->list[i].key = NULL; + tgt->list[i].offset = ~(uintptr_t) 1; + has_firstprivate = true; + 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]; @@ -416,7 +456,7 @@ gomp_map_vars (struct gomp_device_descr splay_tree_key n; if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) { - n = gomp_map_lookup (mem_map, &cur_node); + n = gomp_map_0len_lookup (mem_map, &cur_node); if (!n) { tgt->list[i].key = NULL; @@ -554,6 +594,32 @@ gomp_map_vars (struct gomp_device_descr sizes, kinds); i--; continue; + case GOMP_MAP_ALWAYS_POINTER: + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start + sizeof (void *); + n = splay_tree_lookup (mem_map, &cur_node); + if (n == NULL + || n->host_start > cur_node.host_start + || n->host_end < cur_node.host_end) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("always pointer not mapped"); + } + if ((get_kind (short_mapkind, kinds, i - 1) & typemask) + != GOMP_MAP_ALWAYS_POINTER) + cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i - 1); + if (cur_node.tgt_offset) + cur_node.tgt_offset -= sizes[i]; + devicep->host2dev_func (devicep->target_id, + (void *) (n->tgt->tgt_start + + n->tgt_offset + + cur_node.host_start + - n->host_start), + (void *) &cur_node.tgt_offset, + sizeof (void *)); + cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start - n->host_start; + continue; default: break; } @@ -697,26 +763,7 @@ gomp_map_vars (struct gomp_device_descr { for (i = 0; i < mapnum; i++) { - if (tgt->list[i].key == NULL) - { - if (tgt->list[i].offset == ~(uintptr_t) 0) - 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; - } - else - cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start - + tgt->list[i].key->tgt_offset - + tgt->list[i].offset; + cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); /* FIXME: see above FIXME comment. */ devicep->host2dev_func (devicep->target_id, (void *) (tgt->tgt_start @@ -1551,7 +1598,7 @@ gomp_exit_data (struct gomp_device_descr cur_node.host_end = cur_node.host_start + sizes[i]; splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) - ? gomp_map_lookup (&devicep->mem_map, &cur_node) + ? gomp_map_0len_lookup (&devicep->mem_map, &cur_node) : splay_tree_lookup (&devicep->mem_map, &cur_node); if (!k) continue; @@ -1783,7 +1830,7 @@ omp_target_is_present (void *ptr, int de cur_node.host_start = (uintptr_t) ptr; cur_node.host_end = cur_node.host_start; - splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); + splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node); int ret = n != NULL; gomp_mutex_unlock (&devicep->lock); return ret; --- libgomp/testsuite/libgomp.c/target-12.c.jj 2015-10-14 10:24:10.000000000 +0200 +++ libgomp/testsuite/libgomp.c/target-12.c 2015-11-05 08:57:43.910783553 +0100 @@ -41,7 +41,7 @@ main () if (omp_target_is_present (q, d) != 1 || omp_target_is_present (&q[32], d) != 1 - || omp_target_is_present (&q[128], d) != 1) + || omp_target_is_present (&q[127], d) != 1) abort (); if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0, --- libgomp/testsuite/libgomp.c/target-17.c.jj 2015-10-14 10:24:10.000000000 +0200 +++ libgomp/testsuite/libgomp.c/target-17.c 2015-11-04 17:20:39.441143671 +0100 @@ -37,58 +37,6 @@ foo (int n) } if (err) abort (); - int on = n; - #pragma omp target firstprivate (n) map(tofrom: n) - { - n++; - } - if (on != n) - abort (); - #pragma omp target map(tofrom: n) private (n) - { - n = 25; - } - if (on != n) - abort (); - for (i = 0; i < n; i++) - a[i] += i; - #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i) - { - err = 0; - for (i = 0; i < n; i++) - if (a[i] != 8 * i) - err = 1; - } - if (err) - abort (); - for (i = 0; i < n; i++) - a[i] += i; - #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i) - { - err = 0; - for (i = 0; i < n; i++) - if (a[i] != 9 * i) - err = 1; - } - if (err) - abort (); - for (i = 0; i < n; i++) - a[i] += i; - #pragma omp target map(tofrom:a) map(from:err) private(a, i) - { - err = 0; - for (i = 0; i < n; i++) - a[i] = 7; - #pragma omp parallel for reduction(|:err) - for (i = 0; i < n; i++) - if (a[i] != 7) - err |= 1; - } - if (err) - abort (); - for (i = 0; i < n; i++) - if (a[i] != 10 * i) - abort (); } int --- libgomp/testsuite/libgomp.c/target-19.c.jj 2015-10-14 10:24:10.000000000 +0200 +++ libgomp/testsuite/libgomp.c/target-19.c 2015-11-05 10:07:56.214421909 +0100 @@ -1,21 +1,29 @@ extern void abort (void); -void +__attribute__((noinline, noclone)) void foo (int *p, int *q, int *r, int n, int m) { int i, err, *s = r; + int sep = 1; + #pragma omp target map(to:sep) + sep = 0; #pragma omp target data map(to:p[0:8]) { /* For zero length array sections, p points to the start of - already mapped range, q to the end of it, and r does not point - to an mapped range. */ + already mapped range, q to the end of it (with nothing mapped + after it), and r does not point to an mapped range. */ #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s) { err = 0; for (i = 0; i < 8; i++) - if (p[i] != i + 1 || q[i - 8] != i + 1) + if (p[i] != i + 1) err = 1; - if (p + 8 != q || (r != (int *) 0 && r != s)) + if (sep) + { + if (q != (int *) 0 || r != (int *) 0) + err = 1; + } + else if (p + 8 != q || r != s) err = 1; } if (err) @@ -25,9 +33,14 @@ foo (int *p, int *q, int *r, int n, int { err = 0; for (i = 0; i < 8; i++) - if (p[i] != i + 1 || q[i - 8] != i + 1) + if (p[i] != i + 1) err = 1; - if (p + 8 != q || (r != (int *) 0 && r != s)) + if (sep) + { + if (q != (int *) 0 || r != (int *) 0) + err = 1; + } + else if (p + 8 != q || r != s) err = 1; } if (err) @@ -38,9 +51,14 @@ foo (int *p, int *q, int *r, int n, int { err = 0; for (i = 0; i < 8; i++) - if (p[i] != i + 1 || q[i - 8] != i + 1) + if (p[i] != i + 1) err = 1; - if (p + 8 != q || (r != (int *) 0 && r != s)) + if (sep) + { + if (q != (int *) 0 || r != (int *) 0) + err = 1; + } + else if (p + 8 != q || r != s) err = 1; } if (err) @@ -69,7 +87,14 @@ foo (int *p, int *q, int *r, int n, int for (i = 0; i < 8; i++) if (p[i] != i + 1) err = 1; - if (q[0] != 9 || r != q + 1) + if (q[0] != 9) + err = 1; + else if (sep) + { + if (r != (int *) 0) + err = 1; + } + else if (r != q + 1) err = 1; } if (err) @@ -81,7 +106,14 @@ foo (int *p, int *q, int *r, int n, int for (i = 0; i < 8; i++) if (p[i] != i + 1) err = 1; - if (q[0] != 9 || r != q + 1) + if (q[0] != 9) + err = 1; + else if (sep) + { + if (r != (int *) 0) + err = 1; + } + else if (r != q + 1) err = 1; } if (err) @@ -94,7 +126,14 @@ foo (int *p, int *q, int *r, int n, int for (i = 0; i < 8; i++) if (p[i] != i + 1) err = 1; - if (q[0] != 9 || r != q + 1) + if (q[0] != 9) + err = 1; + else if (sep) + { + if (r != (int *) 0) + err = 1; + } + else if (r != q + 1) err = 1; } if (err) --- libgomp/testsuite/libgomp.c/target-29.c.jj 2015-11-04 16:54:24.544542125 +0100 +++ libgomp/testsuite/libgomp.c/target-29.c 2015-11-04 18:08:41.861051720 +0100 @@ -0,0 +1,112 @@ +#include <omp.h> +#include <stdlib.h> + +struct S { char p[64]; int a; int b[2]; long c[4]; int *d; char q[64]; }; + +__attribute__((noinline, noclone)) void +foo (struct S s) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int sep = 1; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + int err; + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err) + { + err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; + err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; + s.a = 35; s.b[0] = 36; s.b[1] = 37; + s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42; + sep = 0; + } + if (err) abort (); + err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37; + err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42; + if (err) abort (); + s.a = 50; s.b[0] = 49; s.b[1] = 48; + s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43; + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d))) + abort (); + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) + { + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d)) + abort (); + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err) + { + err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; + err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; + s.a = 17; s.b[0] = 18; s.b[1] = 19; + s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24; + } + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3]) + } + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d))) + abort (); + if (err) abort (); + err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19; + err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24; + if (err) abort (); + s.a = 33; s.b[0] = 34; s.b[1] = 35; + s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d)) + abort (); + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err) + { + err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; + err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; + s.a = 49; s.b[0] = 48; s.b[1] = 47; + s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42; + } + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d)) + abort (); + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3]) + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d))) + abort (); + if (err) abort (); + err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47; + err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42; + if (err) abort (); +} + +int +main () +{ + int d[3] = { 18, 19, 20 }; + struct S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, {} }; + foo (s); + return 0; +} --- libgomp/testsuite/libgomp.c/target-30.c.jj 2015-11-04 18:17:50.878194390 +0100 +++ libgomp/testsuite/libgomp.c/target-30.c 2015-11-04 18:17:45.914265082 +0100 @@ -0,0 +1,24 @@ +extern void abort (void); + +#pragma omp declare target +int v = 6; +#pragma omp end declare target + +int +main () +{ + #pragma omp target /* predetermined map(tofrom: v) */ + v++; + #pragma omp target update from (v) + if (v != 7) + abort (); + #pragma omp parallel private (v) num_threads (1) + { + #pragma omp target /* predetermined firstprivate(v) */ + v++; + } + #pragma omp target update from (v) + if (v != 7) + abort (); + return 0; +} --- libgomp/testsuite/libgomp.c++/target-14.C.jj 2015-11-03 10:13:00.620573442 +0100 +++ libgomp/testsuite/libgomp.c++/target-14.C 2015-11-03 10:13:00.620573442 +0100 @@ -0,0 +1,110 @@ +extern "C" void abort (); +int x; + +__attribute__((noinline, noclone)) void +foo (int &a, int (&b)[10], short &c, long (&d)[5], int n) +{ + int err; + int &t = x; + int y[n + 1]; + int (&z)[n + 1] = y; + for (int i = 0; i < n + 1; i++) + z[i] = i + 27; + #pragma omp target enter data map (to: z, c) map (alloc: b, t) + #pragma omp target update to (b, t) + #pragma omp target map (tofrom: a, d) map (from: b, c) map (alloc: t, z) map (from: err) + { + err = a++ != 7; + for (int i = 0; i < 10; i++) + { + err |= b[i] != 10 - i; + b[i] = i - 16; + if (i >= 6) continue; + err |= z[i] != i + 27; + z[i] = 2 * i + 9; + if (i == 5) continue; + err |= d[i] != 12L + i; + d[i] = i + 7; + } + err |= c != 25; + c = 142; + err |= t != 8; + t = 19; + } + if (err) abort (); + #pragma omp target update from (z, c) + #pragma omp target exit data map (from: b, t) map (release: z, c) + if (a != 8 || c != 142 || t != 19) + abort (); + a = 29; + c = 149; + t = 15; + for (int i = 0; i < 10; i++) + { + if (b[i] != i - 16) abort (); + b[i] = i ^ 1; + if (i >= 6) continue; + if (z[i] != 2 * i + 9) abort (); + z[i]++; + if (i == 5) continue; + if (d[i] != i + 7) abort (); + d[i] = 7 - i; + } + #pragma omp target defaultmap(tofrom: scalar) + { + err = a++ != 29; + for (int i = 0; i < 10; i++) + { + err |= b[i] != i ^ 1; + b[i] = i + 5; + if (i >= 6) continue; + err |= z[i] != 2 * i + 10; + z[i] = 9 - 3 * i; + if (i == 5) continue; + err |= d[i] != 7L - i; + d[i] = i; + } + err |= c != 149; + c = -2; + err |= t != 15; + t = 155; + } + if (err || a != 30 || c != -2 || t != 155) + abort (); + for (int i = 0; i < 10; i++) + { + if (b[i] != i + 5) abort (); + if (i >= 6) continue; + if (z[i] != 9 - 3 * i) abort (); + z[i]++; + if (i == 5) continue; + if (d[i] != i) abort (); + } + #pragma omp target data map (alloc: z) + { + #pragma omp target update to (z) + #pragma omp target map(from: err) + { + err = 0; + for (int i = 0; i < 6; i++) + if (z[i] != 10 - 3 * i) err = 1; + else z[i] = i; + } + if (err) abort (); + #pragma omp target update from (z) + } + for (int i = 0; i < 6; i++) + if (z[i] != i) + abort (); +} + +int +main () +{ + int a = 7; + int b[10] = { 10, 9, 8, 7, 6, 5, 4, 3, 2, 1 }; + short c = 25; + long d[5] = { 12, 13, 14, 15, 16 }; + x = 8; + foo (a, b, c, d, 5); +} --- libgomp/testsuite/libgomp.c++/target-15.C.jj 2015-11-04 16:39:37.472162348 +0100 +++ libgomp/testsuite/libgomp.c++/target-15.C 2015-11-04 17:59:21.475097239 +0100 @@ -0,0 +1,168 @@ +#include <omp.h> +#include <stdlib.h> + +struct S { char p[64]; int a; int b[2]; long c[4]; int *d; unsigned char &e; char (&f)[2]; short (&g)[4]; int *&h; char q[64]; }; + +__attribute__((noinline, noclone)) void +foo (S s) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int sep = 1; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + int err; + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err) + { + err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; + err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; + err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26; + err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33; + s.a = 35; s.b[0] = 36; s.b[1] = 37; + s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42; + s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47; + s.h[2] = 48; s.h[3] = 49; s.h[4] = 50; + sep = 0; + } + if (err) abort (); + err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37; + err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42; + err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47; + err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50; + if (err) abort (); + s.a = 50; s.b[0] = 49; s.b[1] = 48; + s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43; + s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38; + s.h[2] = 37; s.h[3] = 36; s.h[4] = 35; + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + { + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + { + err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; + err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; + err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38; + err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35; + s.a = 17; s.b[0] = 18; s.b[1] = 19; + s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24; + s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29; + s.h[2] = 30; s.h[3] = 31; s.h[4] = 32; + } + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + } + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + if (err) abort (); + err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19; + err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24; + err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29; + err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32; + if (err) abort (); + s.a = 33; s.b[0] = 34; s.b[1] = 35; + s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; + s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45; + s.h[2] = 46; s.h[3] = 47; s.h[4] = 48; + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + { + err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; + err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; + err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45; + err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48; + s.a = 49; s.b[0] = 48; s.b[1] = 47; + s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42; + s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37; + s.h[2] = 36; s.h[3] = 35; s.h[4] = 34; + } + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + if (err) abort (); + err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47; + err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42; + err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37; + err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34; + if (err) abort (); +} + +int +main () +{ + int d[3] = { 18, 19, 20 }; + unsigned char e = 21; + char f[2] = { 22, 23 }; + short g[4] = { 24, 25, 26, 27 }; + int hb[7] = { 28, 29, 30, 31, 32, 33, 34 }; + int *h = hb + 1; + S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} }; + foo (s); +} --- libgomp/testsuite/libgomp.c++/target-16.C.jj 2015-11-05 09:55:59.081706150 +0100 +++ libgomp/testsuite/libgomp.c++/target-16.C 2015-11-05 09:58:21.448664482 +0100 @@ -0,0 +1,170 @@ +#include <omp.h> +#include <stdlib.h> + +template <typename C, typename I, typename L, typename UC, typename SH> +struct S { C p[64]; I a; I b[2]; L c[4]; I *d; UC &e; C (&f)[2]; SH (&g)[4]; I *&h; C q[64]; }; + +template <typename C, typename I, typename L, typename UC, typename SH> +__attribute__((noinline, noclone)) void +foo (S<C, I, L, UC, SH> s) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int sep = 1; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + int err; + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err) + { + err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; + err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; + err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26; + err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33; + s.a = 35; s.b[0] = 36; s.b[1] = 37; + s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42; + s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47; + s.h[2] = 48; s.h[3] = 49; s.h[4] = 50; + sep = 0; + } + if (err) abort (); + err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37; + err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42; + err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47; + err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50; + if (err) abort (); + s.a = 50; s.b[0] = 49; s.b[1] = 48; + s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43; + s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38; + s.h[2] = 37; s.h[3] = 36; s.h[4] = 35; + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + { + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + { + err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; + err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; + err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38; + err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35; + s.a = 17; s.b[0] = 18; s.b[1] = 19; + s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24; + s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29; + s.h[2] = 30; s.h[3] = 31; s.h[4] = 32; + } + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + } + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + if (err) abort (); + err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19; + err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24; + err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29; + err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32; + if (err) abort (); + s.a = 33; s.b[0] = 34; s.b[1] = 35; + s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; + s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45; + s.h[2] = 46; s.h[3] = 47; s.h[4] = 48; + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + { + err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; + err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; + err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45; + err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48; + s.a = 49; s.b[0] = 48; s.b[1] = 47; + s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42; + s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37; + s.h[2] = 36; s.h[3] = 35; s.h[4] = 34; + } + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + if (err) abort (); + err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47; + err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42; + err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37; + err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34; + if (err) abort (); +} + +int +main () +{ + int d[3] = { 18, 19, 20 }; + unsigned char e = 21; + char f[2] = { 22, 23 }; + short g[4] = { 24, 25, 26, 27 }; + int hb[7] = { 28, 29, 30, 31, 32, 33, 34 }; + int *h = hb + 1; + S<char, int, long, unsigned char, short> s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} }; + foo (s); +} --- libgomp/testsuite/libgomp.c++/target-17.C.jj 2015-11-05 09:59:26.662729254 +0100 +++ libgomp/testsuite/libgomp.c++/target-17.C 2015-11-05 10:05:17.628696101 +0100 @@ -0,0 +1,173 @@ +#include <omp.h> +#include <stdlib.h> + +template <typename C, typename I, typename L, typename UCR, typename CAR, typename SH, typename IPR> +struct S { C p[64]; I a; I b[2]; L c[4]; I *d; UCR e; CAR f; SH g; IPR h; C q[64]; }; + +template <typename C, typename I, typename L, typename UCR, typename CAR, typename SH, typename IPR> +__attribute__((noinline, noclone)) void +foo (S<C, I, L, UCR, CAR, SH, IPR> s) +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int sep = 1; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + int err; + #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err) + { + err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13; + err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20; + err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26; + err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33; + s.a = 35; s.b[0] = 36; s.b[1] = 37; + s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42; + s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47; + s.h[2] = 48; s.h[3] = 49; s.h[4] = 50; + sep = 0; + } + if (err) abort (); + err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37; + err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42; + err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47; + err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50; + if (err) abort (); + s.a = 50; s.b[0] = 49; s.b[1] = 48; + s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43; + s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38; + s.h[2] = 37; s.h[3] = 36; s.h[4] = 35; + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + { + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + { + err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48; + err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43; + err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38; + err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35; + s.a = 17; s.b[0] = 18; s.b[1] = 19; + s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24; + s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29; + s.h[2] = 30; s.h[3] = 31; s.h[4] = 32; + } + #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + } + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + if (err) abort (); + err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19; + err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24; + err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29; + err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32; + if (err) abort (); + s.a = 33; s.b[0] = 34; s.b[1] = 35; + s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40; + s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45; + s.h[2] = 46; s.h[3] = 47; s.h[4] = 48; + #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err) + { + err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35; + err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40; + err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45; + err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48; + s.a = 49; s.b[0] = 48; s.b[1] = 47; + s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42; + s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37; + s.h[2] = 36; s.h[3] = 35; s.h[4] = 34; + } + #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (!omp_target_is_present (&s.a, d) + || !omp_target_is_present (s.b, d) + || !omp_target_is_present (&s.c[1], d) + || !omp_target_is_present (s.d, d) + || !omp_target_is_present (&s.d[-2], d) + || !omp_target_is_present (&s.e, d) + || !omp_target_is_present (s.f, d) + || !omp_target_is_present (&s.g[1], d) + || !omp_target_is_present (&s.h, d) + || !omp_target_is_present (&s.h[2], d)) + abort (); + #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) + if (sep + && (omp_target_is_present (&s.a, d) + || omp_target_is_present (s.b, d) + || omp_target_is_present (&s.c[1], d) + || omp_target_is_present (s.d, d) + || omp_target_is_present (&s.d[-2], d) + || omp_target_is_present (&s.e, d) + || omp_target_is_present (s.f, d) + || omp_target_is_present (&s.g[1], d) + || omp_target_is_present (&s.h, d) + || omp_target_is_present (&s.h[2], d))) + abort (); + if (err) abort (); + err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47; + err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42; + err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37; + err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34; + if (err) abort (); +} + +int +main () +{ + int d[3] = { 18, 19, 20 }; + unsigned char e = 21; + char f[2] = { 22, 23 }; + short g[4] = { 24, 25, 26, 27 }; + int hb[7] = { 28, 29, 30, 31, 32, 33, 34 }; + int *h = hb + 1; + typedef char (&CAR)[2]; + typedef short (&SH)[4]; + S<char, int, long, unsigned char &, CAR, SH, int *&> s + = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} }; + foo (s); +} --- libgomp/testsuite/libgomp.c++/target-18.C.jj 2015-11-05 10:06:30.699648230 +0100 +++ libgomp/testsuite/libgomp.c++/target-18.C 2015-11-05 10:20:17.084797486 +0100 @@ -0,0 +1,167 @@ +extern "C" void abort (); + +__attribute__((noinline, noclone)) void +foo (int *&p, int *&q, int *&r, int n, int m) +{ + int i, err, *s = r; + int sep = 1; + #pragma omp target map(to:sep) + sep = 0; + #pragma omp target data map(to:p[0:8]) + { + /* For zero length array sections, p points to the start of + already mapped range, q to the end of it (with nothing mapped + after it), and r does not point to an mapped range. */ + #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (sep) + { + if (q != (int *) 0 || r != (int *) 0) + err = 1; + } + else if (p + 8 != q || r != s) + err = 1; + } + if (err) + abort (); + /* Implicit mapping of pointers behaves the same way. */ + #pragma omp target private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (sep) + { + if (q != (int *) 0 || r != (int *) 0) + err = 1; + } + else if (p + 8 != q || r != s) + err = 1; + } + if (err) + abort (); + /* And zero-length array sections, though not known at compile + time, behave the same. */ + #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (sep) + { + if (q != (int *) 0 || r != (int *) 0) + err = 1; + } + else if (p + 8 != q || r != s) + err = 1; + } + if (err) + abort (); + /* Non-zero length array sections, though not known at compile, + behave differently. */ + #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r[0] != 10) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:q[0:1]) + { + /* For zero length array sections, p points to the start of + already mapped range, q points to the start of another one, + and r to the end of the second one. */ + #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9) + err = 1; + else if (sep) + { + if (r != (int *) 0) + err = 1; + } + else if (r != q + 1) + err = 1; + } + if (err) + abort (); + /* Implicit mapping of pointers behaves the same way. */ + #pragma omp target private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9) + err = 1; + else if (sep) + { + if (r != (int *) 0) + err = 1; + } + else if (r != q + 1) + err = 1; + } + if (err) + abort (); + /* And zero-length array sections, though not known at compile + time, behave the same. */ + #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9) + err = 1; + else if (sep) + { + if (r != (int *) 0) + err = 1; + } + else if (r != q + 1) + err = 1; + } + if (err) + abort (); + /* Non-zero length array sections, though not known at compile, + behave differently. */ + #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r[0] != 10) + err = 1; + } + if (err) + abort (); + } + } +} + +int +main () +{ + int a[32], i; + for (i = 0; i < 32; i++) + a[i] = i; + int *p = a + 1, *q = a + 9, *r = a + 10; + foo (p, q, r, 0, 1); + return 0; +} --- libgomp/testsuite/libgomp.c++/target-19.C.jj 2015-11-05 10:18:48.964061178 +0100 +++ libgomp/testsuite/libgomp.c++/target-19.C 2015-11-05 10:30:11.145274934 +0100 @@ -0,0 +1,59 @@ +extern "C" void abort (); +struct S { char a[64]; int (&r)[2]; char b[64]; }; + +__attribute__((noinline, noclone)) void +foo (S s, int (&t)[3], int z) +{ + int err, sep = 1; + // Test that implicit mapping of reference to array does NOT + // behave like zero length array sections. s.r can't be used + // implicitly, as that means implicit mapping of the whole s + // and trying to dereference the references in there is unspecified. + #pragma omp target map(from: err) map(to: sep) + { + err = t[0] != 1 || t[1] != 2 || t[2] != 3; + sep = 0; + } + if (err) abort (); + // But explicit zero length array section mapping does. + #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0]) + { + if (sep) + err = s.r != (int *) 0 || t != (int *) 0; + else + err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; + } + if (err) abort (); + // Similarly zero length array section, but unknown at compile time. + #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z]) + { + if (sep) + err = s.r != (int *) 0 || t != (int *) 0; + else + err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; + } + if (err) abort (); + #pragma omp target enter data map (to: s.r, t) + // But when already mapped, it binds to existing mappings. + #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0]) + { + err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; + sep = 0; + } + if (err) abort (); + #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z]) + { + err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; + sep = 0; + } + if (err) abort (); +} + +int +main () +{ + int t[3] = { 1, 2, 3 }; + int r[2] = { 6, 7 }; + S s = { {}, r, {} }; + foo (s, t, 0); +} Jakub