Hi Jakub,
Thank you for the review.
On 2020/10/13 9:01 PM, Jakub Jelinek wrote:
gcc/c-family/
* c-common.h (c_omp_adjust_clauses): New declaration.
* c-omp.c (c_omp_adjust_clauses): New function.
Besides the naming, I wonder why is it done in a separate function and so
early, can't what the function does be done either in
{,c_}finish_omp_clauses (provided we'd pass separate ORT_OMP vs.
ORT_OMP_TARGET to it to determine if it is target region vs. anything else),
or perhaps even better during gimplification (gimplify_scan_omp_clauses)?
I figured that differentiating with something like "C_ORT_OMP_TARGET" could be
more error prone to adjust changes related to C_ORT_OMP across the code, plus
this has the added benefit of sharing a single place of handling logic across
C/C++.
You're right about the need for early addressable-marking. Learned that the hard
way, one of my prior attempts tried to place this code somewhere in gimplify,
didn't work.
gcc/cp/
* parser.c (cp_parser_omp_target_data): Add use of
new c_omp_adjust_clauses function. Add GOMP_MAP_ATTACH_DETACH as
handled map clause kind.
(cp_parser_omp_target_enter_data): Likewise.
(cp_parser_omp_target_exit_data): Likewise.
(cp_parser_omp_target): Likewise.
* semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to
use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix
interaction between reference case and attach/detach.
(finish_omp_clauses): Adjust bitmap checks to allow struct decl and
same struct field access to co-exist on OpenMP construct.
The changelog has some 8 space indented lines.
I'll take care of that in the final git push.
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE)
+ {
+ tree ptr = OMP_CLAUSE_DECL (c);
+ bool ptr_mapped = false;
+ if (is_target)
+ {
+ for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m))
Isn't this O(n^2) in number of clauses? I mean, e.g. for the equality
comparisons (but see below) it could be dealt with e.g. using some bitmap
with DECL_UIDs.
At this stage, we really don't assume any ordering of the clauses, nor try to
modify its ordering yet, so the base-pointer map (if it exists) could be any
where in the list (building some "visited set" isn't really suitable here).
I don't think this is really that much an issue of concern though.
+ if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_DECL (m) == ptr
Does it really need to be equality? I mean it will be for
map(tofrom:ptr) map(tofrom:ptr[:32])
but what about e.g.
map(tofrom:structx) map(tofrom:structx.ptr[:32])
? It is true that likely we don't parse this yet though.
The code for COMPONENT_REF based expressions are actually handled quite
differently
in gimplify_scan_omp_clauses. Not completely sure there's nothing to handle for
the
code in this patch set, but will have to discover such testcases later.
+ && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM))
What about the always modified mapping kinds?
Took care of that.
+ {
+ ptr_mapped = true;
+ break;
+ }
+
+ if (!ptr_mapped
+ && DECL_P (ptr)
+ && is_global_var (ptr)
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (ptr)))
+ ptr_mapped = true;
+ }
+
+ /* If the pointer variable was mapped, or if this is not an offloaded
+ target region, adjust the map kind to attach/detach. */
+ if (ptr_mapped || !is_target)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH);
+ c_common_mark_addressable_vec (ptr);
Though perhaps this is argument why it needs to be done in the FEs and not
during gimplification, because it is hard to mark something addressable at
that point.
Discussed above.
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13580,16 +13580,17 @@ handle_omp_array_sections (tree c, enum
c_omp_region_type ort)
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)
{
- gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
+ gomp_map_kind k
+ = ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER);
So what kind of C_ORT_* would be left after this change?
C_ORT_*DECLARE_SIMD shouldn't have any kind of array sections in it.
So maybe just
OMP_CLAUSE_SET_MAP_KIND (c2, k);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
?
I have changed this code to just "OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_ATTACH_DETACH);"
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || bitmap_bit_p (&map_head, DECL_UID (t)))
break;
Shall this change apply to OpenACC too?
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE is not a variable in %qs clause", t,
@@ -14751,29 +14753,36 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears both in data and map clauses", t);
remove = true;
}
else
bitmap_set_bit (&generic_head, DECL_UID (t));
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
Ditto. Otherwise, what shall this diagnose now that the restriction that
the same list item may not appear in multiple clauses is gone.
Thanks for catching this, I've added "C_ORT_OMP" tests to these parts.
Attached is the revised patch, dubbed "v2". Entire patch set re-tested with no
regressions
for gcc, g++, gfortran, and libgomp on x86_64-linux with nvptx offloading.
Thanks,
Chung-Lin
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index bb38e6c76a4..35ad417f9cc 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1221,6 +1221,7 @@ extern enum omp_clause_defaultmap_kind
c_omp_predetermined_mapping (tree);
extern tree c_omp_check_context_selector (location_t, tree);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
extern const char *c_omp_map_clause_name (tree, bool);
+extern void c_omp_adjust_clauses (tree, bool);
/* Return next tree in the chain for chain_next walking of tree nodes. */
static inline tree
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index d7cff0f4cca..fd4995d67fb 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2579,3 +2579,50 @@ c_omp_map_clause_name (tree clause, bool oacc)
}
return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
}
+
+/* Adjust map clauses after normal clause parsing, mainly to turn specific
+ base-pointer map cases into attach/detach and mark them addressable. */
+void
+c_omp_adjust_clauses (tree clauses, bool is_target)
+{
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) != ARRAY_TYPE)
+ {
+ tree ptr = OMP_CLAUSE_DECL (c);
+ bool ptr_mapped = false;
+ if (is_target)
+ {
+ for (tree m = clauses; m; m = OMP_CLAUSE_CHAIN (m))
+ if (OMP_CLAUSE_CODE (m) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_DECL (m) == ptr
+ && (OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TO
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_TOFROM
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TO
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_FROM
+ || OMP_CLAUSE_MAP_KIND (m) == GOMP_MAP_ALWAYS_TOFROM))
+ {
+ ptr_mapped = true;
+ break;
+ }
+
+ if (!ptr_mapped
+ && DECL_P (ptr)
+ && is_global_var (ptr)
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (ptr)))
+ ptr_mapped = true;
+ }
+
+ /* If the pointer variable was mapped, or if this is not an offloaded
+ target region, adjust the map kind to attach/detach. */
+ if (ptr_mapped || !is_target)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH);
+ c_common_mark_addressable_vec (ptr);
+ }
+ }
+}
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b6a7ef4c92b..24fd6afb9a3 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19470,6 +19470,7 @@ c_parser_omp_target_data (location_t loc, c_parser
*parser, bool *if_p)
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data");
+ c_omp_adjust_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
@@ -19487,6 +19488,7 @@ c_parser_omp_target_data (location_t loc, c_parser
*parser, bool *if_p)
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
@@ -19610,6 +19612,7 @@ c_parser_omp_target_enter_data (location_t loc,
c_parser *parser,
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
"#pragma omp target enter data");
+ c_omp_adjust_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
@@ -19623,6 +19626,7 @@ c_parser_omp_target_enter_data (location_t loc,
c_parser *parser,
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
@@ -19694,7 +19698,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser
*parser,
tree clauses
= c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data");
-
+ c_omp_adjust_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
@@ -19709,6 +19713,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser
*parser,
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
@@ -19918,6 +19923,8 @@ c_parser_omp_target (c_parser *parser, enum
pragma_context context, bool *if_p)
OMP_TARGET_CLAUSES (stmt)
= c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
"#pragma omp target");
+ c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
pc = &OMP_TARGET_CLAUSES (stmt);
keep_next_level ();
block = c_begin_compound_stmt (true);
@@ -19942,6 +19949,7 @@ check_clauses:
case GOMP_MAP_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 459090e227d..30c48e3a205 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13501,11 +13501,7 @@ handle_omp_array_sections (tree c, enum
c_omp_region_type ort)
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)
- {
- gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, k);
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
@@ -14604,7 +14600,9 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || (ort == C_ORT_OMP
+ && bitmap_bit_p (&map_head, DECL_UID (t))))
break;
}
}
@@ -14673,7 +14671,9 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
else
bitmap_set_bit (&generic_head, DECL_UID (t));
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && (ort != C_ORT_OMP
+ || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -14687,7 +14687,13 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
remove = true;
}
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
- || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ && ort == C_ORT_ACC)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7ec7d42773c..33cc5069ba0 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40510,6 +40510,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token
*pragma_tok, bool *if_p)
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
"#pragma omp target data", pragma_tok);
+ c_omp_adjust_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
@@ -40528,6 +40529,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token
*pragma_tok, bool *if_p)
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
@@ -40611,6 +40613,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser,
cp_token *pragma_tok,
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
"#pragma omp target enter data", pragma_tok);
+ c_omp_adjust_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
@@ -40625,6 +40628,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser,
cp_token *pragma_tok,
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
@@ -40699,6 +40703,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser,
cp_token *pragma_tok,
tree clauses
= cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
"#pragma omp target exit data", pragma_tok);
+ c_omp_adjust_clauses (clauses, false);
int map_seen = 0;
for (tree *pc = &clauses; *pc;)
{
@@ -40714,6 +40719,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser,
cp_token *pragma_tok,
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
@@ -40962,6 +40968,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token
*pragma_tok,
OMP_TARGET_CLAUSES (stmt)
= cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
"#pragma omp target", pragma_tok);
+ c_omp_adjust_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
pc = &OMP_TARGET_CLAUSES (stmt);
keep_next_level (true);
OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
@@ -40985,6 +40993,7 @@ check_clauses:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 1e42cd799c2..a8dc769db34 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5382,11 +5382,7 @@ handle_omp_array_sections (tree c, enum
c_omp_region_type ort)
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)
- {
- gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, k);
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
@@ -5424,8 +5420,12 @@ handle_omp_array_sections (tree c, enum
c_omp_region_type ort)
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
OMP_CLAUSE_DECL (c3) = ptr;
- if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER)
- OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
+ if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH)
+ {
+ OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ }
else
OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
OMP_CLAUSE_SIZE (c3) = size_zero_node;
@@ -7411,7 +7411,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
- if (ort == C_ORT_ACC
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
&& TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
@@ -7457,7 +7457,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
t = TREE_OPERAND (t, 0);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+ || (ort == C_ORT_OMP
+ && bitmap_bit_p (&map_head, DECL_UID (t))))
goto handle_map_references;
}
}
@@ -7551,13 +7553,12 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
bitmap_set_bit (&generic_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
- && (ort != C_ORT_ACC
- || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
+ && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in motion clauses", t);
- if (ort == C_ORT_ACC)
+ else if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
@@ -7566,7 +7567,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
remove = true;
}
else if (bitmap_bit_p (&generic_head, DECL_UID (t))
- || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ && ort == C_ORT_ACC)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
{
if (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -7602,17 +7609,14 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& (OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_ALWAYS_POINTER))
+ != GOMP_MAP_ALWAYS_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_ATTACH_DETACH))
{
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
- {
- gomp_map_kind k
- = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
- : GOMP_MAP_ALWAYS_POINTER;
- OMP_CLAUSE_SET_MAP_KIND (c2, k);
- }
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);
