Hi Jakub,
here is v3 of this patch set.
On 2020/10/29 7:44 PM, Jakub Jelinek wrote:
+extern void c_omp_adjust_clauses (tree, bool);
So, can you please rename the function to either
c_omp_adjust_target_clauses or c_omp_adjust_mapping_clauses or
c_omp_adjust_map_clauses?
I've renamed it to 'c_omp_adjust_map_clauses'.
--- 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
If this is only meant to handle decls, perhaps there should be
&& DECL_P (OMP_CLAUSE_DECL (c))
?
+ && 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;
+ }
What you could e.g. do is have this loop at the start of function, with
&& DECL_P (OMP_CLAUSE_DECL (m))
instead of the == ptr check, and perhaps && POINTER_TYPE_P (TREE_TYPE
(OMP_CLAUSE_DECL (m))) check and set a bit in a bitmap for each such decl,
then in the GOMP_MAP_FIRSTPRIVATE_POINTER loop just check the bitmap.
Or, keep it in the loop like it is above, but populate the bitmap
lazily (upon seeing the first GOMP_MAP_FIRSTPRIVATE_POINTER) and for further
ones just use it.
I re-wrote c_omp_adjust_map_clauses to address the complexity issues you
mentioned,
now it should be limited by a linear pass to collect and merge the firstprivate
base
pointer + existence of a mapping of it, using a hash_map.
Patch set has been re-tested with no regressions for gcc, g++, gfortran, and
libgomp.
Thanks,
Chung-Lin
gcc/c-family/
* c-common.h (c_omp_adjust_map_clauses): New declaration.
* c-omp.c (c_omp_adjust_map_clauses): New function.
gcc/c/
* c-parser.c (c_parser_omp_target_data): Add use of
new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as
handled map clause kind.
(c_parser_omp_target_enter_data): Likewise.
(c_parser_omp_target_exit_data): Likewise.
(c_parser_omp_target): Likewise.
* c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to
use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type.
(c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and
same struct field access to co-exist on OpenMP construct.
gcc/cp/
* parser.c (cp_parser_omp_target_data): Add use of
new c_omp_adjust_map_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.
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index bb38e6c76a4..3eb909a2946 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_map_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..275c6afabe1 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2579,3 +2579,92 @@ 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_map_clauses (tree clauses, bool is_target)
+{
+ if (!is_target)
+ {
+ /* If this is not a target construct, just turn firstprivate pointers
+ into attach/detach, the runtime will check and do the rest. */
+
+ 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
+ && DECL_P (OMP_CLAUSE_DECL (c))
+ && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
+ {
+ tree ptr = OMP_CLAUSE_DECL (c);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH);
+ c_common_mark_addressable_vec (ptr);
+ }
+ return;
+ }
+
+ struct map_clause
+ {
+ tree clause;
+ bool firstprivate_ptr_p;
+ bool decl_mapped;
+ bool omp_declare_target;
+ map_clause (void) : clause (NULL_TREE), firstprivate_ptr_p (false),
+ decl_mapped (false), omp_declare_target (false) { }
+ };
+
+ hash_map<tree, map_clause> maps;
+
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && DECL_P (OMP_CLAUSE_DECL (c)))
+ {
+ /* If this is for a target construct, the firstprivate pointer
+ is changed to attach/detach if either is true:
+ (1) the base-pointer is mapped in this same construct, or
+ (2) the base-pointer is a variable place on the device by
+ "declare target" directives.
+
+ Here we iterate through all map clauses collecting these cases,
+ and merge them with a hash_map to process below. */
+
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
+ {
+ tree ptr = OMP_CLAUSE_DECL (c);
+ map_clause &mc = maps.get_or_insert (ptr);
+ if (mc.clause == NULL_TREE)
+ mc.clause = c;
+ mc.firstprivate_ptr_p = true;
+
+ if (is_global_var (ptr)
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (ptr)))
+ mc.omp_declare_target = true;
+ }
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TOFROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_FROM
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
+ {
+ map_clause &mc = maps.get_or_insert (OMP_CLAUSE_DECL (c));
+ mc.decl_mapped = true;
+ }
+ }
+
+ for (hash_map<tree, map_clause>::iterator i = maps.begin ();
+ i != maps.end (); ++i)
+ {
+ map_clause &mc = (*i).second;
+
+ if (mc.firstprivate_ptr_p
+ && (mc.decl_mapped || mc.omp_declare_target))
+ {
+ OMP_CLAUSE_SET_MAP_KIND (mc.clause, GOMP_MAP_ATTACH_DETACH);
+ c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause));
+ }
+ }
+}
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b6a7ef4c92b..ab7b0bbc29f 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_map_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_map_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_map_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_map_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..ada6662fca7 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..8527a7d0478 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_map_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_map_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_map_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_map_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);
