Hi! This patch adds support for in_reduction clause on target construct, though for now only for synchronous targets (without nowait clause). The encountering thread in that case runs the target task and blocks until the target region ends, so it is implemented by remapping it before entering the target, initializing the private copy if not yet initialized for the current thread and then using the remapped addresses for the mapping addresses. For nowait combined with in_reduction the patch contains a hack where the nowait clause is ignored. To implement it correctly, I think we would need to create a new private variable for the in_reduction and initialize it before doing the async target and adjust the map addresses to that private variable and then pass a function pointer to the library routine with code where the callback would remap the address to the current threads private variable and use in_reduction combiner to combine the private variable we've created into the thread's copy. The library would then need to make sure that the routine is called in some thread participating in the parallel (and not in an unshackeled thread).
Bootstrapped/regtested on x86_64-linux and i686-linux (with no offloading) and regtested on x86_64-linux with offloading to nvptx-none, committed to trunk. 2021-06-24 Jakub Jelinek <ja...@redhat.com> gcc/ * tree.h (OMP_CLAUSE_MAP_IN_REDUCTION): Document meaning for OpenMP. * gimplify.c (gimplify_scan_omp_clauses): For OpenMP map clauses with OMP_CLAUSE_MAP_IN_REDUCTION flag partially defer gimplification of non-decl OMP_CLAUSE_DECL. For OMP_CLAUSE_IN_REDUCTION on OMP_TARGET user outer_ctx instead of ctx for placeholders and initializer/combiner gimplification. * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_MAP_IN_REDUCTION on target constructs. (lower_rec_input_clauses): Likewise. (lower_omp_target): Likewise. * omp-expand.c (expand_omp_target): Temporarily ignore nowait clause on target if in_reduction is present. gcc/c-family/ * c-common.h (enum c_omp_region_type): Add C_ORT_TARGET and C_ORT_OMP_TARGET. * c-omp.c (c_omp_split_clauses): For OMP_CLAUSE_IN_REDUCTION on combined target constructs also add map (always, tofrom:) clause. gcc/c/ * c-parser.c (omp_split_clauses): Pass C_ORT_OMP_TARGET instead of C_ORT_OMP for clauses on target construct. (OMP_TARGET_CLAUSE_MASK): Add in_reduction clause. (c_parser_omp_target): For non-combined target add map (always, tofrom:) clauses for OMP_CLAUSE_IN_REDUCTION. Pass C_ORT_OMP_TARGET to c_finish_omp_clauses. * c-typeck.c (handle_omp_array_sections): Adjust ort handling for addition of C_ORT_OMP_TARGET and simplify, mapping clauses are never present on C_ORT_*DECLARE_SIMD. (c_finish_omp_clauses): Likewise. Handle OMP_CLAUSE_IN_REDUCTION on C_ORT_OMP_TARGET, set OMP_CLAUSE_MAP_IN_REDUCTION on corresponding map clauses. gcc/cp/ * parser.c (cp_omp_split_clauses): Pass C_ORT_OMP_TARGET instead of C_ORT_OMP for clauses on target construct. (OMP_TARGET_CLAUSE_MASK): Add in_reduction clause. (cp_parser_omp_target): For non-combined target add map (always, tofrom:) clauses for OMP_CLAUSE_IN_REDUCTION. Pass C_ORT_OMP_TARGET to finish_omp_clauses. * semantics.c (handle_omp_array_sections_1): Adjust ort handling for addition of C_ORT_OMP_TARGET and simplify, mapping clauses are never present on C_ORT_*DECLARE_SIMD. (handle_omp_array_sections): Likewise. (finish_omp_clauses): Likewise. Handle OMP_CLAUSE_IN_REDUCTION on C_ORT_OMP_TARGET, set OMP_CLAUSE_MAP_IN_REDUCTION on corresponding map clauses. * pt.c (tsubst_expr): Pass C_ORT_OMP_TARGET instead of C_ORT_OMP for clauses on target construct. gcc/testsuite/ * c-c++-common/gomp/target-in-reduction-1.c: New test. * c-c++-common/gomp/clauses-1.c: Add in_reduction clauses on target or combined target constructs. libgomp/ * testsuite/libgomp.c-c++-common/target-in-reduction-1.c: New test. * testsuite/libgomp.c-c++-common/target-in-reduction-2.c: New test. * testsuite/libgomp.c++/target-in-reduction-1.C: New test. * testsuite/libgomp.c++/target-in-reduction-2.C: New test. --- gcc/tree.h.jj 2021-06-23 09:57:34.624714820 +0200 +++ gcc/tree.h 2021-06-23 13:33:00.372434259 +0200 @@ -1651,7 +1651,8 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \ TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) /* Nonzero if this map clause is for an OpenACC compute construct's reduction - variable. */ + variable or OpenMP map clause mentioned also in in_reduction clause on the + same construct. */ #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \ TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) /* Nonzero on map clauses added implicitly for reduction clauses on combined --- gcc/gimplify.c.jj 2021-06-23 09:57:52.047480002 +0200 +++ gcc/gimplify.c 2021-06-23 13:33:00.378434179 +0200 @@ -9566,8 +9566,116 @@ gimplify_scan_omp_clauses (tree *list_p, OMP_CLAUSE_SET_MAP_KIND (c, k); } - if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR) + if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) + { + /* Don't gimplify *pd fully at this point, as the base + will need to be adjusted during omp lowering. */ + auto_vec<tree, 10> expr_stack; + tree *p = pd; + while (handled_component_p (*p) + || TREE_CODE (*p) == INDIRECT_REF + || TREE_CODE (*p) == ADDR_EXPR + || TREE_CODE (*p) == MEM_REF + || TREE_CODE (*p) == NON_LVALUE_EXPR) + { + expr_stack.safe_push (*p); + p = &TREE_OPERAND (*p, 0); + } + for (int i = expr_stack.length () - 1; i >= 0; i--) + { + tree t = expr_stack[i]; + if (TREE_CODE (t) == ARRAY_REF + || TREE_CODE (t) == ARRAY_RANGE_REF) + { + if (TREE_OPERAND (t, 2) == NULL_TREE) + { + tree low = unshare_expr (array_ref_low_bound (t)); + if (!is_gimple_min_invariant (low)) + { + TREE_OPERAND (t, 2) = low; + if (gimplify_expr (&TREE_OPERAND (t, 2), + pre_p, NULL, + is_gimple_reg, + fb_rvalue) == GS_ERROR) + remove = true; + } + } + else if (gimplify_expr (&TREE_OPERAND (t, 2), pre_p, + NULL, is_gimple_reg, + fb_rvalue) == GS_ERROR) + remove = true; + if (TREE_OPERAND (t, 3) == NULL_TREE) + { + tree elmt_size = array_ref_element_size (t); + if (!is_gimple_min_invariant (elmt_size)) + { + elmt_size = unshare_expr (elmt_size); + tree elmt_type + = TREE_TYPE (TREE_TYPE (TREE_OPERAND (t, + 0))); + tree factor + = size_int (TYPE_ALIGN_UNIT (elmt_type)); + elmt_size + = size_binop (EXACT_DIV_EXPR, elmt_size, + factor); + TREE_OPERAND (t, 3) = elmt_size; + if (gimplify_expr (&TREE_OPERAND (t, 3), + pre_p, NULL, + is_gimple_reg, + fb_rvalue) == GS_ERROR) + remove = true; + } + } + else if (gimplify_expr (&TREE_OPERAND (t, 3), pre_p, + NULL, is_gimple_reg, + fb_rvalue) == GS_ERROR) + remove = true; + } + else if (TREE_CODE (t) == COMPONENT_REF) + { + if (TREE_OPERAND (t, 2) == NULL_TREE) + { + tree offset = component_ref_field_offset (t); + if (!is_gimple_min_invariant (offset)) + { + offset = unshare_expr (offset); + tree field = TREE_OPERAND (t, 1); + tree factor + = size_int (DECL_OFFSET_ALIGN (field) + / BITS_PER_UNIT); + offset = size_binop (EXACT_DIV_EXPR, offset, + factor); + TREE_OPERAND (t, 2) = offset; + if (gimplify_expr (&TREE_OPERAND (t, 2), + pre_p, NULL, + is_gimple_reg, + fb_rvalue) == GS_ERROR) + remove = true; + } + } + else if (gimplify_expr (&TREE_OPERAND (t, 2), pre_p, + NULL, is_gimple_reg, + fb_rvalue) == GS_ERROR) + remove = true; + } + } + for (; expr_stack.length () > 0; ) + { + tree t = expr_stack.pop (); + + if (TREE_CODE (t) == ARRAY_REF + || TREE_CODE (t) == ARRAY_RANGE_REF) + { + if (!is_gimple_min_invariant (TREE_OPERAND (t, 1)) + && gimplify_expr (&TREE_OPERAND (t, 1), pre_p, + NULL, is_gimple_val, + fb_rvalue) == GS_ERROR) + remove = true; + } + } + } + else if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + fb_lvalue) == GS_ERROR) { remove = true; break; @@ -9764,17 +9872,21 @@ gimplify_scan_omp_clauses (tree *list_p, || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION) && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)) { - omp_add_variable (ctx, OMP_CLAUSE_REDUCTION_PLACEHOLDER (c), - GOVD_LOCAL | GOVD_SEEN); - if (OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c) + struct gimplify_omp_ctx *pctx + = code == OMP_TARGET ? outer_ctx : ctx; + if (pctx) + omp_add_variable (pctx, OMP_CLAUSE_REDUCTION_PLACEHOLDER (c), + GOVD_LOCAL | GOVD_SEEN); + if (pctx + && OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c) && walk_tree (&OMP_CLAUSE_REDUCTION_INIT (c), find_decl_expr, OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c), NULL) == NULL_TREE) - omp_add_variable (ctx, + omp_add_variable (pctx, OMP_CLAUSE_REDUCTION_DECL_PLACEHOLDER (c), GOVD_LOCAL | GOVD_SEEN); - gimplify_omp_ctxp = ctx; + gimplify_omp_ctxp = pctx; push_gimplify_context (); OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c) = NULL; --- gcc/omp-low.c.jj 2021-06-23 10:03:16.582109219 +0200 +++ gcc/omp-low.c 2021-06-23 17:15:15.113415451 +0200 @@ -1240,6 +1240,7 @@ scan_sharing_clauses (tree clauses, omp_ && ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && (OMP_CLAUSE_REDUCTION_INSCAN (c) || OMP_CLAUSE_REDUCTION_TASK (c))) + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION || is_task_ctx (ctx))) { /* For now. */ @@ -1254,6 +1255,29 @@ scan_sharing_clauses (tree clauses, omp_ if (TREE_CODE (t) == INDIRECT_REF || TREE_CODE (t) == ADDR_EXPR) t = TREE_OPERAND (t, 0); + if (is_omp_target (ctx->stmt)) + { + if (is_variable_sized (t)) + { + gcc_assert (DECL_HAS_VALUE_EXPR_P (t)); + t = DECL_VALUE_EXPR (t); + gcc_assert (TREE_CODE (t) == INDIRECT_REF); + t = TREE_OPERAND (t, 0); + gcc_assert (DECL_P (t)); + } + tree at = t; + if (ctx->outer) + scan_omp_op (&at, ctx->outer); + tree nt = omp_copy_decl_1 (at, ctx); + splay_tree_insert (ctx->field_map, + (splay_tree_key) &DECL_CONTEXT (t), + (splay_tree_value) nt); + if (at != t) + splay_tree_insert (ctx->field_map, + (splay_tree_key) &DECL_CONTEXT (at), + (splay_tree_value) nt); + break; + } install_var_local (t, ctx); if (is_taskreg_ctx (ctx) && (!is_global_var (maybe_lookup_decl_in_outer_ctx (t, ctx)) @@ -1280,6 +1304,21 @@ scan_sharing_clauses (tree clauses, omp_ } break; } + if (is_omp_target (ctx->stmt)) + { + tree at = decl; + if (ctx->outer) + scan_omp_op (&at, ctx->outer); + tree nt = omp_copy_decl_1 (at, ctx); + splay_tree_insert (ctx->field_map, + (splay_tree_key) &DECL_CONTEXT (decl), + (splay_tree_value) nt); + if (at != decl) + splay_tree_insert (ctx->field_map, + (splay_tree_key) &DECL_CONTEXT (at), + (splay_tree_value) nt); + break; + } if (is_task_ctx (ctx) || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && OMP_CLAUSE_REDUCTION_TASK (c) @@ -1546,7 +1585,8 @@ scan_sharing_clauses (tree clauses, omp_ else install_var_field (decl, true, 3, ctx); if (is_gimple_omp_offloaded (ctx->stmt) - && !OMP_CLAUSE_MAP_IN_REDUCTION (c)) + && !(is_gimple_omp_oacc (ctx->stmt) + && OMP_CLAUSE_MAP_IN_REDUCTION (c))) install_var_local (decl, ctx); } } @@ -1692,7 +1732,7 @@ scan_sharing_clauses (tree clauses, omp_ case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: decl = OMP_CLAUSE_DECL (c); - if (TREE_CODE (decl) != MEM_REF) + if (TREE_CODE (decl) != MEM_REF && !is_omp_target (ctx->stmt)) { if (is_variable_sized (decl)) install_var_local (decl, ctx); @@ -1844,8 +1884,11 @@ scan_sharing_clauses (tree clauses, omp_ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION) && OMP_CLAUSE_REDUCTION_PLACEHOLDER (c)) { - scan_omp (&OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c), ctx); - scan_omp (&OMP_CLAUSE_REDUCTION_GIMPLE_MERGE (c), ctx); + omp_context *rctx = ctx; + if (is_omp_target (ctx->stmt)) + rctx = ctx->outer; + scan_omp (&OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c), rctx); + scan_omp (&OMP_CLAUSE_REDUCTION_GIMPLE_MERGE (c), rctx); } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE && OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c)) @@ -4828,7 +4871,9 @@ lower_rec_input_clauses (tree clauses, g break; case OMP_CLAUSE_REDUCTION: case OMP_CLAUSE_IN_REDUCTION: - if (is_task_ctx (ctx) || OMP_CLAUSE_REDUCTION_TASK (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION + || is_task_ctx (ctx) + || OMP_CLAUSE_REDUCTION_TASK (c)) { task_reduction_p = true; if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) @@ -4958,7 +5003,12 @@ lower_rec_input_clauses (tree clauses, g } new_var = var; } - if (c_kind != OMP_CLAUSE_COPYIN) + if (c_kind == OMP_CLAUSE_IN_REDUCTION && is_omp_target (ctx->stmt)) + { + splay_tree_key key = (splay_tree_key) &DECL_CONTEXT (var); + new_var = (tree) splay_tree_lookup (ctx->field_map, key)->value; + } + else if (c_kind != OMP_CLAUSE_COPYIN) new_var = lookup_decl (var, ctx); if (c_kind == OMP_CLAUSE_SHARED || c_kind == OMP_CLAUSE_COPYIN) @@ -4980,7 +5030,10 @@ lower_rec_input_clauses (tree clauses, g if (TREE_CODE (orig_var) == POINTER_PLUS_EXPR) { tree b = TREE_OPERAND (orig_var, 1); - b = maybe_lookup_decl (b, ctx); + if (is_omp_target (ctx->stmt)) + b = NULL_TREE; + else + b = maybe_lookup_decl (b, ctx); if (b == NULL) { b = TREE_OPERAND (orig_var, 1); @@ -5006,6 +5059,8 @@ lower_rec_input_clauses (tree clauses, g || (TREE_CODE (TREE_TYPE (TREE_TYPE (out))) != POINTER_TYPE))) x = var; + else if (is_omp_target (ctx->stmt)) + x = out; else { bool by_ref = use_pointer_for_field (var, NULL); @@ -5049,7 +5104,11 @@ lower_rec_input_clauses (tree clauses, g const char *name = get_name (orig_var); if (pass != 3 && !TREE_CONSTANT (v)) { - tree t = maybe_lookup_decl (v, ctx); + tree t; + if (is_omp_target (ctx->stmt)) + t = NULL_TREE; + else + t = maybe_lookup_decl (v, ctx); if (t) v = t; else @@ -5100,7 +5159,11 @@ lower_rec_input_clauses (tree clauses, g TYPE_SIZE_UNIT (type)); else { - tree t = maybe_lookup_decl (v, ctx); + tree t; + if (is_omp_target (ctx->stmt)) + t = NULL_TREE; + else + t = maybe_lookup_decl (v, ctx); if (t) v = t; else @@ -5410,8 +5473,11 @@ lower_rec_input_clauses (tree clauses, g } else if (pass == 2) { - if (is_global_var (maybe_lookup_decl_in_outer_ctx (var, ctx))) + tree out = maybe_lookup_decl_in_outer_ctx (var, ctx); + if (is_global_var (out)) x = var; + else if (is_omp_target (ctx->stmt)) + x = out; else { bool by_ref = use_pointer_for_field (var, ctx); @@ -6345,7 +6411,27 @@ lower_rec_input_clauses (tree clauses, g if (OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c)) { tseq = OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c); - lower_omp (&tseq, ctx); + if (c_kind == OMP_CLAUSE_IN_REDUCTION + && is_omp_target (ctx->stmt)) + { + tree d = maybe_lookup_decl_in_outer_ctx (var, ctx); + tree oldv = NULL_TREE; + gcc_assert (d); + if (DECL_HAS_VALUE_EXPR_P (d)) + oldv = DECL_VALUE_EXPR (d); + SET_DECL_VALUE_EXPR (d, new_vard); + DECL_HAS_VALUE_EXPR_P (d) = 1; + lower_omp (&tseq, ctx); + if (oldv) + SET_DECL_VALUE_EXPR (d, oldv); + else + { + SET_DECL_VALUE_EXPR (d, NULL_TREE); + DECL_HAS_VALUE_EXPR_P (d) = 0; + } + } + else + lower_omp (&tseq, ctx); gimple_seq_add_seq (ilist, tseq); } OMP_CLAUSE_REDUCTION_GIMPLE_INIT (c) = NULL; @@ -12184,11 +12270,26 @@ lower_omp_target (gimple_stmt_iterator * location_t loc = gimple_location (stmt); bool offloaded, data_region; unsigned int map_cnt = 0; + tree in_reduction_clauses = NULL_TREE; offloaded = is_gimple_omp_offloaded (stmt); switch (gimple_omp_target_kind (stmt)) { case GF_OMP_TARGET_KIND_REGION: + tree *p, *q; + q = &in_reduction_clauses; + for (p = gimple_omp_target_clauses_ptr (stmt); *p; ) + if (OMP_CLAUSE_CODE (*p) == OMP_CLAUSE_IN_REDUCTION) + { + *q = *p; + q = &OMP_CLAUSE_CHAIN (*q); + *p = OMP_CLAUSE_CHAIN (*p); + } + else + p = &OMP_CLAUSE_CHAIN (*p); + *q = NULL_TREE; + *p = in_reduction_clauses; + /* FALLTHRU */ case GF_OMP_TARGET_KIND_UPDATE: case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_DATA: @@ -12217,12 +12318,17 @@ lower_omp_target (gimple_stmt_iterator * gimple_seq dep_ilist = NULL; gimple_seq dep_olist = NULL; - if (omp_find_clause (clauses, OMP_CLAUSE_DEPEND)) + bool has_depend = omp_find_clause (clauses, OMP_CLAUSE_DEPEND) != NULL_TREE; + if (has_depend || in_reduction_clauses) { push_gimplify_context (); dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK)); - lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt), - &dep_ilist, &dep_olist); + if (has_depend) + lower_depend_clauses (gimple_omp_target_clauses_ptr (stmt), + &dep_ilist, &dep_olist); + if (in_reduction_clauses) + lower_rec_input_clauses (in_reduction_clauses, &dep_ilist, &dep_olist, + ctx, NULL); } tgt_bind = NULL; @@ -12348,6 +12454,7 @@ lower_omp_target (gimple_stmt_iterator * /* Don't remap compute constructs' reduction variables, because the intermediate result must be local to each gang. */ if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && is_gimple_omp_oacc (ctx->stmt) && OMP_CLAUSE_MAP_IN_REDUCTION (c))) { x = build_receiver_ref (var, true, ctx); @@ -12565,16 +12672,46 @@ lower_omp_target (gimple_stmt_iterator * if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) { - gcc_checking_assert (OMP_CLAUSE_DECL (OMP_CLAUSE_CHAIN (c)) - == get_base_address (ovar)); nc = OMP_CLAUSE_CHAIN (c); + gcc_checking_assert (OMP_CLAUSE_DECL (nc) + == get_base_address (ovar)); ovar = OMP_CLAUSE_DECL (nc); } else { tree x = build_sender_ref (ovar, ctx); - tree v - = build_fold_addr_expr_with_type (ovar, ptr_type_node); + tree v = ovar; + if (in_reduction_clauses + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IN_REDUCTION (c)) + { + v = unshare_expr (v); + tree *p = &v; + while (handled_component_p (*p) + || TREE_CODE (*p) == INDIRECT_REF + || TREE_CODE (*p) == ADDR_EXPR + || TREE_CODE (*p) == MEM_REF + || TREE_CODE (*p) == NON_LVALUE_EXPR) + p = &TREE_OPERAND (*p, 0); + tree d = *p; + if (is_variable_sized (d)) + { + gcc_assert (DECL_HAS_VALUE_EXPR_P (d)); + d = DECL_VALUE_EXPR (d); + gcc_assert (TREE_CODE (d) == INDIRECT_REF); + d = TREE_OPERAND (d, 0); + gcc_assert (DECL_P (d)); + } + splay_tree_key key + = (splay_tree_key) &DECL_CONTEXT (d); + tree nd = (tree) splay_tree_lookup (ctx->field_map, + key)->value; + if (d == *p) + *p = nd; + else + *p = build_fold_indirect_ref (nd); + } + v = build_fold_addr_expr_with_type (v, ptr_type_node); gimplify_assign (x, v, &ilist); nc = NULL_TREE; } @@ -12601,19 +12738,45 @@ lower_omp_target (gimple_stmt_iterator * if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); + var = NULL_TREE; + if (nc) + { + if (in_reduction_clauses + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_IN_REDUCTION (c)) + { + tree d = ovar; + if (is_variable_sized (d)) + { + gcc_assert (DECL_HAS_VALUE_EXPR_P (d)); + d = DECL_VALUE_EXPR (d); + gcc_assert (TREE_CODE (d) == INDIRECT_REF); + d = TREE_OPERAND (d, 0); + gcc_assert (DECL_P (d)); + } + splay_tree_key key + = (splay_tree_key) &DECL_CONTEXT (d); + tree nd = (tree) splay_tree_lookup (ctx->field_map, + key)->value; + if (d == ovar) + var = nd; + else + var = build_fold_indirect_ref (nd); + } + else + var = lookup_decl_in_outer_ctx (ovar, ctx); + } if (nc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH) && is_omp_target (stmt)) { - var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (c, ctx); gimplify_assign (x, build_fold_addr_expr (var), &ilist); } else if (nc) { - var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP --- gcc/omp-expand.c.jj 2021-06-23 09:57:34.304719132 +0200 +++ gcc/omp-expand.c 2021-06-23 13:33:00.379434165 +0200 @@ -9615,6 +9615,10 @@ expand_omp_target (struct omp_region *re } c = omp_find_clause (clauses, OMP_CLAUSE_NOWAIT); + /* FIXME: in_reduction(...) nowait is unimplemented yet, pretend + nowait doesn't appear. */ + if (c && omp_find_clause (clauses, OMP_CLAUSE_IN_REDUCTION)) + c = NULL; if (c) flags_i |= GOMP_TARGET_FLAG_NOWAIT; } --- gcc/c-family/c-common.h.jj 2021-06-23 09:57:33.857725157 +0200 +++ gcc/c-family/c-common.h 2021-06-23 13:33:00.375434219 +0200 @@ -1208,7 +1208,9 @@ enum c_omp_region_type C_ORT_OMP = 1 << 0, C_ORT_ACC = 1 << 1, C_ORT_DECLARE_SIMD = 1 << 2, - C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD + C_ORT_TARGET = 1 << 3, + C_ORT_OMP_DECLARE_SIMD = C_ORT_OMP | C_ORT_DECLARE_SIMD, + C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET }; extern tree c_finish_omp_master (location_t, tree); --- gcc/c-family/c-omp.c.jj 2021-06-23 09:57:33.916724362 +0200 +++ gcc/c-family/c-omp.c 2021-06-23 13:33:00.376434206 +0200 @@ -2092,6 +2092,19 @@ c_omp_split_clauses (location_t loc, enu s = C_OMP_CLAUSE_SPLIT_TEAMS; break; case OMP_CLAUSE_IN_REDUCTION: + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)) != 0) + { + /* When on target, map(always, tofrom: item) is added as + well. For non-combined target it is added in the FEs. */ + c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses); + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_TOFROM); + OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; + cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = c; + s = C_OMP_CLAUSE_SPLIT_TARGET; + break; + } /* in_reduction on taskloop simd becomes reduction on the simd and keeps being in_reduction on taskloop. */ if (code == OMP_SIMD) --- gcc/c/c-parser.c.jj 2021-06-23 09:57:33.992723338 +0200 +++ gcc/c/c-parser.c 2021-06-23 13:33:00.374434233 +0200 @@ -18701,7 +18701,9 @@ omp_split_clauses (location_t loc, enum c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = c_finish_omp_clauses (cclauses[i], C_ORT_OMP); + cclauses[i] = c_finish_omp_clauses (cclauses[i], + i == C_OMP_CLAUSE_SPLIT_TARGET + ? C_ORT_OMP_TARGET : C_ORT_OMP); } /* OpenMP 5.0: @@ -20013,6 +20015,7 @@ c_parser_omp_target_exit_data (location_ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) static bool @@ -20179,7 +20182,18 @@ c_parser_omp_target (c_parser *parser, e OMP_TARGET_CLAUSES (stmt) = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, - "#pragma omp target"); + "#pragma omp target", false); + for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) + { + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_DECL (c); + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_TOFROM); + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = nc; + } + OMP_TARGET_CLAUSES (stmt) + = c_finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); pc = &OMP_TARGET_CLAUSES (stmt); --- gcc/c/c-typeck.c.jj 2021-06-23 09:57:34.040722690 +0200 +++ gcc/c/c-typeck.c 2021-06-23 13:33:00.375434219 +0200 @@ -13648,32 +13648,29 @@ handle_omp_array_sections (tree c, enum && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) return false; gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR); - if (ort == C_ORT_OMP || ort == C_ORT_ACC) - switch (OMP_CLAUSE_MAP_KIND (c)) - { - case GOMP_MAP_ALLOC: - case GOMP_MAP_IF_PRESENT: - case GOMP_MAP_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_RELEASE: - case GOMP_MAP_DELETE: - case GOMP_MAP_FORCE_TO: - case GOMP_MAP_FORCE_FROM: - case GOMP_MAP_FORCE_TOFROM: - case GOMP_MAP_FORCE_PRESENT: - OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; - break; - default: - break; - } + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + break; + default: + break; + } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - if (ort != C_ORT_OMP && ort != C_ORT_ACC) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - else if (TREE_CODE (t) == COMPONENT_REF) + if (TREE_CODE (t) == COMPONENT_REF) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); @@ -13970,6 +13967,7 @@ c_finish_omp_clauses (tree clauses, enum int reduction_seen = 0; bool allocate_seen = false; bool implicit_moved = false; + bool target_in_reduction_seen = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -13981,7 +13979,7 @@ c_finish_omp_clauses (tree clauses, enum bitmap_initialize (&map_field_head, &bitmap_default_obstack); bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack); /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head - instead. */ + instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head. */ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); if (ort & C_ORT_ACC) @@ -14374,8 +14372,22 @@ c_finish_omp_clauses (tree clauses, enum || (ort == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR || (OMP_CLAUSE_CODE (c) - == OMP_CLAUSE_USE_DEVICE_ADDR)))) - { + == OMP_CLAUSE_USE_DEVICE_ADDR))) + || (ort == C_ORT_OMP_TARGET + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION + && (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data-sharing " + "clauses", t); + remove = true; + break; + } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) + target_in_reduction_seen = true; if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), @@ -14390,7 +14402,8 @@ c_finish_omp_clauses (tree clauses, enum } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) - || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) + || bitmap_bit_p (&lastprivate_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qE appears more than once in data clauses", t); @@ -14457,7 +14470,8 @@ c_finish_omp_clauses (tree clauses, enum && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) remove = true; else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qE appears more than once in data clauses", t); @@ -14861,7 +14875,7 @@ c_finish_omp_clauses (tree clauses, enum if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { if (bitmap_bit_p (&map_field_head, DECL_UID (t)) - || (ort == C_ORT_OMP + || (ort != C_ORT_ACC && bitmap_bit_p (&map_head, DECL_UID (t)))) break; } @@ -14918,7 +14932,8 @@ c_finish_omp_clauses (tree clauses, enum && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) { if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -14935,13 +14950,10 @@ c_finish_omp_clauses (tree clauses, enum remove = true; } else - { - bitmap_set_bit (&generic_head, DECL_UID (t)); - bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); - } + bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) - && (ort != C_ORT_OMP + && (ort == C_ORT_ACC || !bitmap_bit_p (&map_field_head, DECL_UID (t)))) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) @@ -14955,8 +14967,8 @@ c_finish_omp_clauses (tree clauses, enum "%qD appears more than once in map clauses", t); remove = true; } - else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - && ort == C_ORT_ACC) + else if (ort == C_ORT_ACC + && bitmap_bit_p (&generic_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -15050,7 +15062,7 @@ c_finish_omp_clauses (tree clauses, enum if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) { if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR - && ort == C_ORT_OMP) + && ort != C_ORT_ACC) { error_at (OMP_CLAUSE_LOCATION (c), "%qs variable is not a pointer", @@ -15335,7 +15347,10 @@ c_finish_omp_clauses (tree clauses, enum reduction_seen = -2; } - if (linear_variable_step_check || reduction_seen == -2 || allocate_seen) + if (linear_variable_step_check + || reduction_seen == -2 + || allocate_seen + || target_in_reduction_seen) for (pc = &clauses, c = clauses; c ; c = *pc) { bool remove = false; @@ -15383,6 +15398,20 @@ c_finish_omp_clauses (tree clauses, enum else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION && reduction_seen == -2) OMP_CLAUSE_REDUCTION_INSCAN (c) = 0; + if (target_in_reduction_seen + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + { + tree t = OMP_CLAUSE_DECL (c); + while (handled_component_p (t) + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ADDR_EXPR + || TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == NON_LVALUE_EXPR) + t = TREE_OPERAND (t, 0); + if (DECL_P (t) + && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + OMP_CLAUSE_MAP_IN_REDUCTION (c) = 1; + } if (remove) *pc = OMP_CLAUSE_CHAIN (c); --- gcc/cp/parser.c.jj 2021-06-23 09:57:34.070722286 +0200 +++ gcc/cp/parser.c 2021-06-23 13:33:00.368434313 +0200 @@ -40877,7 +40877,9 @@ cp_omp_split_clauses (location_t loc, en c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = finish_omp_clauses (cclauses[i], C_ORT_OMP); + cclauses[i] = finish_omp_clauses (cclauses[i], + i == C_OMP_CLAUSE_SPLIT_TARGET + ? C_ORT_OMP_TARGET : C_ORT_OMP); } /* OpenMP 5.0: @@ -42219,6 +42221,7 @@ cp_parser_omp_target_update (cp_parser * | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEFAULTMAP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ALLOCATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IN_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR)) static bool @@ -42381,7 +42384,18 @@ cp_parser_omp_target (cp_parser *parser, OMP_TARGET_CLAUSES (stmt) = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK, - "#pragma omp target", pragma_tok); + "#pragma omp target", pragma_tok, false); + for (tree c = OMP_TARGET_CLAUSES (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) + { + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = OMP_CLAUSE_DECL (c); + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_ALWAYS_TOFROM); + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = nc; + } + OMP_TARGET_CLAUSES (stmt) + = finish_omp_clauses (OMP_TARGET_CLAUSES (stmt), C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true); pc = &OMP_TARGET_CLAUSES (stmt); --- gcc/cp/semantics.c.jj 2021-06-23 09:57:34.199720547 +0200 +++ gcc/cp/semantics.c 2021-06-23 13:33:00.369434300 +0200 @@ -5042,7 +5042,7 @@ handle_omp_array_sections_1 (tree c, tre omp_clause_code_name[OMP_CLAUSE_CODE (c)]); return error_mark_node; } - else if (ort == C_ORT_OMP + else if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP && TREE_CODE (t) == PARM_DECL && DECL_ARTIFICIAL (t) && DECL_NAME (t) == this_identifier @@ -5069,7 +5069,7 @@ handle_omp_array_sections_1 (tree c, tre return ret; } - if (ort == C_ORT_OMP + if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TASK_REDUCTION) @@ -5571,33 +5571,30 @@ handle_omp_array_sections (tree c, enum || (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)) return false; - if (ort == C_ORT_OMP || ort == C_ORT_ACC) - switch (OMP_CLAUSE_MAP_KIND (c)) - { - case GOMP_MAP_ALLOC: - case GOMP_MAP_IF_PRESENT: - case GOMP_MAP_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_RELEASE: - case GOMP_MAP_DELETE: - case GOMP_MAP_FORCE_TO: - case GOMP_MAP_FORCE_FROM: - case GOMP_MAP_FORCE_TOFROM: - case GOMP_MAP_FORCE_PRESENT: - OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; - break; - default: - break; - } + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_IF_PRESENT: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + break; + default: + break; + } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC) - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - else if (TREE_CODE (t) == COMPONENT_REF) + if (TREE_CODE (t) == COMPONENT_REF) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH); else if (REFERENCE_REF_P (t) && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF) @@ -6592,6 +6589,7 @@ finish_omp_clauses (tree clauses, enum c tree detach_seen = NULL_TREE; bool mergeable_seen = false; bool implicit_moved = false; + bool target_in_reduction_seen = false; bitmap_obstack_initialize (NULL); bitmap_initialize (&generic_head, &bitmap_default_obstack); @@ -6603,7 +6601,7 @@ finish_omp_clauses (tree clauses, enum c bitmap_initialize (&map_field_head, &bitmap_default_obstack); bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack); /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head - instead. */ + instead and for ort == C_ORT_OMP_TARGET used as in_reduction_head. */ bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack); if (ort & C_ORT_ACC) @@ -6866,8 +6864,22 @@ finish_omp_clauses (tree clauses, enum c || (ort == C_ORT_OMP && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR || (OMP_CLAUSE_CODE (c) - == OMP_CLAUSE_USE_DEVICE_ADDR)))) - { + == OMP_CLAUSE_USE_DEVICE_ADDR))) + || (ort == C_ORT_OMP_TARGET + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION + && (bitmap_bit_p (&generic_head, DECL_UID (t)) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))) + { + error_at (OMP_CLAUSE_LOCATION (c), + "%qD appears more than once in data-sharing " + "clauses", t); + remove = true; + break; + } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IN_REDUCTION) + target_in_reduction_seen = true; if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), @@ -6882,7 +6894,8 @@ finish_omp_clauses (tree clauses, enum c } else if (bitmap_bit_p (&generic_head, DECL_UID (t)) || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) - || bitmap_bit_p (&lastprivate_head, DECL_UID (t))) + || bitmap_bit_p (&lastprivate_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -6982,7 +6995,8 @@ finish_omp_clauses (tree clauses, enum c && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) remove = true; else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -7795,13 +7809,10 @@ finish_omp_clauses (tree clauses, enum c t = TREE_OPERAND (t, 0); OMP_CLAUSE_DECL (c) = t; } - if ((ort == C_ORT_ACC || ort == C_ORT_OMP) - && TREE_CODE (t) == COMPONENT_REF + if (TREE_CODE (t) == COMPONENT_REF && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF) t = TREE_OPERAND (TREE_OPERAND (t, 0), 0); if (TREE_CODE (t) == COMPONENT_REF - && ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP - || ort == C_ORT_ACC) && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_) { if (type_dependent_expression_p (t)) @@ -7842,7 +7853,7 @@ finish_omp_clauses (tree clauses, enum c if (VAR_P (t) || TREE_CODE (t) == PARM_DECL) { if (bitmap_bit_p (&map_field_head, DECL_UID (t)) - || (ort == C_ORT_OMP + || (ort != C_ORT_ACC && bitmap_bit_p (&map_head, DECL_UID (t)))) goto handle_map_references; } @@ -7924,7 +7935,8 @@ finish_omp_clauses (tree clauses, enum c && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) { if (bitmap_bit_p (&generic_head, DECL_UID (t)) - || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) + || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) + || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -7941,10 +7953,7 @@ finish_omp_clauses (tree clauses, enum c remove = true; } else - { - bitmap_set_bit (&generic_head, DECL_UID (t)); - bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); - } + bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); } else if (bitmap_bit_p (&map_head, DECL_UID (t)) && !bitmap_bit_p (&map_field_head, DECL_UID (t))) @@ -7960,8 +7969,8 @@ finish_omp_clauses (tree clauses, enum c "%qD appears more than once in map clauses", t); remove = true; } - else if (bitmap_bit_p (&generic_head, DECL_UID (t)) - && ort == C_ORT_ACC) + else if (ort == C_ORT_ACC + && bitmap_bit_p (&generic_head, DECL_UID (t))) { error_at (OMP_CLAUSE_LOCATION (c), "%qD appears more than once in data clauses", t); @@ -8511,6 +8520,22 @@ finish_omp_clauses (tree clauses, enum c } pc = &OMP_CLAUSE_CHAIN (c); continue; + case OMP_CLAUSE_MAP: + if (target_in_reduction_seen && !processing_template_decl) + { + t = OMP_CLAUSE_DECL (c); + while (handled_component_p (t) + || TREE_CODE (t) == INDIRECT_REF + || TREE_CODE (t) == ADDR_EXPR + || TREE_CODE (t) == MEM_REF + || TREE_CODE (t) == NON_LVALUE_EXPR) + t = TREE_OPERAND (t, 0); + if (DECL_P (t) + && bitmap_bit_p (&oacc_reduction_head, DECL_UID (t))) + OMP_CLAUSE_MAP_IN_REDUCTION (c) = 1; + } + pc = &OMP_CLAUSE_CHAIN (c); + continue; case OMP_CLAUSE_NOWAIT: if (copyprivate_seen) { --- gcc/cp/pt.c.jj 2021-06-23 09:57:34.121721599 +0200 +++ gcc/cp/pt.c 2021-06-23 13:33:00.371434273 +0200 @@ -18880,9 +18880,12 @@ tsubst_expr (tree t, tree args, tsubst_f case OACC_DATA: case OMP_TARGET_DATA: case OMP_TARGET: - tmp = tsubst_omp_clauses (OMP_CLAUSES (t), (TREE_CODE (t) == OACC_DATA) - ? C_ORT_ACC : C_ORT_OMP, args, complain, - in_decl); + tmp = tsubst_omp_clauses (OMP_CLAUSES (t), + TREE_CODE (t) == OACC_DATA + ? C_ORT_ACC + : TREE_CODE (t) == OMP_TARGET + ? C_ORT_OMP_TARGET : C_ORT_OMP, + args, complain, in_decl); keep_next_level (true); stmt = begin_omp_structured_block (); --- gcc/testsuite/c-c++-common/gomp/target-in-reduction-1.c.jj 2021-06-23 13:33:00.371434273 +0200 +++ gcc/testsuite/c-c++-common/gomp/target-in-reduction-1.c 2021-06-23 13:33:00.371434273 +0200 @@ -0,0 +1,12 @@ +void +foo (int i, int j, int k) +{ + #pragma omp target in_reduction (+:i) private (i) /* { dg-error "'i' appears more than once in data-sharing clauses" } */ + ; + #pragma omp target private (i) in_reduction (+:i) /* { dg-error "'i' appears both in data and map clauses" } */ + ; + #pragma omp target in_reduction (+:i) firstprivate (i) /* { dg-error "'i' appears more than once in data-sharing clauses" } */ + ; /* { dg-error "'i' appears both in data and map clauses" "" { target *-*-* } .-1 } */ + #pragma omp target firstprivate (i) in_reduction (+:i) /* { dg-error "'i' appears both in data and map clauses" } */ + ; +} --- gcc/testsuite/c-c++-common/gomp/clauses-1.c.jj 2021-06-23 09:57:34.474716841 +0200 +++ gcc/testsuite/c-c++-common/gomp/clauses-1.c 2021-06-23 13:33:00.372434259 +0200 @@ -125,20 +125,20 @@ bar (int d, int m, int i1, int i2, int i #pragma omp target parallel \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ - nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) + nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) ; #pragma omp target parallel for \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ lastprivate (l) linear (ll:1) ordered schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp target parallel for \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) nowait depend(inout: dd[0]) order(concurrent) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp target parallel for simd \ @@ -146,18 +146,18 @@ bar (int d, int m, int i1, int i2, int i if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ lastprivate (l) linear (ll:1) schedule(static, 4) collapse(1) \ safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) order(concurrent) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp target teams \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) ; #pragma omp target teams distribute \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \ - collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) + collapse(1) dist_schedule(static, 16) nowait depend(inout: dd[0]) allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ; #pragma omp target teams distribute parallel for \ @@ -166,7 +166,7 @@ bar (int d, int m, int i1, int i2, int i collapse(1) dist_schedule(static, 16) \ if (parallel: i2) num_threads (nth) proc_bind(spread) \ lastprivate (l) schedule(static, 4) nowait depend(inout: dd[0]) order(concurrent) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp target teams distribute parallel for simd \ @@ -176,7 +176,7 @@ bar (int d, int m, int i1, int i2, int i if (parallel: i2) num_threads (nth) proc_bind(spread) \ lastprivate (l) schedule(static, 4) order(concurrent) \ safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) if (simd: i3) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp target teams distribute simd \ @@ -184,14 +184,14 @@ bar (int d, int m, int i1, int i2, int i shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \ collapse(1) dist_schedule(static, 16) order(concurrent) \ safelen(8) simdlen(4) aligned(q: 32) nowait depend(inout: dd[0]) nontemporal(ntm) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp target simd \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ safelen(8) simdlen(4) lastprivate (l) linear(ll: 1) aligned(q: 32) reduction(+:r) \ nowait depend(inout: dd[0]) nontemporal(ntm) if(simd:i3) order(concurrent) \ - allocate (omp_default_mem_alloc:f) + allocate (omp_default_mem_alloc:f) in_reduction(+:r2) for (int i = 0; i < 64; i++) ll++; #pragma omp taskgroup task_reduction(+:r2) allocate (r2) @@ -215,7 +215,7 @@ bar (int d, int m, int i1, int i2, int i order(concurrent) allocate (f) for (int i = 0; i < 64; i++) ll++; - #pragma omp target nowait depend(inout: dd[0]) + #pragma omp target nowait depend(inout: dd[0]) in_reduction(+:r2) #pragma omp teams distribute \ private(p) firstprivate (f) shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) \ collapse(1) dist_schedule(static, 16) allocate (omp_default_mem_alloc: f) @@ -349,28 +349,28 @@ bar (int d, int m, int i1, int i2, int i device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ nowait depend(inout: dd[0]) lastprivate (l) bind(parallel) order(concurrent) collapse(1) \ - allocate (omp_default_mem_alloc: f) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) for (l = 0; l < 64; ++l) ; #pragma omp target parallel loop \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ if (parallel: i2) default(shared) shared(s) reduction(+:r) num_threads (nth) proc_bind(spread) \ nowait depend(inout: dd[0]) lastprivate (l) order(concurrent) collapse(1) \ - allocate (omp_default_mem_alloc: f) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) for (l = 0; l < 64; ++l) ; #pragma omp target teams loop \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \ lastprivate (l) bind(teams) collapse(1) \ - allocate (omp_default_mem_alloc: f) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) for (l = 0; l < 64; ++l) ; #pragma omp target teams loop \ device(d) map (tofrom: m) if (target: i1) private (p) firstprivate (f) defaultmap(tofrom: scalar) is_device_ptr (idp) \ shared(s) default(shared) reduction(+:r) num_teams(nte) thread_limit(tl) nowait depend(inout: dd[0]) \ lastprivate (l) order(concurrent) collapse(1) \ - allocate (omp_default_mem_alloc: f) + allocate (omp_default_mem_alloc: f) in_reduction(+:r2) for (l = 0; l < 64; ++l) ; } --- libgomp/testsuite/libgomp.c-c++-common/target-in-reduction-1.c.jj 2021-06-23 13:33:00.379434165 +0200 +++ libgomp/testsuite/libgomp.c-c++-common/target-in-reduction-1.c 2021-06-23 13:33:00.379434165 +0200 @@ -0,0 +1,104 @@ +void +foo (int x, int *y, int n, int v) +{ + int z[3] = { 45, 46, 47 }; + int u[n], w[n], i; + for (i = 0; i < n; i++) + w[i] = u[i] = n + i; + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x++; + y[0] += 2; + y[1] += 3; + z[1] += 4; + u[0] += 5; + w[1] += 6; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x += 4; + y[0] += 5; + y[1] += 6; + z[2] += 7; + u[1] += 8; + w[2] += 7; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x += 9; + y[0] += 10; + y[1] += 11; + z[1] += 12; + u[2] += 13; + w[1] += 14; + } + } + if (x != 56 || y[0] != 60 || y[1] != 64) + __builtin_abort (); + if (z[0] != 45 || z[1] != 62 || z[2] != 54) + __builtin_abort (); + if (u[0] != 8 || u[1] != 12 || u[2] != 18) + __builtin_abort (); + if (w[0] != 3 || w[1] != 24 || w[2] != 12) + __builtin_abort (); +} + +void +bar (int x, int *y, int n, int v) +{ + int z[3] = { 45, 46, 47 }; + int u[n], w[n], i; + for (i = 0; i < n; i++) + w[i] = u[i] = n + i; + #pragma omp parallel master + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x++; + y[0] += 2; + y[1] += 3; + z[1] += 4; + u[0] += 5; + w[1] += 6; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x += 4; + y[0] += 5; + y[1] += 6; + z[2] += 7; + u[1] += 8; + w[2] += 7; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x += 9; + y[0] += 10; + y[1] += 11; + z[1] += 12; + u[2] += 13; + w[1] += 14; + } + } + if (x != 56 || y[0] != 77 || y[1] != 84) + __builtin_abort (); + if (z[0] != 45 || z[1] != 62 || z[2] != 54) + __builtin_abort (); + if (u[0] != 8 || u[1] != 12 || u[2] != 18) + __builtin_abort (); + if (w[0] != 3 || w[1] != 24 || w[2] != 12) + __builtin_abort (); +} + +int +main () +{ + int y[2] = { 43, 44 }; + #pragma omp parallel master + foo (42, y, 3, 2); + bar (42, y, 3, 2); + return 0; +} --- libgomp/testsuite/libgomp.c-c++-common/target-in-reduction-2.c.jj 2021-06-23 13:33:00.379434165 +0200 +++ libgomp/testsuite/libgomp.c-c++-common/target-in-reduction-2.c 2021-06-23 17:09:12.538281676 +0200 @@ -0,0 +1,173 @@ +struct S { int a, b, c[2]; }; +#pragma omp declare reduction (+: struct S : (omp_out.a += omp_in.a, omp_out.b += omp_in.b)) \ + initializer (omp_priv = { 0, 0, { 0, 0 } }) + +void +foo (struct S x, struct S *y, int n, int v) +{ + struct S z[3] = { { 45, 47, {} }, { 46, 48, {} }, { 47, 49, {} } }; + struct S u[n], w[n]; + int i; + for (i = 0; i < n; i++) + { + w[i].a = u[i].a = n + i; + w[i].b = u[i].b = n - i; + w[i].c[0] = u[i].c[0] = 0; + w[i].c[1] = u[i].c[1] = 0; + } + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x.a++; + x.b++; + y[0].a += 2; + y[0].b += 12; + y[1].a += 3; + y[1].b += 13; + z[1].a += 4; + z[1].b += 14; + u[0].a += 5; + u[0].b += 15; + w[1].a += 6; + w[1].b += 16; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) map(tofrom: x.a, x.b, x.c[:2]) + { + x.a += 4; + x.b += 14; + y[0].a += 5; + y[0].b += 15; + y[1].a += 6; + y[1].b += 16; + z[2].a += 7; + z[2].b += 17; + u[1].a += 8; + u[1].b += 18; + w[2].a += 7; + w[2].b += 17; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x.a += 9; + x.b += 19; + y[0].a += 10; + y[0].b += 20; + y[1].a += 11; + y[1].b += 21; + z[1].a += 12; + z[1].b += 22; + u[2].a += 13; + u[2].b += 23; + w[1].a += 14; + w[1].b += 24; + } + } + if (x.a != 56 || y[0].a != 60 || y[1].a != 64) + __builtin_abort (); + if (x.b != 86 || y[0].b != 100 || y[1].b != 104) + __builtin_abort (); + if (z[0].a != 45 || z[1].a != 62 || z[2].a != 54) + __builtin_abort (); + if (z[0].b != 47 || z[1].b != 84 || z[2].b != 66) + __builtin_abort (); + if (u[0].a != 8 || u[1].a != 12 || u[2].a != 18) + __builtin_abort (); + if (u[0].b != 18 || u[1].b != 20 || u[2].b != 24) + __builtin_abort (); + if (w[0].a != 3 || w[1].a != 24 || w[2].a != 12) + __builtin_abort (); + if (w[0].b != 3 || w[1].b != 42 || w[2].b != 18) + __builtin_abort (); +} + +void +bar (struct S x, struct S *y, int n, int v) +{ + struct S z[3] = { { 45, 47, {} }, { 46, 48, {} }, { 47, 49, {} } }; + struct S u[n], w[n]; + int i; + for (i = 0; i < n; i++) + { + w[i].a = u[i].a = n + i; + w[i].b = u[i].b = n - i; + w[i].c[0] = u[i].c[0] = 0; + w[i].c[1] = u[i].c[1] = 0; + } + #pragma omp parallel master + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x.a++; + x.b++; + y[0].a += 2; + y[0].b += 12; + y[1].a += 3; + y[1].b += 13; + z[1].a += 4; + z[1].b += 14; + u[0].a += 5; + u[0].b += 15; + w[1].a += 6; + w[1].b += 16; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) map(tofrom: x.a, x.b, x.c[:2]) + { + x.a += 4; + x.b += 14; + y[0].a += 5; + y[0].b += 15; + y[1].a += 6; + y[1].b += 16; + z[2].a += 7; + z[2].b += 17; + u[1].a += 8; + u[1].b += 18; + w[2].a += 7; + w[2].b += 17; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x.a += 9; + x.b += 19; + y[0].a += 10; + y[0].b += 20; + y[1].a += 11; + y[1].b += 21; + z[1].a += 12; + z[1].b += 22; + u[2].a += 13; + u[2].b += 23; + w[1].a += 14; + w[1].b += 24; + } + } + if (x.a != 56 || y[0].a != 77 || y[1].a != 84) + __builtin_abort (); + if (x.b != 86 || y[0].b != 147 || y[1].b != 154) + __builtin_abort (); + if (z[0].a != 45 || z[1].a != 62 || z[2].a != 54) + __builtin_abort (); + if (z[0].b != 47 || z[1].b != 84 || z[2].b != 66) + __builtin_abort (); + if (u[0].a != 8 || u[1].a != 12 || u[2].a != 18) + __builtin_abort (); + if (u[0].b != 18 || u[1].b != 20 || u[2].b != 24) + __builtin_abort (); + if (w[0].a != 3 || w[1].a != 24 || w[2].a != 12) + __builtin_abort (); + if (w[0].b != 3 || w[1].b != 42 || w[2].b != 18) + __builtin_abort (); +} + +int +main () +{ + struct S x = { 42, 52 }; + struct S y[2] = { { 43, 53 }, { 44, 54 } }; + #pragma omp parallel master + foo (x, y, 3, 2); + bar (x, y, 3, 2); + return 0; +} --- libgomp/testsuite/libgomp.c++/target-in-reduction-1.C.jj 2021-06-23 17:21:25.748440669 +0200 +++ libgomp/testsuite/libgomp.c++/target-in-reduction-1.C 2021-06-23 17:21:25.748440669 +0200 @@ -0,0 +1,113 @@ +void +foo (int &x, int *&y, int n, int v) +{ + int zu[3] = { 45, 46, 47 }; + int uu[n], wu[n], i; + int (&z)[3] = zu; + int (&u)[n] = uu; + int (&w)[n] = wu; + for (i = 0; i < n; i++) + w[i] = u[i] = n + i; + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x++; + y[0] += 2; + y[1] += 3; + z[1] += 4; + u[0] += 5; + w[1] += 6; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x += 4; + y[0] += 5; + y[1] += 6; + z[2] += 7; + u[1] += 8; + w[2] += 7; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x += 9; + y[0] += 10; + y[1] += 11; + z[1] += 12; + u[2] += 13; + w[1] += 14; + } + } + if (x != 56 || y[0] != 60 || y[1] != 64) + __builtin_abort (); + if (z[0] != 45 || z[1] != 62 || z[2] != 54) + __builtin_abort (); + if (u[0] != 8 || u[1] != 12 || u[2] != 18) + __builtin_abort (); + if (w[0] != 3 || w[1] != 24 || w[2] != 12) + __builtin_abort (); +} + +void +bar (int &x, int *&y, int n, int v) +{ + int zu[3] = { 45, 46, 47 }; + int uu[n], wu[n], i; + int (&z)[3] = zu; + int (&u)[n] = uu; + int (&w)[n] = wu; + for (i = 0; i < n; i++) + w[i] = u[i] = n + i; + #pragma omp parallel master + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x++; + y[0] += 2; + y[1] += 3; + z[1] += 4; + u[0] += 5; + w[1] += 6; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x += 4; + y[0] += 5; + y[1] += 6; + z[2] += 7; + u[1] += 8; + w[2] += 7; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x += 9; + y[0] += 10; + y[1] += 11; + z[1] += 12; + u[2] += 13; + w[1] += 14; + } + } + if (x != 56 || y[0] != 77 || y[1] != 84) + __builtin_abort (); + if (z[0] != 45 || z[1] != 62 || z[2] != 54) + __builtin_abort (); + if (u[0] != 8 || u[1] != 12 || u[2] != 18) + __builtin_abort (); + if (w[0] != 3 || w[1] != 24 || w[2] != 12) + __builtin_abort (); +} + +int +main () +{ + int x = 42; + int yu[2] = { 43, 44 }; + int *y = yu; + #pragma omp parallel master + foo (x, y, 3, 2); + x = 42; + bar (x, y, 3, 2); + return 0; +} --- libgomp/testsuite/libgomp.c++/target-in-reduction-2.C.jj 2021-06-23 17:21:25.748440669 +0200 +++ libgomp/testsuite/libgomp.c++/target-in-reduction-2.C 2021-06-23 18:46:25.078073385 +0200 @@ -0,0 +1,182 @@ +struct S { int a, b, c[2]; }; +#pragma omp declare reduction (+: S : (omp_out.a += omp_in.a, omp_out.b += omp_in.b)) \ + initializer (omp_priv = { 0, 0, { 0, 0 } }) + +void +foo (S &x, S *&y, int n, int v) +{ + S zu[3] = { { 45, 47, {} }, { 46, 48, {} }, { 47, 49, {} } }; + S uu[n], wu[n]; + S (&z)[3] = zu; + S (&u)[n] = uu; + S (&w)[n] = wu; + int i; + for (i = 0; i < n; i++) + { + w[i].a = u[i].a = n + i; + w[i].b = u[i].b = n - i; + w[i].c[0] = u[i].c[0] = 0; + w[i].c[1] = u[i].c[1] = 0; + } + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x.a++; + x.b++; + y[0].a += 2; + y[0].b += 12; + y[1].a += 3; + y[1].b += 13; + z[1].a += 4; + z[1].b += 14; + u[0].a += 5; + u[0].b += 15; + w[1].a += 6; + w[1].b += 16; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x.a += 4; + x.b += 14; + y[0].a += 5; + y[0].b += 15; + y[1].a += 6; + y[1].b += 16; + z[2].a += 7; + z[2].b += 17; + u[1].a += 8; + u[1].b += 18; + w[2].a += 7; + w[2].b += 17; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x.a += 9; + x.b += 19; + y[0].a += 10; + y[0].b += 20; + y[1].a += 11; + y[1].b += 21; + z[1].a += 12; + z[1].b += 22; + u[2].a += 13; + u[2].b += 23; + w[1].a += 14; + w[1].b += 24; + } + } + if (x.a != 56 || y[0].a != 60 || y[1].a != 64) + __builtin_abort (); + if (x.b != 86 || y[0].b != 100 || y[1].b != 104) + __builtin_abort (); + if (z[0].a != 45 || z[1].a != 62 || z[2].a != 54) + __builtin_abort (); + if (z[0].b != 47 || z[1].b != 84 || z[2].b != 66) + __builtin_abort (); + if (u[0].a != 8 || u[1].a != 12 || u[2].a != 18) + __builtin_abort (); + if (u[0].b != 18 || u[1].b != 20 || u[2].b != 24) + __builtin_abort (); + if (w[0].a != 3 || w[1].a != 24 || w[2].a != 12) + __builtin_abort (); + if (w[0].b != 3 || w[1].b != 42 || w[2].b != 18) + __builtin_abort (); +} + +void +bar (S &x, S *&y, int n, int v) +{ + S zu[3] = { { 45, 47, {} }, { 46, 48, {} }, { 47, 49, {} } }; + S uu[n], wu[n]; + S (&z)[3] = zu; + S (&u)[n] = uu; + S (&w)[n] = wu; + int i; + for (i = 0; i < n; i++) + { + w[i].a = u[i].a = n + i; + w[i].b = u[i].b = n - i; + w[i].c[0] = u[i].c[0] = 0; + w[i].c[1] = u[i].c[1] = 0; + } + #pragma omp parallel master + #pragma omp taskgroup task_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + #pragma omp task in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x.a++; + x.b++; + y[0].a += 2; + y[0].b += 12; + y[1].a += 3; + y[1].b += 13; + z[1].a += 4; + z[1].b += 14; + u[0].a += 5; + u[0].b += 15; + w[1].a += 6; + w[1].b += 16; + } + #pragma omp target in_reduction (+: x, y[:2], z[1:2], u, w[1:v]) + { + x.a += 4; + x.b += 14; + y[0].a += 5; + y[0].b += 15; + y[1].a += 6; + y[1].b += 16; + z[2].a += 7; + z[2].b += 17; + u[1].a += 8; + u[1].b += 18; + w[2].a += 7; + w[2].b += 17; + } + #pragma omp target in_reduction (+: x, y[:v], z[1:v], u, w[1:2]) + { + x.a += 9; + x.b += 19; + y[0].a += 10; + y[0].b += 20; + y[1].a += 11; + y[1].b += 21; + z[1].a += 12; + z[1].b += 22; + u[2].a += 13; + u[2].b += 23; + w[1].a += 14; + w[1].b += 24; + } + } + if (x.a != 56 || y[0].a != 77 || y[1].a != 84) + __builtin_abort (); + if (x.b != 86 || y[0].b != 147 || y[1].b != 154) + __builtin_abort (); + if (z[0].a != 45 || z[1].a != 62 || z[2].a != 54) + __builtin_abort (); + if (z[0].b != 47 || z[1].b != 84 || z[2].b != 66) + __builtin_abort (); + if (u[0].a != 8 || u[1].a != 12 || u[2].a != 18) + __builtin_abort (); + if (u[0].b != 18 || u[1].b != 20 || u[2].b != 24) + __builtin_abort (); + if (w[0].a != 3 || w[1].a != 24 || w[2].a != 12) + __builtin_abort (); + if (w[0].b != 3 || w[1].b != 42 || w[2].b != 18) + __builtin_abort (); +} + +int +main () +{ + S x = { 42, 52 }; + S yu[2] = { { 43, 53 }, { 44, 54 } }; + S *y = yu; + #pragma omp parallel master + foo (x, y, 3, 2); + x.a = 42; + x.b = 52; + bar (x, y, 3, 2); + return 0; +} Jakub