* [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
@ 2020-09-01 13:16 Chung-Lin Tang
2020-09-16 14:11 ` Chung-Lin Tang
` (2 more replies)
0 siblings, 3 replies; 10+ messages in thread
From: Chung-Lin Tang @ 2020-09-01 13:16 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine Moore,
Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 1810 bytes --]
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.
[-- Attachment #2: omp5-tgtmapping.01.cfrontends.patch --]
[-- Type: text/plain, Size: 20487 bytes --]
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;
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-09-01 13:16 [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts Chung-Lin Tang
@ 2020-09-16 14:11 ` Chung-Lin Tang
2020-09-29 10:16 ` Jakub Jelinek
2020-10-13 13:01 ` Jakub Jelinek
2 siblings, 0 replies; 10+ messages in thread
From: Chung-Lin Tang @ 2020-09-16 14:11 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Tobias Burnus, Catherine Moore,
Thomas Schwinge
Ping this patch set.
Thanks,
Chung-Lin
On 2020/9/1 9:16 PM, Chung-Lin Tang wrote:
> 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.
^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-09-01 13:16 [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts Chung-Lin Tang
2020-09-16 14:11 ` Chung-Lin Tang
@ 2020-09-29 10:16 ` Jakub Jelinek
2020-10-06 14:54 ` Chung-Lin Tang
2020-10-13 13:01 ` Jakub Jelinek
2 siblings, 1 reply; 10+ messages in thread
From: Jakub Jelinek @ 2020-09-29 10:16 UTC (permalink / raw)
To: Chung-Lin Tang
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
On Tue, Sep 01, 2020 at 09:16:23PM +0800, Chung-Lin Tang wrote:
> 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.
Do you think you could mention in detail which exact target mapping changes
in the spec is the patchset attempting to implement?
5.0 unfortunately contains many target mapping changes and this patchset
can't implement them all and it would be easier to see the list of rules
(e.g. from openmp-diff-full-4.5-5.0.pdf, if you don't have that one, I can
send it to you), rather than trying to guess them from the patchset.
Thanks.
> gcc/c-family/
> * c-common.h (c_omp_adjust_clauses): New declaration.
> * c-omp.c (c_omp_adjust_clauses): New function.
This function name is too broad, it should have target in it as it is
for processing target* construct clauses only.
Jakub
^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-09-29 10:16 ` Jakub Jelinek
@ 2020-10-06 14:54 ` Chung-Lin Tang
0 siblings, 0 replies; 10+ messages in thread
From: Chung-Lin Tang @ 2020-10-06 14:54 UTC (permalink / raw)
To: Jakub Jelinek
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
On 2020/9/29 6:16 PM, Jakub Jelinek wrote:
> On Tue, Sep 01, 2020 at 09:16:23PM +0800, Chung-Lin Tang wrote:
>> 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.
>
> Do you think you could mention in detail which exact target mapping changes
> in the spec is the patchset attempting to implement?
> 5.0 unfortunately contains many target mapping changes and this patchset
> can't implement them all and it would be easier to see the list of rules
> (e.g. from openmp-diff-full-4.5-5.0.pdf, if you don't have that one, I can
> send it to you), rather than trying to guess them from the patchset.
>
> Thanks.
Hi Jakub,
the main implemented features are the clause ordering rules:
"For a given construct, the effect of a map clause with the to, from, or tofrom map-type is
ordered before the effect of a map clause with the alloc, release, or delete map-type."
"If item1 is a list item in a map clause, and item2 is another list item in a map clause on
the same construct that has a base pointer that is, or is part of, item1, then:
* If the map clause(s) appear on a target, target data, or target enter data construct,
then on entry to the corresponding region the effect of the map clause on item1 is ordered
to occur before the effect of the map clause on item2.
\x0f * If the map clause(s) appear on a target, target data, or target exit data construct then
on exit from the corresponding region the effect of the map clause on item2 is ordered to
occur before the effect of the map clause on item1."
and the base-pointer attachment behavior:
"If a list item in a map clause has a base pointer, and a pointer variable is present in the device data
environment that corresponds to the base pointer when the effect of the map clause occurs, then if
the corresponding pointer or the corresponding list item is created in the device data environment
on entry to the construct, then:
...
2. The corresponding pointer variable becomes an attached pointer for the corresponding list item."
(these passages are all in the "2.19.7.1 map Clause" section of the 5.0 spec, all are new as
also verified from the diff PDFs you sent us)
Also, because of the these new features, having multiple maps of the same variables now have meaning
in OpenMP, so changes in the C/C++ frontends to relax the no-duplicate rules are also included.
>> gcc/c-family/
>> * c-common.h (c_omp_adjust_clauses): New declaration.
>> * c-omp.c (c_omp_adjust_clauses): New function.
>
> This function name is too broad, it should have target in it as it is
> for processing target* construct clauses only.
>
> Jakub
Sure, I'll update this naming in a later version.
Thanks,
Chung-Lin
^ permalink raw reply [flat|nested] 10+ messages in thread
* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-09-01 13:16 [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts Chung-Lin Tang
2020-09-16 14:11 ` Chung-Lin Tang
2020-09-29 10:16 ` Jakub Jelinek
@ 2020-10-13 13:01 ` Jakub Jelinek
2020-10-28 10:31 ` Chung-Lin Tang
2 siblings, 1 reply; 10+ messages in thread
From: Jakub Jelinek @ 2020-10-13 13:01 UTC (permalink / raw)
To: Chung-Lin Tang
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
On Tue, Sep 01, 2020 at 09:16:23PM +0800, Chung-Lin Tang wrote:
> 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.
Sorry for the delay in patch review and thanks for the standard citations,
that really helps.
> 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)?
> 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.
The changelog has some 8 space indented lines.
> + 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.
> + 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.
> + && (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?
> + {
> + 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.
> --- 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);
?
> 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.
> {
> 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;
So what is this supposed to diagnose now that the restriction that
C++ ditto.
Jakub
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-10-13 13:01 ` Jakub Jelinek
@ 2020-10-28 10:31 ` Chung-Lin Tang
2020-10-29 11:44 ` Jakub Jelinek
0 siblings, 1 reply; 10+ messages in thread
From: Chung-Lin Tang @ 2020-10-28 10:31 UTC (permalink / raw)
To: Jakub Jelinek
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 6643 bytes --]
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
[-- Attachment #2: omp5-tgtmapping.01.cfrontends.v2.patch --]
[-- Type: text/plain, Size: 13993 bytes --]
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);
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-10-28 10:31 ` Chung-Lin Tang
@ 2020-10-29 11:44 ` Jakub Jelinek
2020-11-03 18:02 ` Chung-Lin Tang
0 siblings, 1 reply; 10+ messages in thread
From: Jakub Jelinek @ 2020-10-29 11:44 UTC (permalink / raw)
To: Chung-Lin Tang
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
On Wed, Oct 28, 2020 at 06:31:22PM +0800, Chung-Lin Tang wrote:
> > > + 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.
Many functions try hard to avoid O(n^2) issues, see e.g. all the bitmap
handling in *finish_omp_clauses etc.
One can have tens of thousands of clauses and then the quadraticness will
hit hard. This does a mere OMP_CLAUSE_DECL (c) == ptr comparison, so it
is only about the decls and decls can be very easily handled through
DECL_UID (bitmaps, hash sets/maps/tables).
> +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?
> --- 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.
Jakub
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-10-29 11:44 ` Jakub Jelinek
@ 2020-11-03 18:02 ` Chung-Lin Tang
2020-11-06 9:52 ` Jakub Jelinek
2021-06-04 12:18 ` [PATCH] openmp: Call c_omp_adjust_map_clauses even for combined target [PR100902] Jakub Jelinek
0 siblings, 2 replies; 10+ messages in thread
From: Chung-Lin Tang @ 2020-11-03 18:02 UTC (permalink / raw)
To: Jakub Jelinek
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
[-- Attachment #1: Type: text/plain, Size: 3967 bytes --]
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.
[-- Attachment #2: omp5-tgtmapping.01.cfrontends.v3.patch --]
[-- Type: text/plain, Size: 15447 bytes --]
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);
^ permalink raw reply [flat|nested] 10+ messages in thread* Re: [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts
2020-11-03 18:02 ` Chung-Lin Tang
@ 2020-11-06 9:52 ` Jakub Jelinek
2021-06-04 12:18 ` [PATCH] openmp: Call c_omp_adjust_map_clauses even for combined target [PR100902] Jakub Jelinek
1 sibling, 0 replies; 10+ messages in thread
From: Jakub Jelinek @ 2020-11-06 9:52 UTC (permalink / raw)
To: Chung-Lin Tang
Cc: gcc-patches, Tobias Burnus, Catherine Moore, Thomas Schwinge
On Wed, Nov 04, 2020 at 02:02:25AM +0800, Chung-Lin Tang wrote:
> 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.
Ok, thanks.
Jakub
^ permalink raw reply [flat|nested] 10+ messages in thread
* [PATCH] openmp: Call c_omp_adjust_map_clauses even for combined target [PR100902]
2020-11-03 18:02 ` Chung-Lin Tang
2020-11-06 9:52 ` Jakub Jelinek
@ 2021-06-04 12:18 ` Jakub Jelinek
1 sibling, 0 replies; 10+ messages in thread
From: Jakub Jelinek @ 2021-06-04 12:18 UTC (permalink / raw)
To: Chung-Lin Tang; +Cc: Tobias Burnus, gcc-patches, Thomas Schwinge
Hi!
When looking at in_reduction support for target, I've noticed that
c_omp_adjust_map_clauses is not called for the combined target case.
The following patch fixes it, will commit once tested.
Unfortunately, there are other issues.
One is (also mentioned in the PR) that currently the pointer attachment
stuff seems to be clause ordering dependent (the standard says that clause
ordering on the same construct does not matter), the baz and qux cases
in the PR are rejected while when swapped it is accepted.
Note, the order of clauses in GCC really is treated as insignificant
initially and only later on the compiler can adjust the ordering (e.g. when
we sort map clauses based on what they refer to etc.) and in particular,
clauses from parsing is reverse of the order in user code, while
c_omp_split_clauses performed for combined/composite constructs typically
reverses that ordering, i.e. makes it follow the user code ordering.
And another one is I'm slightly afraid c_omp_adjust_map_clauses might
misbehave in templates, though haven't tried to verify it with testcases.
When processing_template_decl, the non-dependent clauses will be handled
usually the same as when not in a template, but dependent clauses aren't
processed or only limited processing is done there, and rest is deferred
till later. From quick skimming of c_omp_adjust_map_clauses, it seems
it might not be very happy about non-processed map clauses that might
still have the TREE_LIST representation of array sections, or might
not have finalized decls or base decls etc.
So, for this I wonder if cp_parser_omp_target (and other cp/parser.c
callers of c_omp_adjust_map_clauses) shouldn't call it only
if (!processing_template_decl) - perhaps you could add
cp_omp_adjust_map_clauses wrapper that would be
if (!processing_template_decl)
c_omp_adjust_map_clauses (...);
- and call c_omp_adjust_map_clauses from within pt.c after the clauses
are tsubsted and finish_omp_clauses is called again.
2021-06-04 Jakub Jelinek <jakub@redhat.com>
PR c/100902
* c-parser.c (c_parser_omp_target): Call c_omp_adjust_map_clauses
even when target is combined with other constructs.
* parser.c (cp_parser_omp_target): Call c_omp_adjust_map_clauses
even when target is combined with other constructs.
* c-c++-common/gomp/pr100902-1.c: New test.
--- gcc/c/c-parser.c.jj 2021-06-04 11:15:11.616690819 +0200
+++ gcc/c/c-parser.c 2021-06-04 13:15:42.788471162 +0200
@@ -20133,6 +20133,7 @@ c_parser_omp_target (c_parser *parser, e
tree stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+ c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
OMP_TARGET_BODY (stmt) = block;
OMP_TARGET_COMBINED (stmt) = 1;
SET_EXPR_LOCATION (stmt, loc);
--- gcc/cp/parser.c.jj 2021-05-31 10:11:15.145978965 +0200
+++ gcc/cp/parser.c 2021-06-04 13:17:00.952392489 +0200
@@ -42233,6 +42233,7 @@ cp_parser_omp_target (cp_parser *parser,
tree stmt = make_node (OMP_TARGET);
TREE_TYPE (stmt) = void_type_node;
OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+ c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
OMP_TARGET_BODY (stmt) = body;
OMP_TARGET_COMBINED (stmt) = 1;
SET_EXPR_LOCATION (stmt, pragma_tok->location);
--- gcc/testsuite/c-c++-common/gomp/pr100902-1.c.jj 2021-06-04 13:13:48.471048762 +0200
+++ gcc/testsuite/c-c++-common/gomp/pr100902-1.c 2021-06-04 13:02:07.655723606 +0200
@@ -0,0 +1,17 @@
+/* PR c/100902 */
+
+void
+foo (int *ptr)
+{
+ #pragma omp target map (ptr, ptr[:4])
+ #pragma omp parallel master
+ ptr[0] = 1;
+}
+
+void
+bar (int *ptr)
+{
+ #pragma omp target parallel map (ptr[:4], ptr)
+ #pragma omp master
+ ptr[0] = 1;
+}
Jakub
^ permalink raw reply [flat|nested] 10+ messages in thread
end of thread, other threads:[~2021-06-04 12:18 UTC | newest]
Thread overview: 10+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-09-01 13:16 [PATCH, 1/3, OpenMP] Target mapping changes for OpenMP 5.0, front-end parts Chung-Lin Tang
2020-09-16 14:11 ` Chung-Lin Tang
2020-09-29 10:16 ` Jakub Jelinek
2020-10-06 14:54 ` Chung-Lin Tang
2020-10-13 13:01 ` Jakub Jelinek
2020-10-28 10:31 ` Chung-Lin Tang
2020-10-29 11:44 ` Jakub Jelinek
2020-11-03 18:02 ` Chung-Lin Tang
2020-11-06 9:52 ` Jakub Jelinek
2021-06-04 12:18 ` [PATCH] openmp: Call c_omp_adjust_map_clauses even for combined target [PR100902] Jakub Jelinek
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).