Hi Jakub,
this patch set implements parts of the target mapping changes introduced
in OpenMP 5.0, mainly the attachment requirements for pointer-based
list items, and the clause ordering.
The first patch here are the C/C++ front-end changes.
The entire set of changes has been tested for without regressions for
the compiler and libgomp. Hope this is ready to commit to master.
Thanks,
Chung-Lin
gcc/c-family/
* c-common.h (c_omp_adjust_clauses): New declaration.
* c-omp.c (c_omp_adjust_clauses): New function.
gcc/c/
* c-parser.c (c_parser_omp_target_data): Add use of
new c_omp_adjust_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_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 4fc64bc4aa6..9ef85b401f0 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1208,14 +1208,15 @@ extern tree c_omp_declare_simd_clauses_to_numbers
(tree, tree);
extern void c_omp_declare_simd_clauses_to_decls (tree, tree);
extern bool c_omp_predefined_variable (tree);
extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree);
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
c_tree_chain_next (tree t)
{
/* TREE_CHAIN of a type is TYPE_STUB_DECL, which is different
kind of object, never a long chain of nodes. Prefer
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index d7cff0f4cca..596f33cebfb 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2575,7 +2575,51 @@ c_omp_map_clause_name (tree clause, bool oacc)
case GOMP_MAP_DEVICE_RESIDENT: return "device_resident";
case GOMP_MAP_LINK: return "link";
case GOMP_MAP_FORCE_DEVICEPTR: return "deviceptr";
default: break;
}
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))
+ {
+ 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 a8bc301ffad..92dfe3b6a4a 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19452,14 +19452,15 @@ c_parser_omp_teams (location_t loc, c_parser *parser,
static tree
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;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (*pc))
{
case GOMP_MAP_TO:
@@ -19469,14 +19470,15 @@ c_parser_omp_target_data (location_t loc, c_parser
*parser, bool *if_p)
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target data%> with map-type other "
"than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
"on %<map%> clause");
@@ -19592,27 +19594,29 @@ c_parser_omp_target_enter_data (location_t loc,
c_parser *parser,
c_parser_skip_to_pragma_eol (parser, false);
return NULL_TREE;
}
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;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (*pc))
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target enter data%> with map-type other "
"than %<to%> or %<alloc%> on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
@@ -19676,29 +19680,30 @@ c_parser_omp_target_exit_data (location_t loc,
c_parser *parser,
c_parser_skip_to_pragma_eol (parser, false);
return NULL_TREE;
}
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;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (*pc))
{
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH_DETACH:
break;
default:
map_seen |= 1;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target exit data%> with map-type other "
"than %<from%>, %<release%> or %<delete%> on %<map%>"
" clause");
@@ -19900,14 +19905,16 @@ c_parser_omp_target (c_parser *parser, enum
pragma_context context, bool *if_p)
stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
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);
add_stmt (c_parser_omp_structured_block (parser, if_p));
OMP_TARGET_BODY (stmt) = c_end_compound_stmt (loc, block, true);
SET_EXPR_LOCATION (stmt, loc);
@@ -19924,14 +19931,15 @@ check_clauses:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
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),
"%<#pragma omp target%> with map-type other "
"than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
"on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 0d639b60ea3..17ac2f566da 100644
--- 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);
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& !c_mark_addressable (t))
return false;
@@ -14682,15 +14683,16 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
t = TREE_OPERAND (t, 0);
}
}
if (remove)
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))
+ || bitmap_bit_p (&map_head, DECL_UID (t)))
break;
}
}
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)))
{
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)
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", t);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in map clauses", t);
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),
"%qD appears more than once in data clauses", t);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears both in data and map clauses", t);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7cc2dbed5fe..7773f9d4f79 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40449,14 +40449,15 @@ cp_parser_omp_teams (cp_parser *parser, cp_token
*pragma_tok,
static tree
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;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (*pc))
{
case GOMP_MAP_TO:
@@ -40467,14 +40468,15 @@ cp_parser_omp_target_data (cp_parser *parser,
cp_token *pragma_tok, bool *if_p)
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
map_seen = 3;
break;
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;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target data%> with map-type other "
"than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
"on %<map%> clause");
@@ -40550,28 +40552,30 @@ cp_parser_omp_target_enter_data (cp_parser *parser,
cp_token *pragma_tok,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return NULL_TREE;
}
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;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (*pc))
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALLOC:
map_seen = 3;
break;
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;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target enter data%> with map-type other "
"than %<to%> or %<alloc%> on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
@@ -40638,14 +40642,15 @@ cp_parser_omp_target_exit_data (cp_parser *parser,
cp_token *pragma_tok,
cp_parser_skip_to_pragma_eol (parser, pragma_tok);
return NULL_TREE;
}
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;)
{
if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (*pc))
{
case GOMP_MAP_FROM:
@@ -40653,14 +40658,15 @@ cp_parser_omp_target_exit_data (cp_parser *parser,
cp_token *pragma_tok,
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
map_seen = 3;
break;
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;
error_at (OMP_CLAUSE_LOCATION (*pc),
"%<#pragma omp target exit data%> with map-type other "
"than %<from%>, %<release%> or %<delete%> on %<map%>"
" clause");
@@ -40901,14 +40907,16 @@ cp_parser_omp_target (cp_parser *parser, cp_token
*pragma_tok,
stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
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);
SET_EXPR_LOCATION (stmt, pragma_tok->location);
add_stmt (stmt);
@@ -40924,14 +40932,15 @@ check_clauses:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
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),
"%<#pragma omp target%> with map-type other "
"than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
"on %<map%> clause");
*pc = OMP_CLAUSE_CHAIN (*pc);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index b71ca0729a8..0f6b36f2dab 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5373,16 +5373,17 @@ handle_omp_array_sections (tree c, enum
c_omp_region_type ort)
}
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)
{
- 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);
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
t = TREE_OPERAND (t, 0);
gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
@@ -5414,16 +5415,20 @@ handle_omp_array_sections (tree c, enum
c_omp_region_type ort)
&& TYPE_REF_P (TREE_TYPE (ptr))
&& INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
{
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
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;
OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2);
OMP_CLAUSE_CHAIN (c2) = c3;
}
}
@@ -7400,15 +7405,15 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
OMP_CLAUSE_SIZE (c) = size_zero_node;
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
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);
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_)
@@ -7446,15 +7451,16 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
}
if (remove)
break;
if (REFERENCE_REF_P (t))
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))
+ || bitmap_bit_p (&map_head, DECL_UID (t)))
goto handle_map_references;
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
break;
@@ -7540,30 +7546,35 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
"%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))
- && (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
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in map clauses", t);
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),
"%qD appears more than once in data clauses", t);
else
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears both in data and map clauses", t);
@@ -7591,23 +7602,25 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
= TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
}
else if (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))
+ != 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;
+ = ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ ? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER);
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);
OMP_CLAUSE_DECL (c2) = t;
OMP_CLAUSE_SIZE (c2) = size_zero_node;
