mirrors/gcc
1
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-11-23 19:03:59 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

openmp: Implement OpenMP 5.0 base-pointer attachement and clause ordering

Browse Source 
This patch implements some parts of the target variable mapping changes
specified in OpenMP 5.0, including base-pointer attachment/detachment
behavior for array section list-items in map clauses, and ordering of
map clauses according to map kind.

2020-11-10  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/c-family/ChangeLog:

	* c-common.h (c_omp_adjust_map_clauses): New declaration.
	* c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses.
	(c_omp_adjust_map_clauses): New function.

gcc/c/ChangeLog:

	* 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/ChangeLog:

	* 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.

gcc/ChangeLog:

	* gimplify.c (is_or_contains_p): New static helper function.
	(omp_target_reorder_clauses): New function.
	(gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to
	reorder clause list according to OpenMP 5.0 rules. Add handling of
	GOMP_MAP_ATTACH_DETACH for OpenMP cases.
	* omp-low.c (is_omp_target): New static helper function.
	(scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH
	for OpenMP cases.
	(lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for
	OpenMP cases.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid.
	* gfortran.dg/gomp/map-2.f90: Likewise.
	* c-c++-common/gomp/map-5.c: New testcase.

libgomp/ChangeLog:

	* libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag
	usable.
	* oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to
	'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'.
	(goacc_enter_datum): Likewise for call to gomp_map_vars_async.
	(goacc_enter_data_internal): Likewise.
	* target.c (gomp_map_vars_internal):
	Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use
	of gomp_attach_pointer for OpenMP cases.
	(gomp_exit_data): Add handling of GOMP_MAP_DETACH.
	(GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH.
	* testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
This commit is contained in:
Chung-Lin Tang 2020-11-10 03:36:58 -08:00
parent cba3d03da6
commit 9e62802422
15 changed files with 616 additions and 103 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
gcc/c-family/c-common.h
View File

@ -1224,6 +1224,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

90
gcc/c-family/c-omp.c
View File

@ -2771,3 +2771,93 @@ c_omp_map_clause_name (tree clause, bool oacc)
}
return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
}
/* Used to merge map clause information in c_omp_adjust_map_clauses. */
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) { }
};
/* 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;
}
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));
}
}
}

10
gcc/c/c-parser.c
View File

@ -19511,6 +19511,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;)
{
@ -19528,6 +19529,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;
@ -19651,6 +19653,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;)
{
@ -19664,6 +19667,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;
@ -19735,7 +19739,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;)
{
@ -19750,6 +19754,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;
@ -19960,6 +19965,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);
@ -19984,6 +19991,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),

22
gcc/c/c-typeck.c
View File

@ -13584,11 +13584,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
@ -14711,7 +14707,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;
}
}
@ -14780,7 +14778,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),
@ -14794,7 +14794,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),

9
gcc/cp/parser.c
View File

@ -40785,6 +40785,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;)
{
@ -40803,6 +40804,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;
@ -40886,6 +40888,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;)
{
@ -40900,6 +40903,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;
@ -40974,6 +40978,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;)
{
@ -40989,6 +40994,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;
@ -41238,6 +41244,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);
@ -41261,6 +41269,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),

44
gcc/cp/semantics.c
View File

@ -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;
@ -7486,7 +7486,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);
@ -7532,7 +7532,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;
}
}
@ -7626,13 +7628,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
@ -7641,7 +7642,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),
@ -7677,17 +7684,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);

272
gcc/gimplify.c
View File

@ -8364,6 +8364,113 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
return base;
}
/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */
static bool
is_or_contains_p (tree expr, tree base_ptr)
{
while (expr != base_ptr)
if (TREE_CODE (base_ptr) == COMPONENT_REF)
base_ptr = TREE_OPERAND (base_ptr, 0);
else
break;
return expr == base_ptr;
}
/* Implement OpenMP 5.x map ordering rules for target directives. There are
several rules, and with some level of ambiguity, hopefully we can at least
collect the complexity here in one place. */
static void
omp_target_reorder_clauses (tree *list_p)
{
/* Collect refs to alloc/release/delete maps. */
auto_vec<tree, 32> ard;
tree *cp = list_p;
while (*cp != NULL_TREE)
if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
{
/* Unlink cp and push to ard. */
tree c = *cp;
tree nc = OMP_CLAUSE_CHAIN (c);
*cp = nc;
ard.safe_push (c);
/* Any associated pointer type maps should also move along. */
while (*cp != NULL_TREE
&& OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
|| OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
{
c = *cp;
nc = OMP_CLAUSE_CHAIN (c);
*cp = nc;
ard.safe_push (c);
}
}
else
cp = &OMP_CLAUSE_CHAIN (*cp);
/* Link alloc/release/delete maps to the end of list. */
for (unsigned int i = 0; i < ard.length (); i++)
{
*cp = ard[i];
cp = &OMP_CLAUSE_CHAIN (ard[i]);
}
*cp = NULL_TREE;
/* OpenMP 5.0 requires that pointer variables are mapped before
its use as a base-pointer. */
auto_vec<tree *, 32> atf;
for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
{
/* Collect alloc, to, from, to/from clause tree pointers. */
gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
if (k == GOMP_MAP_ALLOC
|| k == GOMP_MAP_TO
|| k == GOMP_MAP_FROM
|| k == GOMP_MAP_TOFROM
|| k == GOMP_MAP_ALWAYS_TO
|| k == GOMP_MAP_ALWAYS_FROM
|| k == GOMP_MAP_ALWAYS_TOFROM)
atf.safe_push (cp);
}
for (unsigned int i = 0; i < atf.length (); i++)
if (atf[i])
{
tree *cp = atf[i];
tree decl = OMP_CLAUSE_DECL (*cp);
if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
{
tree base_ptr = TREE_OPERAND (decl, 0);
STRIP_TYPE_NOPS (base_ptr);
for (unsigned int j = i + 1; j < atf.length (); j++)
{
tree *cp2 = atf[j];
tree decl2 = OMP_CLAUSE_DECL (*cp2);
if (is_or_contains_p (decl2, base_ptr))
{
/* Move *cp2 to before *cp. */
tree c = *cp2;
*cp2 = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *cp;
*cp = c;
atf[j] = NULL;
}
}
}
}
}
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
@ -8405,6 +8512,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
if (code == OMP_TARGET
|| code == OMP_TARGET_DATA
|| code == OMP_TARGET_ENTER_DATA
|| code == OMP_TARGET_EXIT_DATA)
omp_target_reorder_clauses (list_p);
while ((c = *list_p) != NULL)
{
bool remove = false;
@ -8845,15 +8958,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
== GOMP_MAP_FIRSTPRIVATE_REFERENCE)
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
&& TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
{
OMP_CLAUSE_SIZE (c)
= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
false);
omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
GOVD_FIRSTPRIVATE | GOVD_SEEN);
if ((region_type & ORT_TARGET) != 0)
omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
GOVD_FIRSTPRIVATE | GOVD_SEEN);
}
if (!DECL_P (decl))
{
tree d = decl, *pd;
@ -8878,7 +8994,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
bool indir_p = false;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
if ((region_type & ORT_ACC) != 0
if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
&& TREE_CODE (*pd) == COMPONENT_REF
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
&& code != OACC_UPDATE)
@ -8886,9 +9002,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
while (TREE_CODE (decl) == COMPONENT_REF)
{
decl = TREE_OPERAND (decl, 0);
if ((TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1)))
|| INDIRECT_REF_P (decl))
if (((TREE_CODE (decl) == MEM_REF
&& integer_zerop (TREE_OPERAND (decl, 1)))
|| INDIRECT_REF_P (decl))
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== POINTER_TYPE))
{
indir_p = true;
decl = TREE_OPERAND (decl, 0);
@ -8915,8 +9033,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (decl != orig_decl && DECL_P (decl) && indir_p)
{
gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
: GOMP_MAP_ATTACH;
gomp_map_kind k
= ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
/* We have a dereference of a struct member. Make this an
attach/detach operation, and ensure the base pointer is
mapped as a FIRSTPRIVATE_POINTER. */
@ -8925,6 +9044,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree next_clause = OMP_CLAUSE_CHAIN (c);
if (k == GOMP_MAP_ATTACH
&& code != OACC_ENTER_DATA
&& code != OMP_TARGET_ENTER_DATA
&& (!next_clause
|| (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
|| (OMP_CLAUSE_MAP_KIND (next_clause)
@ -8972,17 +9092,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if (code == OACC_UPDATE
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
remove = true;
break;
}
if (DECL_P (decl)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
&& code != OACC_UPDATE)
&& code != OACC_UPDATE
&& code != OMP_TARGET_UPDATE)
{
if (error_operand_p (decl))
{
@ -9044,15 +9159,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
bool has_attachments = false;
/* For OpenACC, pointers in structs should trigger an
attach action. */
if (attach_detach && (region_type & ORT_ACC) != 0)
if (attach_detach
&& ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
|| code == OMP_TARGET_ENTER_DATA
|| code == OMP_TARGET_EXIT_DATA))
{
/* Turn a GOMP_MAP_ATTACH_DETACH clause into a
GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we
have detected a case that needs a GOMP_MAP_STRUCT
mapping added. */
gomp_map_kind k
= (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
: GOMP_MAP_ATTACH;
= ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
@ -9148,33 +9267,38 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
if (scp)
continue;
tree d1 = OMP_CLAUSE_DECL (*sc);
tree d2 = OMP_CLAUSE_DECL (c);
while (TREE_CODE (d1) == ARRAY_REF)
d1 = TREE_OPERAND (d1, 0);
while (TREE_CODE (d2) == ARRAY_REF)
d2 = TREE_OPERAND (d2, 0);
if (TREE_CODE (d1) == INDIRECT_REF)
d1 = TREE_OPERAND (d1, 0);
if (TREE_CODE (d2) == INDIRECT_REF)
d2 = TREE_OPERAND (d2, 0);
while (TREE_CODE (d1) == COMPONENT_REF)
if (TREE_CODE (d2) == COMPONENT_REF
&& TREE_OPERAND (d1, 1)
== TREE_OPERAND (d2, 1))
{
d1 = TREE_OPERAND (d1, 0);
d2 = TREE_OPERAND (d2, 0);
}
else
break;
if (d1 == d2)
if ((region_type & ORT_ACC) != 0)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE appears more than once in map "
"clauses", OMP_CLAUSE_DECL (c));
remove = true;
break;
/* This duplicate checking code is currently only
enabled for OpenACC. */
tree d1 = OMP_CLAUSE_DECL (*sc);
tree d2 = OMP_CLAUSE_DECL (c);
while (TREE_CODE (d1) == ARRAY_REF)
d1 = TREE_OPERAND (d1, 0);
while (TREE_CODE (d2) == ARRAY_REF)
d2 = TREE_OPERAND (d2, 0);
if (TREE_CODE (d1) == INDIRECT_REF)
d1 = TREE_OPERAND (d1, 0);
if (TREE_CODE (d2) == INDIRECT_REF)
d2 = TREE_OPERAND (d2, 0);
while (TREE_CODE (d1) == COMPONENT_REF)
if (TREE_CODE (d2) == COMPONENT_REF
&& TREE_OPERAND (d1, 1)
== TREE_OPERAND (d2, 1))
{
d1 = TREE_OPERAND (d1, 0);
d2 = TREE_OPERAND (d2, 0);
}
else
break;
if (d1 == d2)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qE appears more than once in map "
"clauses", OMP_CLAUSE_DECL (c));
remove = true;
break;
}
}
if (maybe_lt (offset1, offsetn)
|| (known_eq (offset1, offsetn)
@ -9220,6 +9344,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
}
}
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
remove = true;
break;
}
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@ -9236,10 +9368,60 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
else
{
/* DECL_P (decl) == true */
tree *sc;
if (struct_map_to_clause
&& (sc = struct_map_to_clause->get (decl)) != NULL
&& OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
&& decl == OMP_CLAUSE_DECL (*sc))
{
/* We have found a map of the whole structure after a
leading GOMP_MAP_STRUCT has been created, so refill the
leading clause into a map of the whole structure
variable, and remove the current one.
TODO: we should be able to remove some maps of the
following structure element maps if they are of
compatible TO/FROM/ALLOC type. */
OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
remove = true;
break;
}
}
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
if ((code == OMP_TARGET
|| code == OMP_TARGET_DATA
|| code == OMP_TARGET_ENTER_DATA
|| code == OMP_TARGET_EXIT_DATA)
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
octx = octx->outer_context)
{
splay_tree_node n
= splay_tree_lookup (octx->variables,
(splay_tree_key) OMP_CLAUSE_DECL (c));
/* If this is contained in an outer OpenMP region as a
firstprivate value, remove the attach/detach. */
if (n && (n->value & GOVD_FIRSTPRIVATE))
{
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
goto do_add;
}
}
enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
? GOMP_MAP_DETACH
: GOMP_MAP_ATTACH);
OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
}
goto do_add;
case OMP_CLAUSE_DEPEND:

90
gcc/omp-low.c
View File

@ -214,6 +214,21 @@ is_oacc_kernels (omp_context *ctx)
== GF_OMP_TARGET_KIND_OACC_KERNELS));
}
/* Return true if STMT corresponds to an OpenMP target region. */
static bool
is_omp_target (gimple *stmt)
{
if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
{
int kind = gimple_omp_target_kind (stmt);
return (kind == GF_OMP_TARGET_KIND_REGION
|| kind == GF_OMP_TARGET_KIND_DATA
|| kind == GF_OMP_TARGET_KIND_ENTER_DATA
|| kind == GF_OMP_TARGET_KIND_EXIT_DATA);
}
return false;
}
/* If DECL is the artificial dummy VAR_DECL created for non-static
data member privatization, return the underlying "this" parameter,
otherwise return NULL. */
@ -1346,7 +1361,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& DECL_P (decl)
&& ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE))
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
@ -1367,6 +1384,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
break;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& DECL_P (decl)
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& is_omp_target (ctx->stmt))
{
/* If this is an offloaded region, an attach operation should
only exist when the pointer variable is mapped in a prior
clause. */
if (is_gimple_omp_offloaded (ctx->stmt))
gcc_assert
(maybe_lookup_decl (decl, ctx)
|| (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl))));
/* By itself, attach/detach is generated as part of pointer
variable mapping and should not create new variables in the
offloaded region, however sender refs for it must be created
for its address to be passed to the runtime. */
tree field
= build_decl (OMP_CLAUSE_LOCATION (c),
FIELD_DECL, NULL_TREE, ptr_type_node);
SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
insert_field_into_struct (ctx->record_type, field);
/* To not clash with a map of the pointer variable itself,
attach/detach maps have their field looked up by the *clause*
tree expression, not the decl. */
gcc_assert (!splay_tree_lookup (ctx->field_map,
(splay_tree_key) c));
splay_tree_insert (ctx->field_map, (splay_tree_key) c,
(splay_tree_value) field);
break;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
@ -1607,6 +1658,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
break;
if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& is_omp_target (ctx->stmt)
&& !is_gimple_omp_offloaded (ctx->stmt))
break;
if (DECL_P (decl))
{
if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@ -11471,6 +11527,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
@ -11481,8 +11539,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
case GOMP_MAP_LINK:
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
case GOMP_MAP_FORCE_DETACH:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
@ -11537,6 +11593,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
continue;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& is_omp_target (stmt))
{
gcc_assert (maybe_lookup_field (c, ctx));
map_cnt++;
continue;
}
if (!maybe_lookup_field (var, ctx))
continue;
@ -11769,14 +11835,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
if (!maybe_lookup_field (ovar, ctx))
if (!maybe_lookup_field (ovar, ctx)
&& !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
continue;
}
talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
if (nc)
if (nc
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
&& is_omp_target (stmt))
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (c, ctx);
gimplify_assign (x, build_fold_addr_expr (var), &ilist);
}
else if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);

20
gcc/testsuite/c-c++-common/gomp/clauses-2.c
View File

@ -13,35 +13,35 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
bar (p);
#pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
bar (p);
#pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
#pragma omp target map (p) , map (p[0])
bar (p);
#pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
bar (&q);
#pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
bar (p);
#pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t) map (t.r)
bar (&t.r);
#pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t.r) map (t)
bar (&t.r);
#pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t.r) map (t.r)
bar (&t.r);
#pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
#pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
#pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t.s[0]) map (t)
bar (t.s);
#pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t) map(t.s[0])
bar (t.s);
#pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t.s[0]) map (t.s[2])
bar (t.s);
#pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t.t[0:2]) map (t.t[4:6])
bar (t.t);
#pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
#pragma omp target map (t.t[i:j]) map (t.t[k:l])
bar (t.t);
#pragma omp target map (t.s[0]) map (t.r)
bar (t.s);
@ -50,5 +50,5 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
#pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */
bar (t.s);
}

24
gcc/testsuite/c-c++-common/gomp/map-5.c Normal file
View File

@ -0,0 +1,24 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
void foo (void)
{
/* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses. */
int a, b, c;
#pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c)
#pragma omp target exit data map(from:a) map(release:b) map(from:c)
#pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c)
a = b = c = 1;
#pragma omp target enter data map(to:a) map(alloc:b) map(to:c)
#pragma omp target exit data map(from:a) map(delete:b) map(from:c)
}
/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */
/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */
/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */

2
gcc/testsuite/gfortran.dg/gomp/map-2.f90
View File

@ -2,5 +2,5 @@ type t
integer :: i
end type t
type(t) v
!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" }
!$omp target enter data map(to:v%i, v%i)
end

8
libgomp/libgomp.h
View File

@ -1162,10 +1162,10 @@ struct gomp_device_descr
/* Kind of the pragma, for which gomp_map_vars () is called. */
enum gomp_map_vars_kind
{
GOMP_MAP_VARS_OPENACC,
GOMP_MAP_VARS_TARGET,
GOMP_MAP_VARS_DATA,
GOMP_MAP_VARS_ENTER_DATA
GOMP_MAP_VARS_OPENACC = 1,
GOMP_MAP_VARS_TARGET = 2,
GOMP_MAP_VARS_DATA = 4,
GOMP_MAP_VARS_ENTER_DATA = 8
};
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,

9
libgomp/oacc-mem.c
View File

@ -403,7 +403,8 @@ acc_map_data (void *h, void *d, size_t s)
struct target_mem_desc *tgt
= gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
&kinds, true, GOMP_MAP_VARS_ENTER_DATA);
&kinds, true,
GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA);
assert (tgt);
assert (tgt->list_count == 1);
splay_tree_key n = tgt->list[0].key;
@ -572,7 +573,8 @@ goacc_enter_datum (void **hostaddrs, size_t *sizes, void *kinds, int async)
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs, NULL, sizes,
kinds, true, GOMP_MAP_VARS_ENTER_DATA);
kinds, true, (GOMP_MAP_VARS_OPENACC
| GOMP_MAP_VARS_ENTER_DATA));
assert (tgt);
assert (tgt->list_count == 1);
n = tgt->list[0].key;
@ -1202,7 +1204,8 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
struct target_mem_desc *tgt
= gomp_map_vars_async (acc_dev, aq, groupnum, &hostaddrs[i], NULL,
&sizes[i], &kinds[i], true,
GOMP_MAP_VARS_ENTER_DATA);
(GOMP_MAP_VARS_OPENACC
| GOMP_MAP_VARS_ENTER_DATA));
assert (tgt);
gomp_mutex_lock (&acc_dev->lock);

36
libgomp/target.c
View File

@ -683,7 +683,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
tgt->refcount = pragma_kind == GOMP_MAP_VARS_ENTER_DATA ? 0 : 1;
tgt->refcount = (pragma_kind & GOMP_MAP_VARS_ENTER_DATA) ? 0 : 1;
tgt->device_descr = devicep;
tgt->prev = NULL;
struct gomp_coalesce_buf cbuf, *cbufp = NULL;
@ -1212,15 +1212,16 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
/* OpenACC 'attach'/'detach' doesn't affect
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
gomp_attach_pointer (devicep, aq, mem_map, n,
(uintptr_t) hostaddrs[i], sizes[i],
cbufp);
}
else
else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
{
gomp_mutex_unlock (&devicep->lock);
gomp_fatal ("outer struct not mapped for attach");
}
gomp_attach_pointer (devicep, aq, mem_map, n,
(uintptr_t) hostaddrs[i], sizes[i],
cbufp);
continue;
}
default:
@ -1415,7 +1416,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
/* If the variable from "omp target enter data" map-list was already mapped,
tgt is not needed. Otherwise tgt will be freed by gomp_unmap_vars or
gomp_exit_data. */
if (pragma_kind == GOMP_MAP_VARS_ENTER_DATA && tgt->refcount == 0)
if ((pragma_kind & GOMP_MAP_VARS_ENTER_DATA) && tgt->refcount == 0)
{
free (tgt);
tgt = NULL;
@ -2475,6 +2476,19 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
return;
}
for (i = 0; i < mapnum; i++)
if ((kinds[i] & typemask) == GOMP_MAP_DETACH)
{
struct splay_tree_key_s cur_node;
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start + sizeof (void *);
splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &cur_node);
if (n)
gomp_detach_pointer (devicep, NULL, n, (uintptr_t) hostaddrs[i],
false, NULL);
}
for (i = 0; i < mapnum; i++)
{
struct splay_tree_key_s cur_node;
@ -2512,7 +2526,9 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
cur_node.host_end - cur_node.host_start);
if (k->refcount == 0)
gomp_remove_var (devicep, k);
break;
case GOMP_MAP_DETACH:
break;
default:
gomp_mutex_unlock (&devicep->lock);
@ -2621,6 +2637,14 @@ GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs,
&kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
i += j - i - 1;
}
else if (i + 1 < mapnum && (kinds[i + 1] & 0xff) == GOMP_MAP_ATTACH)
{
/* An attach operation must be processed together with the mapped
base-pointer list item. */
gomp_map_vars (devicep, 2, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
true, GOMP_MAP_VARS_ENTER_DATA);
i += 1;
}
else
gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
true, GOMP_MAP_VARS_ENTER_DATA);

82
libgomp/testsuite/libgomp.c-c++-common/ptr-attach-1.c Normal file
View File

@ -0,0 +1,82 @@
#include <stdlib.h>
struct S
{
int a, b;
int *ptr;
int c, d;
};
typedef struct S S;
#pragma omp declare target
int *gp;
#pragma omp end declare target
#define N 10
int main (void)
{
/* Test to see if pointer attachment works, for scalar pointers,
and pointer fields in structures. */
int *ptr = (int *) malloc (sizeof (int) * N);
int *orig_ptr = ptr;
#pragma omp target map (ptr, ptr[:N])
{
for (int i = 0; i < N; i++)
ptr[i] = N - i;
}
if (ptr != orig_ptr)
abort ();
for (int i = 0; i < N; i++)
if (ptr[i] != N - i)
abort ();
S s = { 0 };
s.ptr = ptr;
#pragma omp target map (s, s.ptr[:N])
{
for (int i = 0; i < N; i++)
s.ptr[i] = i;
s.a = 1;
s.b = 2;
}
if (s.ptr != ptr)
abort ();
for (int i = 0; i < N; i++)
if (s.ptr[i] != i)
abort ();
if (s.a != 1 || s.b != 2 || s.c != 0 || s.d != 0)
abort ();
gp = (int *) malloc (sizeof (int) * N);
orig_ptr = gp;
for (int i = 0; i < N; i++)
gp[i] = i - 1;
#pragma omp target map (gp[:N])
{
for (int i = 0; i < N; i++)
gp[i] += 1;
}
if (gp != orig_ptr)
abort ();
for (int i = 0; i < N; i++)
if (gp[i] != i)
abort ();
free (ptr);
free (gp);
return 0;
}