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);
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)
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:
| (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
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);
&& 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);
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);
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)
|| (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),
}
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);
&& 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);
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;
}
&& 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);
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)
"%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);
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",
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;
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);
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:
| (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
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);
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 ();
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
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)
|| (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)
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);
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)
|| (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),
}
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);
&& 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);
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))
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;
}
&& 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);
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)))
"%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);
}
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)
{
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;
|| 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;
}
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;
}
&& ((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. */
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))
}
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)
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);
}
}
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);
|| 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))
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)
}
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)
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);
|| (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);
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
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
}
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);
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;
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:
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;
/* 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);
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;
}
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
#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 \
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 \
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 \
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 \
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)
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)
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)
;
}
--- /dev/null
+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" } */
+ ;
+}
#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
--- /dev/null
+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;
+}
--- /dev/null
+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;
+}
--- /dev/null
+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;
+}
--- /dev/null
+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;
+}
projects / platform / upstream / gcc.git / commitdiff