@@ -1221,7 +1221,8 @@ enum c_omp_region_type
C_ORT_DECLARE_SIMD = 1 << 2,
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
+ C_ORT_OMP_TARGET = C_ORT_OMP | C_ORT_TARGET,
+ C_ORT_ACC_TARGET = C_ORT_ACC | C_ORT_TARGET
};
extern tree c_finish_omp_master (location_t, tree);
@@ -1321,10 +1322,11 @@ public:
bool maybe_zero_length_array_section (tree);
tree expand_array_base (tree, vec<omp_addr_token *> &, tree, unsigned *,
- bool, bool);
+ c_omp_region_type, bool);
tree expand_component_selector (tree, vec<omp_addr_token *> &, tree,
- unsigned *, bool);
- tree expand_map_clause (tree, tree, vec<omp_addr_token *> &, bool);
+ unsigned *, c_omp_region_type);
+ tree expand_map_clause (tree, tree, vec<omp_addr_token *> &,
+ c_omp_region_type);
};
enum c_omp_directive_kind {
@@ -3369,7 +3369,8 @@ tree
c_omp_address_inspector::expand_array_base (tree c,
vec<omp_addr_token *> &addr_tokens,
tree expr, unsigned *idx,
- bool target, bool decl_p)
+ c_omp_region_type ort,
+ bool decl_p)
{
using namespace omp_addr_tokenizer;
location_t loc = OMP_CLAUSE_LOCATION (c);
@@ -3379,14 +3380,26 @@ c_omp_address_inspector::expand_array_base (tree c,
&& is_global_var (decl)
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl)));
+ bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
bool implicit_p = (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c));
bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
tree c2 = NULL_TREE, c3 = NULL_TREE;
unsigned consume_tokens = 2;
+ bool target = (ort & C_ORT_TARGET) != 0;
+ bool openmp = (ort & C_ORT_OMP) != 0;
gcc_assert (i == 0);
+ if (!openmp
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ {
+ *idx = ++i;
+ return c;
+ }
+
switch (addr_tokens[i + 1]->u.access_kind)
{
case ACCESS_DIRECT:
@@ -3396,11 +3409,19 @@ c_omp_address_inspector::expand_array_base (tree c,
case ACCESS_REF:
{
- /* Copy the referenced object. */
+ /* Copy the referenced object. Note that we do this even for !MAP_P
+ clauses. */
tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
OMP_CLAUSE_DECL (c) = obj;
OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
+ if (!map_p)
+ {
+ if (decl_p)
+ c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+ break;
+ }
+
/* If we have a reference to a pointer, avoid using
FIRSTPRIVATE_REFERENCE here in case the pointer is modified in the
offload region (we can only do that if the pointer does not point
@@ -3440,6 +3461,13 @@ c_omp_address_inspector::expand_array_base (tree c,
case ACCESS_INDEXED_REF_TO_ARRAY:
{
+ if (!map_p)
+ {
+ if (decl_p)
+ c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+ break;
+ }
+
tree virtual_origin
= convert_from_reference (addr_tokens[i + 1]->expr);
virtual_origin = build_fold_addr_expr (virtual_origin);
@@ -3466,6 +3494,13 @@ c_omp_address_inspector::expand_array_base (tree c,
case ACCESS_INDEXED_ARRAY:
{
+ if (!map_p)
+ {
+ if (decl_p)
+ c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+ break;
+ }
+
/* The code handling "firstprivatize_array_bases" in gimplify.cc is
relevant here. What do we need to create for arrays at this
stage? (This condition doesn't feel quite right. FIXME?) */
@@ -3500,6 +3535,13 @@ c_omp_address_inspector::expand_array_base (tree c,
case ACCESS_POINTER:
case ACCESS_POINTER_OFFSET:
{
+ if (!map_p)
+ {
+ if (decl_p)
+ c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+ break;
+ }
+
unsigned last_access = i + 1;
tree virtual_origin;
@@ -3523,7 +3565,11 @@ c_omp_address_inspector::expand_array_base (tree c,
addr_tokens[last_access]->expr);
tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
- if (decl_p && target && !chain_p && !declare_target_p)
+ /* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute
+ regions (e.g. "acc data" constructs). It'll be removed anyway in
+ gimplify.cc, but doing it this way maintains diagnostic
+ behaviour. */
+ if (decl_p && (target || !openmp) && !chain_p && !declare_target_p)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
else
{
@@ -3543,6 +3589,13 @@ c_omp_address_inspector::expand_array_base (tree c,
case ACCESS_REF_TO_POINTER:
case ACCESS_REF_TO_POINTER_OFFSET:
{
+ if (!map_p)
+ {
+ if (decl_p)
+ c_common_mark_addressable_vec (addr_tokens[i + 1]->expr);
+ break;
+ }
+
unsigned last_access = i + 1;
tree virtual_origin;
@@ -3617,7 +3670,7 @@ c_omp_address_inspector::expand_array_base (tree c,
i += consume_tokens;
*idx = i;
- if (target && chain_p)
+ if (target && chain_p && map_p)
return omp_expand_access_chain (c, expr, addr_tokens, idx);
else if (chain_p)
while (*idx < addr_tokens.length ()
@@ -3634,13 +3687,15 @@ c_omp_address_inspector::expand_component_selector (tree c,
vec<omp_addr_token *>
&addr_tokens,
tree expr, unsigned *idx,
- bool target)
+ c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
location_t loc = OMP_CLAUSE_LOCATION (c);
unsigned i = *idx;
tree c2 = NULL_TREE, c3 = NULL_TREE;
bool chain_p = omp_access_chain_p (addr_tokens, i + 1);
+ bool map_p = OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP;
+ bool target = (ort & C_ORT_TARGET) != 0;
switch (addr_tokens[i + 1]->u.access_kind)
{
@@ -3650,11 +3705,15 @@ c_omp_address_inspector::expand_component_selector (tree c,
case ACCESS_REF:
{
- /* Copy the referenced object. */
+ /* Copy the referenced object. Note that we also do this for !MAP_P
+ clauses. */
tree obj = convert_from_reference (addr_tokens[i + 1]->expr);
OMP_CLAUSE_DECL (c) = obj;
OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (obj));
+ if (!map_p)
+ break;
+
c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_DECL (c2) = addr_tokens[i + 1]->expr;
@@ -3664,6 +3723,9 @@ c_omp_address_inspector::expand_component_selector (tree c,
case ACCESS_INDEXED_REF_TO_ARRAY:
{
+ if (!map_p)
+ break;
+
tree virtual_origin
= convert_from_reference (addr_tokens[i + 1]->expr);
virtual_origin = build_fold_addr_expr (virtual_origin);
@@ -3685,6 +3747,9 @@ c_omp_address_inspector::expand_component_selector (tree c,
case ACCESS_POINTER:
case ACCESS_POINTER_OFFSET:
{
+ if (!map_p)
+ break;
+
tree virtual_origin
= fold_convert_loc (loc, ptrdiff_type_node,
addr_tokens[i + 1]->expr);
@@ -3704,6 +3769,9 @@ c_omp_address_inspector::expand_component_selector (tree c,
case ACCESS_REF_TO_POINTER:
case ACCESS_REF_TO_POINTER_OFFSET:
{
+ if (!map_p)
+ break;
+
tree ptr = convert_from_reference (addr_tokens[i + 1]->expr);
tree virtual_origin = fold_convert_loc (loc, ptrdiff_type_node,
ptr);
@@ -3749,7 +3817,7 @@ c_omp_address_inspector::expand_component_selector (tree c,
i += 2;
*idx = i;
- if (target && chain_p)
+ if (target && chain_p && map_p)
return omp_expand_access_chain (c, expr, addr_tokens, idx);
else if (chain_p)
while (*idx < addr_tokens.length ()
@@ -3765,7 +3833,7 @@ c_omp_address_inspector::expand_component_selector (tree c,
tree
c_omp_address_inspector::expand_map_clause (tree c, tree expr,
vec<omp_addr_token *> &addr_tokens,
- bool target)
+ c_omp_region_type ort)
{
using namespace omp_addr_tokenizer;
unsigned i, length = addr_tokens.length ();
@@ -3779,7 +3847,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
&& addr_tokens[i]->u.structure_base_kind == BASE_DECL
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
- c = expand_array_base (c, addr_tokens, expr, &i, target, true);
+ c = expand_array_base (c, addr_tokens, expr, &i, ort, true);
if (c == error_mark_node)
return error_mark_node;
}
@@ -3788,7 +3856,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
&& addr_tokens[i]->u.structure_base_kind == BASE_ARBITRARY_EXPR
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
- c = expand_array_base (c, addr_tokens, expr, &i, target, false);
+ c = expand_array_base (c, addr_tokens, expr, &i, ort, false);
if (c == error_mark_node)
return error_mark_node;
}
@@ -3824,7 +3892,7 @@ c_omp_address_inspector::expand_map_clause (tree c, tree expr,
&& addr_tokens[i]->type == COMPONENT_SELECTOR
&& addr_tokens[i + 1]->type == ACCESS_METHOD)
{
- c = expand_component_selector (c, addr_tokens, expr, &i, target);
+ c = expand_component_selector (c, addr_tokens, expr, &i, ort);
/* We used 'expr', so these must have been the last tokens. */
gcc_assert (i == length);
if (c == error_mark_node)
@@ -17135,7 +17135,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list)
static tree
c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
- const char *where, bool finish_p = true)
+ const char *where, bool finish_p = true,
+ bool target = false)
{
tree clauses = NULL;
bool first = true;
@@ -17336,7 +17337,8 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_parser_skip_to_pragma_eol (parser);
if (finish_p)
- return c_finish_omp_clauses (clauses, C_ORT_ACC);
+ return c_finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET
+ : C_ORT_ACC);
return clauses;
}
@@ -18061,12 +18063,13 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
mask |= OACC_LOOP_CLAUSE_MASK;
tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
- cclauses == NULL);
+ /*finish_p=*/cclauses == NULL,
+ /*target=*/is_parallel);
if (cclauses)
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel);
if (*cclauses)
- *cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC);
+ *cclauses = c_finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET);
if (clauses)
clauses = c_finish_omp_clauses (clauses, C_ORT_ACC);
}
@@ -18191,7 +18194,9 @@ c_parser_oacc_compute (location_t loc, c_parser *parser,
}
}
- tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
+ tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
+ /*finish_p=*/true,
+ /*target=*/true);
tree block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser, if_p));
@@ -13314,6 +13314,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
+ bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
@@ -13433,7 +13434,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
error_at (OMP_CLAUSE_LOCATION (c),
"expected single pointer in %qs clause",
- user_omp_clause_code_name (c, ort == C_ORT_ACC));
+ user_omp_clause_code_name (c, openacc));
return error_mark_node;
}
}
@@ -13881,9 +13882,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
c_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
- tree nc = ai.expand_map_clause (c, first, addr_tokens,
- (ort == C_ORT_OMP_TARGET
- || ort == C_ORT_ACC));
+ tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
if (nc != error_mark_node)
{
if (ai.maybe_zero_length_array_section (c))
@@ -14166,6 +14165,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bool allocate_seen = false;
bool implicit_moved = false;
bool target_in_reduction_seen = false;
+ bool openacc = (ort & C_ORT_ACC) != 0;
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -14181,7 +14181,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
- if (ort & C_ORT_ACC)
+ if (openacc)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
{
@@ -14575,8 +14575,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if ((ort == C_ORT_ACC
- && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+ else if ((openacc && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
|| (ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
|| (OMP_CLAUSE_CODE (c)
@@ -14599,7 +14598,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
- ort == C_ORT_ACC
+ openacc
? "%qD appears more than once in reduction clauses"
: "%qD appears more than once in data clauses",
t);
@@ -14622,7 +14621,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
- if (ort == C_ORT_ACC)
+ if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -14690,7 +14689,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&map_head, DECL_UID (t))
|| bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
- if (ort == C_ORT_ACC)
+ if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
@@ -15037,7 +15036,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion "
"clauses", rt);
- else if (ort == C_ORT_ACC)
+ else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data "
"clauses", rt);
@@ -15125,8 +15124,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
to be written. */
if (addr_tokens[0]->type == STRUCTURE_BASE
&& (bitmap_bit_p (&map_field_head, DECL_UID (t))
- || (ort != C_ORT_ACC
- && bitmap_bit_p (&map_head, DECL_UID (t)))))
+ || (!openacc && bitmap_bit_p (&map_head, DECL_UID (t)))))
goto skip_decl_checks;
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@@ -15192,7 +15190,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
- && ort == C_ORT_ACC)
+ && openacc)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
@@ -15207,7 +15205,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion clauses", t);
- else if (ort == C_ORT_ACC)
+ else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -15215,8 +15213,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in map clauses", t);
remove = true;
}
- else if (ort == C_ORT_ACC
- && bitmap_bit_p (&generic_head, DECL_UID (t)))
+ else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
@@ -15225,7 +15222,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&is_on_device_head, DECL_UID (t)))
{
- if (ort == C_ORT_ACC)
+ if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -15258,9 +15255,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c),
- addr_tokens,
- (ort == C_ORT_OMP_TARGET
- || ort == C_ORT_ACC));
+ addr_tokens, ort);
if (nc != error_mark_node)
c = nc;
}
@@ -15341,7 +15336,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
- && ort != C_ORT_ACC)
+ && !openacc)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is not a pointer",
@@ -40200,7 +40200,7 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list)
static tree
cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
const char *where, cp_token *pragma_tok,
- bool finish_p = true)
+ bool finish_p = true, bool target = false)
{
tree clauses = NULL;
bool first = true;
@@ -40410,7 +40410,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
if (finish_p)
- return finish_omp_clauses (clauses, C_ORT_ACC);
+ return finish_omp_clauses (clauses, target ? C_ORT_ACC_TARGET : C_ORT_ACC);
return clauses;
}
@@ -45028,7 +45028,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
bool found_in_scope = global_bindings_p ();
clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
- "#pragma acc declare", pragma_tok, true);
+ "#pragma acc declare", pragma_tok);
if (omp_find_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
@@ -45276,12 +45276,13 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
mask |= OACC_LOOP_CLAUSE_MASK;
tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
- cclauses == NULL);
+ /*finish_p=*/cclauses == NULL,
+ /*target=*/is_parallel);
if (cclauses)
{
clauses = c_oacc_split_loop_clauses (clauses, cclauses, is_parallel);
if (*cclauses)
- *cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC);
+ *cclauses = finish_omp_clauses (*cclauses, C_ORT_ACC_TARGET);
if (clauses)
clauses = finish_omp_clauses (clauses, C_ORT_ACC);
}
@@ -45406,7 +45407,9 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok,
}
}
- tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok);
+ tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok,
+ /*finish_p=*/true,
+ /*target=*/true);
tree block = begin_omp_parallel ();
unsigned int save = cp_parser_begin_omp_structured_block (parser);
@@ -19196,8 +19196,8 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
case OACC_KERNELS:
case OACC_PARALLEL:
case OACC_SERIAL:
- tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain,
- in_decl);
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC_TARGET, args,
+ complain, in_decl);
stmt = begin_omp_parallel ();
RECUR (OMP_BODY (t));
finish_omp_construct (TREE_CODE (t), stmt, tmp);
@@ -5139,6 +5139,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
enum c_omp_region_type ort)
{
tree ret, low_bound, length, type;
+ bool openacc = (ort & C_ORT_ACC) != 0;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
@@ -5255,7 +5256,7 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
{
error_at (OMP_CLAUSE_LOCATION (c),
"expected single pointer in %qs clause",
- user_omp_clause_code_name (c, ort == C_ORT_ACC));
+ user_omp_clause_code_name (c, openacc));
return error_mark_node;
}
}
@@ -5744,9 +5745,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
cp_omp_address_inspector ai (OMP_CLAUSE_LOCATION (c), t);
- tree nc = ai.expand_map_clause (c, first, addr_tokens,
- (ort == C_ORT_OMP_TARGET
- || ort == C_ORT_ACC));
+ tree nc = ai.expand_map_clause (c, first, addr_tokens, ort);
if (nc != error_mark_node)
{
using namespace omp_addr_tokenizer;
@@ -6693,6 +6692,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_head oacc_reduction_head, is_on_device_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
+ bool openacc = (ort & C_ORT_ACC) != 0;
bool branch_seen = false;
bool copyprivate_seen = false;
bool ordered_seen = false;
@@ -6725,7 +6725,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
bitmap_initialize (&is_on_device_head, &bitmap_default_obstack);
- if (ort & C_ORT_ACC)
+ if (openacc)
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_ASYNC)
{
@@ -6974,7 +6974,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
check_dup_generic_t:
if (t == current_class_ptr
- && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
+ && ((ort != C_ORT_OMP_DECLARE_SIMD && !openacc)
|| (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
{
@@ -6999,7 +6999,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
- else if ((ort == C_ORT_ACC
+ else if ((openacc
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
|| (ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
@@ -7023,7 +7023,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
- ort == C_ORT_ACC
+ openacc
? "%qD appears more than once in reduction clauses"
: "%qD appears more than once in data clauses",
t);
@@ -7046,7 +7046,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& bitmap_bit_p (&map_head, DECL_UID (t)))
{
- if (ort == C_ORT_ACC)
+ if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -7108,7 +7108,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
- if (ort != C_ORT_ACC && t == current_class_ptr)
+ if (!openacc && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7147,7 +7147,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&map_head, DECL_UID (t))
|| bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
- if (ort == C_ORT_ACC)
+ if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
@@ -7168,7 +7168,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
- if (ort != C_ORT_ACC && t == current_class_ptr)
+ if (!openacc && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -8048,7 +8048,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion"
" clauses", rt);
- else if (ort == C_ORT_ACC)
+ else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data"
" clauses", rt);
@@ -8145,8 +8145,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
to be written. */
if (addr_tokens[0]->type == STRUCTURE_BASE
&& (bitmap_bit_p (&map_field_head, DECL_UID (t))
- || (ort != C_ORT_ACC
- && bitmap_bit_p (&map_head, DECL_UID (t)))))
+ || (!openacc && bitmap_bit_p (&map_head, DECL_UID (t)))))
goto skip_decl_checks;
if (!processing_template_decl && TREE_CODE (t) == FIELD_DECL)
@@ -8242,7 +8241,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
- && ort == C_ORT_ACC)
+ && openacc)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
@@ -8261,7 +8260,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion clauses", t);
- else if (ort == C_ORT_ACC)
+ else if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -8269,8 +8268,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in map clauses", t);
remove = true;
}
- else if (ort == C_ORT_ACC
- && bitmap_bit_p (&generic_head, DECL_UID (t)))
+ else if (openacc && bitmap_bit_p (&generic_head, DECL_UID (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
@@ -8279,7 +8277,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t))
|| bitmap_bit_p (&is_on_device_head, DECL_UID (t)))
{
- if (ort == C_ORT_ACC)
+ if (openacc)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -8303,7 +8301,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
skip_decl_checks:
- /* If we call omp_expand_map_clause in handle_omp_array_sections,
+ /* If we call ai.expand_map_clause in handle_omp_array_sections,
the containing loop (here) iterates through the new nodes
created by that expansion. Avoid expanding those again (just
by checking the node type). */
@@ -8311,8 +8309,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
&& !processing_template_decl
&& ort != C_ORT_DECLARE_SIMD
&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
- || ((OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_FIRSTPRIVATE_POINTER)
+ || (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
@@ -8321,9 +8318,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
grp_start_p = pc;
grp_sentinel = OMP_CLAUSE_CHAIN (c);
tree nc = ai.expand_map_clause (c, OMP_CLAUSE_DECL (c),
- addr_tokens,
- (ort == C_ORT_OMP_TARGET
- || ort == C_ORT_ACC));
+ addr_tokens, ort);
if (nc != error_mark_node)
c = nc;
}