@@ -1245,7 +1245,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);
@@ -1345,10 +1346,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 {
@@ -3370,7 +3370,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);
@@ -3380,14 +3381,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:
@@ -3397,11 +3410,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
@@ -3441,6 +3462,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);
@@ -3467,6 +3495,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?) */
@@ -3501,6 +3536,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;
@@ -3524,7 +3566,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
{
@@ -3544,6 +3590,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;
@@ -3618,7 +3671,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 ()
@@ -3635,13 +3688,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)
{
@@ -3651,11 +3706,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;
@@ -3665,6 +3724,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);
@@ -3686,6 +3748,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);
@@ -3705,6 +3770,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);
@@ -3750,7 +3818,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 ()
@@ -3766,7 +3834,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 ();
@@ -3780,7 +3848,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;
}
@@ -3789,7 +3857,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;
}
@@ -3825,7 +3893,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)
@@ -17514,7 +17514,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;
@@ -17715,7 +17716,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;
}
@@ -18440,12 +18442,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);
}
@@ -18570,7 +18573,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));
@@ -13634,6 +13634,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))
@@ -13753,7 +13754,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;
}
}
@@ -14201,9 +14202,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))
@@ -14486,6 +14485,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);
@@ -14501,7 +14501,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)
{
@@ -14895,8 +14895,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)
@@ -14919,7 +14918,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);
@@ -14942,7 +14941,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
@@ -15010,7 +15009,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)
@@ -15357,7 +15356,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);
@@ -15445,8 +15444,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)
@@ -15512,7 +15510,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);
@@ -15527,7 +15525,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
@@ -15535,8 +15533,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);
@@ -15545,7 +15542,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
@@ -15578,9 +15575,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;
}
@@ -15661,7 +15656,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",
@@ -40764,7 +40764,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;
@@ -40974,7 +40974,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;
}
@@ -45587,7 +45587,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)
@@ -45835,12 +45835,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);
}
@@ -45965,7 +45966,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);
@@ -19258,8 +19258,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);
@@ -5167,6 +5167,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))
@@ -5283,7 +5284,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;
}
}
@@ -5772,9 +5773,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;
@@ -6721,6 +6720,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;
@@ -6753,7 +6753,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)
{
@@ -7002,7 +7002,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)))
{
@@ -7027,7 +7027,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
@@ -7051,7 +7051,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);
@@ -7074,7 +7074,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
@@ -7136,7 +7136,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%>"
@@ -7175,7 +7175,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)
@@ -7196,7 +7196,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%>"
@@ -8076,7 +8076,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);
@@ -8173,8 +8173,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)
@@ -8270,7 +8269,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);
@@ -8289,7 +8288,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
@@ -8297,8 +8296,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);
@@ -8307,7 +8305,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
@@ -8331,7 +8329,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). */
@@ -8339,8 +8337,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
@@ -8349,9 +8346,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;
}