mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-05 17:00:58 +08:00
Code Issues Packages Projects Releases Wiki Activity

OpenMP 5.0: Remove array section base-pointer mapping semantics and other front-end adjustments

Browse Source 
This patch implements three pieces of functionality:

(1) Adjust array section mapping to have standards conforming behavior,
mapping array sections should *NOT* also map the base-pointer:

struct S { int *ptr; ... };
struct S s;

Instead of generating this during gimplify:
                              map(to:*_1 [len: 400]) map(attach:s.ptr [bias: 0])

Now, adjust to:

(i.e. do not map the base-pointer together. The attach operation is still
generated, and if s.ptr is already mapped prior, attachment will happen)

The correct way of achieving the base-pointer-also-mapped behavior would be to
use:

(A small Fortran front-end patch to trans-openmp.c:gfc_trans_omp_array_section
 is also included, which removes generation of a GOMP_MAP_ALWAYS_POINTER for
 array types, which appears incorrect and causes a regression in
 libgomp.fortranlibgomp.fortran/struct-elem-map-1.f90)

(2) Related to the first item above, are fixes in libgomp/target.c to not
overwrite attached pointers when handling device<->host copies, mainly for the
"always" case.

(3) The third is a set of changes to the C/C++ front-ends to extend the allowed
component access syntax in map clauses. These changes are enabled for both
OpenACC and OpenMP.

gcc/c/ChangeLog:

	* c-parser.c (struct omp_dim): New struct type for use inside
	c_parser_omp_variable_list.
	(c_parser_omp_variable_list): Allow multiple levels of array and
	component accesses in array section base-pointer expression.
	(c_parser_omp_clause_to): Set 'allow_deref' to true in call to
	c_parser_omp_var_list_parens.
	(c_parser_omp_clause_from): Likewise.
	* c-typeck.c (handle_omp_array_sections_1): Extend allowed range
	of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
	POINTER_PLUS_EXPR.
	(c_finish_omp_clauses): Extend allowed ranged of expressions
	involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.

gcc/cp/ChangeLog:

	* parser.c (struct omp_dim): New struct type for use inside
	cp_parser_omp_var_list_no_open.
	(cp_parser_omp_var_list_no_open): Allow multiple levels of array and
	component accesses in array section base-pointer expression.
	(cp_parser_omp_all_clauses): Set 'allow_deref' to true in call to
	cp_parser_omp_var_list for to/from clauses.
	* semantics.c (handle_omp_array_sections_1): Extend allowed range
	of base-pointer expressions involving INDIRECT/MEM/ARRAY_REF and
	POINTER_PLUS_EXPR.
	(handle_omp_array_sections): Adjust pointer map generation of
	references.
	(finish_omp_clauses): Extend allowed ranged of expressions
	involving INDIRECT/MEM/ARRAY_REF and POINTER_PLUS_EXPR.

gcc/fortran/ChangeLog:

	* trans-openmp.c (gfc_trans_omp_array_section): Do not generate
	GOMP_MAP_ALWAYS_POINTER map for main array maps of ARRAY_TYPE type.

gcc/ChangeLog:

	* gimplify.c (extract_base_bit_offset): Add 'tree *offsetp' parameter,
	accomodate case where 'offset' return of get_inner_reference is
	non-NULL.
	(is_or_contains_p): Further robustify conditions.
	(omp_target_reorder_clauses): In alloc/to/from sorting phase, also
	move following GOMP_MAP_ALWAYS_POINTER maps along.  Add new sorting
	phase where we make sure pointers with an attach/detach map are ordered
	correctly.
	(gimplify_scan_omp_clauses): Add modifications to avoid creating
	GOMP_MAP_STRUCT and associated alloc map for attach/detach maps.

gcc/testsuite/ChangeLog:

	* c-c++-common/goacc/deep-copy-arrayofstruct.c: Adjust testcase.
	* c-c++-common/gomp/target-enter-data-1.c: New testcase.
	* c-c++-common/gomp/target-implicit-map-2.c: New testcase.

libgomp/ChangeLog:

	* target.c (gomp_map_vars_existing): Make sure attached pointer is
	not overwritten during cross-host/device copying.
	(gomp_update): Likewise.
	(gomp_exit_data): Likewise.
	* testsuite/libgomp.c++/target-11.C: Adjust testcase.
	* testsuite/libgomp.c++/target-12.C: Likewise.
	* testsuite/libgomp.c++/target-15.C: Likewise.
	* testsuite/libgomp.c++/target-16.C: Likewise.
	* testsuite/libgomp.c++/target-17.C: Likewise.
	* testsuite/libgomp.c++/target-21.C: Likewise.
	* testsuite/libgomp.c++/target-23.C: Likewise.
	* testsuite/libgomp.c/target-23.c: Likewise.
	* testsuite/libgomp.c/target-29.c: Likewise.
	* testsuite/libgomp.c-c++-common/target-implicit-map-2.c: New testcase.
This commit is contained in:
Chung-Lin Tang 2021-12-08 23:58:55 +08:00
parent 6b49d50a27
commit 6c0399378e
20 changed files with 751 additions and 155 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -12989,19 +12989,29 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
The optional ALLOW_DEREF argument is true if list items can use the deref
(->) operator. */
struct omp_dim
{
tree low_bound, length;
location_t loc;
bool no_colon;
omp_dim (tree lb, tree len, location_t lo, bool nc)
: low_bound (lb), length (len), loc (lo), no_colon (nc) {}
};
static tree
c_parser_omp_variable_list (c_parser *parser,
location_t clause_loc,
enum omp_clause_code kind, tree list,
bool allow_deref = false)
{
auto_vec<omp_dim> dims;
bool array_section_p;
auto_vec<c_token> tokens;
unsigned int tokens_avail = 0;
bool first = true;
while (1)
{
bool array_section_p = false;
if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
{
if (c_parser_next_token_is_not (parser, CPP_NAME)
@ -13120,6 +13130,7 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
start_component_ref:
while (c_parser_next_token_is (parser, CPP_DOT)
|| (allow_deref
&& c_parser_next_token_is (parser, CPP_DEREF)))
@ -13147,9 +13158,13 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
array_section_p = false;
dims.truncate (0);
while (c_parser_next_token_is (parser, CPP_OPEN_SQUARE))
{
location_t loc = UNKNOWN_LOCATION;
tree low_bound = NULL_TREE, length = NULL_TREE;
bool no_colon = false;
c_parser_consume_token (parser);
if (!c_parser_next_token_is (parser, CPP_COLON))
@ -13160,9 +13175,13 @@ c_parser_omp_variable_list (c_parser *parser,
expr = convert_lvalue_to_rvalue (expr_loc, expr,
false, true);
low_bound = expr.value;
loc = expr_loc;
}
if (c_parser_next_token_is (parser, CPP_CLOSE_SQUARE))
length = integer_one_node;
{
length = integer_one_node;
no_colon = true;
}
else
{
/* Look for `:'. */
@ -13191,8 +13210,33 @@ c_parser_omp_variable_list (c_parser *parser,
break;
}
t = tree_cons (low_bound, length, t);
dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
}
if (t != error_mark_node)
{
if ((kind == OMP_CLAUSE_MAP
|| kind == OMP_CLAUSE_FROM
|| kind == OMP_CLAUSE_TO)
&& !array_section_p
&& (c_parser_next_token_is (parser, CPP_DOT)
|| (allow_deref
&& c_parser_next_token_is (parser,
CPP_DEREF))))
{
for (unsigned i = 0; i < dims.length (); i++)
{
gcc_assert (dims[i].length == integer_one_node);
t = build_array_ref (dims[i].loc,
t, dims[i].low_bound);
}
goto start_component_ref;
}
else
for (unsigned i = 0; i < dims.length (); i++)
t = tree_cons (dims[i].low_bound, dims[i].length, t);
}
if ((kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
&& t != error_mark_node
&& parser->tokens_avail != 2)
@ -16439,7 +16483,7 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
static tree
c_parser_omp_clause_to (c_parser *parser, tree list)
{
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list);
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list, true);
}
/* OpenMP 4.0:
@ -16448,7 +16492,7 @@ c_parser_omp_clause_to (c_parser *parser, tree list)
static tree
c_parser_omp_clause_from (c_parser *parser, tree list)
{
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list);
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list, true);
}
/* OpenMP 4.0:

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

@ -13220,6 +13220,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
while (TREE_CODE (t) == INDIRECT_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
@ -13241,10 +13253,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF)
while (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
{
@ -13533,15 +13549,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
array-section-subscript, the array section can't be contiguous. */
array-section-subscript, the array section could be non-contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
error_at (OMP_CLAUSE_LOCATION (c),
"array section is not contiguous in %qs clause",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
/* If any prior dimension has a non-one length, then deem this
array section as non-contiguous. */
for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST;
d = TREE_CHAIN (d))
{
tree d_length = TREE_VALUE (d);
if (d_length == NULL_TREE || !integer_onep (d_length))
{
error_at (OMP_CLAUSE_LOCATION (c),
"array section is not contiguous in %qs clause",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
}
}
}
else
@ -14890,13 +14916,20 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
{
while (TREE_CODE (t) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF)
do
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
while (TREE_CODE (t) == COMPONENT_REF);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
@ -14963,14 +14996,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bias) to zero here, so it is not set erroneously to the pointer
size later on in gimplify.c. */
OMP_CLAUSE_SIZE (c) = size_zero_node;
while (TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
indir_component_ref_p = false;
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
&& (TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF
|| TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
|| TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
@ -15006,7 +15057,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
t = TREE_OPERAND (t, 0);
if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
if (TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
@ -15015,6 +15066,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
if (remove)
break;
@ -15086,7 +15146,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in data clauses", t);
remove = true;
}
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 (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),

54
gcc/cp/parser.c
View File

@ -36406,11 +36406,22 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
The optional ALLOW_DEREF argument is true if list items can use the deref
(->) operator. */
struct omp_dim
{
tree low_bound, length;
location_t loc;
bool no_colon;
omp_dim (tree lb, tree len, location_t lo, bool nc)
: low_bound (lb), length (len), loc (lo), no_colon (nc) {}
};
static tree
cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
tree list, bool *colon,
bool allow_deref = false)
{
auto_vec<omp_dim> dims;
bool array_section_p;
cp_token *token;
bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
if (colon)
@ -36491,6 +36502,7 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
start_component_ref:
while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
|| (allow_deref
&& cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
@ -36514,14 +36526,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
array_section_p = false;
dims.truncate (0);
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
{
location_t loc = UNKNOWN_LOCATION;
tree low_bound = NULL_TREE, length = NULL_TREE;
bool no_colon = false;
parser->colon_corrects_to_scope_p = false;
cp_lexer_consume_token (parser->lexer);
if (!cp_lexer_next_token_is (parser->lexer, CPP_COLON))
{
loc = cp_lexer_peek_token (parser->lexer)->location;
low_bound = cp_parser_expression (parser);
/* Later handling is not prepared to see through these. */
gcc_checking_assert (!location_wrapper_p (low_bound));
@ -36530,7 +36547,10 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
parser->colon_corrects_to_scope_p
= saved_colon_corrects_to_scope_p;
if (cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_SQUARE))
length = integer_one_node;
{
length = integer_one_node;
no_colon = true;
}
else
{
/* Look for `:'. */
@ -36543,6 +36563,8 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
}
if (kind == OMP_CLAUSE_DEPEND || kind == OMP_CLAUSE_AFFINITY)
cp_parser_commit_to_tentative_parse (parser);
else
array_section_p = true;
if (!cp_lexer_next_token_is (parser->lexer,
CPP_CLOSE_SQUARE))
{
@ -36561,8 +36583,30 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
goto skip_comma;
}
decl = tree_cons (low_bound, length, decl);
dims.safe_push (omp_dim (low_bound, length, loc, no_colon));
}
if ((kind == OMP_CLAUSE_MAP
|| kind == OMP_CLAUSE_FROM
|| kind == OMP_CLAUSE_TO)
&& !array_section_p
&& (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
|| (allow_deref
&& cp_lexer_next_token_is (parser->lexer,
CPP_DEREF))))
{
for (unsigned i = 0; i < dims.length (); i++)
{
gcc_assert (dims[i].length == integer_one_node);
decl = build_array_ref (dims[i].loc,
decl, dims[i].low_bound);
}
goto start_component_ref;
}
else
for (unsigned i = 0; i < dims.length (); i++)
decl = tree_cons (dims[i].low_bound, dims[i].length, decl);
break;
default:
break;
@ -40064,11 +40108,13 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO_DECLARE,
clauses);
else
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses);
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
true);
c_name = "to";
break;
case PRAGMA_OMP_CLAUSE_FROM:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses);
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
true);
c_name = "from";
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:

144
gcc/cp/semantics.c
View File

@ -5025,6 +5025,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
ret = t;
while (TREE_CODE (t) == INDIRECT_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
if (TREE_CODE (t) == COMPONENT_REF
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
@ -5049,10 +5061,14 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == INDIRECT_REF)
while (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
if (REFERENCE_REF_P (t))
@ -5336,15 +5352,25 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
/* If there is a pointer type anywhere but in the very first
array-section-subscript, the array section can't be contiguous. */
array-section-subscript, the array section could be non-contiguous. */
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_AFFINITY
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
&& TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
{
error_at (OMP_CLAUSE_LOCATION (c),
"array section is not contiguous in %qs clause",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
/* If any prior dimension has a non-one length, then deem this
array section as non-contiguous. */
for (tree d = TREE_CHAIN (t); TREE_CODE (d) == TREE_LIST;
d = TREE_CHAIN (d))
{
tree d_length = TREE_VALUE (d);
if (d_length == NULL_TREE || !integer_onep (d_length))
{
error_at (OMP_CLAUSE_LOCATION (c),
"array section is not contiguous in %qs clause",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
}
}
}
else
@ -5615,16 +5641,37 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
default:
break;
}
bool reference_always_pointer = true;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
{
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
&& TYPE_REF_P (TREE_TYPE (t)))
{
if (TREE_CODE (TREE_TYPE (TREE_TYPE (t))) == ARRAY_TYPE)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
else
t = convert_from_reference (t);
reference_always_pointer = false;
}
}
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
: GOMP_MAP_ALWAYS_POINTER;
gomp_map_kind k;
if ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
&& TREE_CODE (TREE_TYPE (t)) == POINTER_TYPE)
k = GOMP_MAP_ATTACH_DETACH;
else
{
t = TREE_OPERAND (t, 0);
k = (ort == C_ORT_ACC
? GOMP_MAP_ATTACH_DETACH : GOMP_MAP_ALWAYS_POINTER);
}
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
@ -5648,8 +5695,10 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
OMP_CLAUSE_SIZE (c2) = t;
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
ptr = OMP_CLAUSE_DECL (c2);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
if (reference_always_pointer
&& OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& TYPE_REF_P (TREE_TYPE (ptr))
&& INDIRECT_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
{
@ -7850,15 +7899,22 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
{
while (TREE_CODE (t) == COMPONENT_REF)
t = TREE_OPERAND (t, 0);
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == INDIRECT_REF)
do
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
while (TREE_CODE (t) == COMPONENT_REF);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IMPLICIT (c)
&& (bitmap_bit_p (&map_head, DECL_UID (t))
@ -7929,15 +7985,33 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH)
OMP_CLAUSE_DECL (c) = t;
}
while (TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == COMPOUND_EXPR)
{
t = TREE_OPERAND (t, 1);
STRIP_NOPS (t);
}
indir_component_ref_p = false;
if (TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
&& (TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF
|| TREE_CODE (TREE_OPERAND (t, 0)) == ARRAY_REF))
{
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
indir_component_ref_p = true;
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
if (TREE_CODE (t) == COMPONENT_REF
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
@ -7972,6 +8046,24 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
t = TREE_OPERAND (t, 0);
if (TREE_CODE (t) == MEM_REF)
{
if (maybe_ne (mem_ref_offset (t), 0))
error_at (OMP_CLAUSE_LOCATION (c),
"cannot dereference %qE in %qs clause", t,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
else
t = TREE_OPERAND (t, 0);
}
while (TREE_CODE (t) == MEM_REF
|| TREE_CODE (t) == INDIRECT_REF
|| TREE_CODE (t) == ARRAY_REF)
{
t = TREE_OPERAND (t, 0);
STRIP_NOPS (t);
if (TREE_CODE (t) == POINTER_PLUS_EXPR)
t = TREE_OPERAND (t, 0);
}
}
if (remove)
break;
@ -8069,7 +8161,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
"%qD appears more than once in data clauses", t);
remove = true;
}
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 (ort == C_ORT_ACC)
error_at (OMP_CLAUSE_LOCATION (c),
@ -8116,8 +8209,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
if (t != OMP_CLAUSE_DECL (c)
&& TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
tree decl = OMP_CLAUSE_DECL (c);
if (t != decl
&& (TREE_CODE (decl) == COMPONENT_REF
|| (INDIRECT_REF_P (decl)
&& TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
&& TYPE_REF_P (TREE_TYPE (TREE_OPERAND (decl, 0))))))
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
handle_map_references:
@ -8146,7 +8244,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);

3
gcc/fortran/trans-openmp.c
View File

@ -2460,6 +2460,9 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
TREE_TYPE (TREE_TYPE (decl)),
decl, offset, NULL_TREE, NULL_TREE);
OMP_CLAUSE_DECL (node) = offset;
if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
return;
}
else
{

220
gcc/gimplify.c
View File

@ -8660,7 +8660,7 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
static tree
extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
poly_offset_int *poffsetp)
poly_offset_int *poffsetp, tree *offsetp)
{
tree offset;
poly_int64 bitsize, bitpos;
@ -8707,10 +8707,11 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
&& TREE_CODE (TREE_TYPE (TREE_OPERAND (base, 0))) == REFERENCE_TYPE)
base = TREE_OPERAND (base, 0);
gcc_assert (offset == NULL_TREE || poly_int_tree_p (offset));
if (offset)
poffset = wi::to_poly_offset (offset);
if (offset && poly_int_tree_p (offset))
{
poffset = wi::to_poly_offset (offset);
offset = NULL_TREE;
}
else
poffset = 0;
@ -8719,6 +8720,7 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
*bitposp = bitpos;
*poffsetp = poffset;
*offsetp = offset;
/* Set *BASE_REF if BASE was a dereferenced reference variable. */
if (base_ref && orig_base != base)
@ -8732,12 +8734,22 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
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;
if ((TREE_CODE (expr) == INDIRECT_REF && TREE_CODE (base_ptr) == MEM_REF)
|| (TREE_CODE (expr) == MEM_REF && TREE_CODE (base_ptr) == INDIRECT_REF))
return operand_equal_p (TREE_OPERAND (expr, 0),
TREE_OPERAND (base_ptr, 0));
while (!operand_equal_p (expr, base_ptr))
{
if (TREE_CODE (base_ptr) == COMPOUND_EXPR)
base_ptr = TREE_OPERAND (base_ptr, 1);
if (TREE_CODE (base_ptr) == COMPONENT_REF
|| TREE_CODE (base_ptr) == POINTER_PLUS_EXPR
|| TREE_CODE (base_ptr) == SAVE_EXPR)
base_ptr = TREE_OPERAND (base_ptr, 0);
else
break;
}
return operand_equal_p (expr, base_ptr);
}
/* Implement OpenMP 5.x map ordering rules for target directives. There are
@ -8817,21 +8829,107 @@ omp_target_reorder_clauses (tree *list_p)
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;
if (atf[j])
{
tree *cp2 = atf[j];
tree decl2 = OMP_CLAUSE_DECL (*cp2);
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;
if (*cp2 != NULL_TREE
&& OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (*cp2) == GOMP_MAP_ALWAYS_POINTER)
{
tree c2 = *cp2;
*cp2 = OMP_CLAUSE_CHAIN (c2);
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
}
atf[j] = NULL;
}
}
}
}
}
/* For attach_detach map clauses, if there is another map that maps the
attached/detached pointer, make sure that map is ordered before the
attach_detach. */
atf.truncate (0);
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 clauses, and
always_pointer/attach_detach clauses. */
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
|| k == GOMP_MAP_ATTACH_DETACH
|| k == GOMP_MAP_ALWAYS_POINTER)
atf.safe_push (cp);
}
for (unsigned int i = 0; i < atf.length (); i++)
if (atf[i])
{
tree *cp = atf[i];
tree ptr = OMP_CLAUSE_DECL (*cp);
STRIP_TYPE_NOPS (ptr);
if (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH)
for (unsigned int j = i + 1; j < atf.length (); j++)
{
tree *cp2 = atf[j];
tree decl2 = OMP_CLAUSE_DECL (*cp2);
if (OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ATTACH_DETACH
&& OMP_CLAUSE_MAP_KIND (*cp2) != GOMP_MAP_ALWAYS_POINTER
&& is_or_contains_p (decl2, ptr))
{
/* Move *cp2 to before *cp. */
tree c = *cp2;
*cp2 = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *cp;
*cp = c;
atf[j] = NULL;
/* If decl2 is of the form '*decl2_opnd0', and followed by an
ALWAYS_POINTER or ATTACH_DETACH of 'decl2_opnd0', move the
pointer operation along with *cp2. This can happen for C++
reference sequences. */
if (j + 1 < atf.length ()
&& (TREE_CODE (decl2) == INDIRECT_REF
|| TREE_CODE (decl2) == MEM_REF))
{
tree *cp3 = atf[j + 1];
tree decl3 = OMP_CLAUSE_DECL (*cp3);
tree decl2_opnd0 = TREE_OPERAND (decl2, 0);
if ((OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ALWAYS_POINTER
|| OMP_CLAUSE_MAP_KIND (*cp3) == GOMP_MAP_ATTACH_DETACH)
&& operand_equal_p (decl3, decl2_opnd0))
{
/* Also move *cp3 to before *cp. */
c = *cp3;
*cp2 = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *cp;
*cp = c;
atf[j + 1] = NULL;
j += 1;
}
}
}
}
}
}
/* DECL is supposed to have lastprivate semantics in the outer contexts
@ -8923,6 +9021,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
hash_map<tree_operand_hash, tree *> *struct_seen_clause = NULL;
hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
@ -9398,6 +9497,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
bool indir_p = false;
bool component_ref_p = false;
tree indir_base = NULL_TREE;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@ -9416,6 +9516,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
== POINTER_TYPE))
{
indir_p = true;
indir_base = decl;
decl = TREE_OPERAND (decl, 0);
STRIP_NOPS (decl);
}
@ -9462,7 +9563,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
!= GOMP_MAP_POINTER)
|| OMP_CLAUSE_DECL (next_clause) != decl)
&& (!struct_deref_set
|| !struct_deref_set->contains (decl)))
|| !struct_deref_set->contains (decl))
&& (!struct_map_to_clause
|| !struct_map_to_clause->get (indir_base)))
{
if (!struct_deref_set)
struct_deref_set = new hash_set<tree> ();
@ -9506,7 +9609,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
if ((DECL_P (decl)
|| (component_ref_p
&& (INDIRECT_REF_P (decl)
|| TREE_CODE (decl) == MEM_REF)))
|| TREE_CODE (decl) == MEM_REF
|| TREE_CODE (decl) == ARRAY_REF)))
&& 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
@ -9541,7 +9645,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
}
if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
/* The below prev_list_p based error recovery code is
currently no longer valid for OpenMP. */
if (code != OMP_TARGET
&& code != OMP_TARGET_DATA
&& code != OMP_TARGET_UPDATE
&& code != OMP_TARGET_ENTER_DATA
&& code != OMP_TARGET_EXIT_DATA
&& OMP_CLAUSE_CHAIN (*prev_list_p) != c)
{
tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
@ -9554,13 +9666,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
poly_offset_int offset1;
poly_int64 bitpos1;
tree tree_offset1;
tree base_ref;
tree base
= extract_base_bit_offset (OMP_CLAUSE_DECL (c), &base_ref,
&bitpos1, &offset1);
&bitpos1, &offset1,
&tree_offset1);
gcc_assert (base == decl);
bool do_map_struct = (base == decl && !tree_offset1);
splay_tree_node n
= (DECL_P (decl)
@ -9592,6 +9706,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
/* We currently don't handle non-constant offset accesses wrt to
GOMP_MAP_STRUCT elements. */
if (!do_map_struct)
goto skip_map_struct;
/* Nor for attach_detach for OpenMP. */
if ((code == OMP_TARGET
|| code == OMP_TARGET_DATA
|| code == OMP_TARGET_UPDATE
|| code == OMP_TARGET_ENTER_DATA
|| code == OMP_TARGET_EXIT_DATA)
&& attach_detach)
{
if (DECL_P (decl))
{
if (struct_seen_clause == NULL)
struct_seen_clause
= new hash_map<tree_operand_hash, tree *>;
if (!struct_seen_clause->get (decl))
struct_seen_clause->put (decl, list_p);
}
goto skip_map_struct;
}
if ((DECL_P (decl)
&& (n == NULL || (n->value & GOVD_MAP) == 0))
|| (!DECL_P (decl)
@ -9631,9 +9771,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
struct_map_to_clause->put (decl, l);
if (ptr || attach_detach)
{
insert_struct_comp_map (code, c, l, *prev_list_p,
tree **sc = (struct_seen_clause
? struct_seen_clause->get (decl)
: NULL);
tree *insert_node_pos = sc ? *sc : prev_list_p;
insert_struct_comp_map (code, c, l, *insert_node_pos,
NULL);
*prev_list_p = l;
*insert_node_pos = l;
prev_list_p = NULL;
}
else
@ -9719,9 +9864,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree sc_decl = OMP_CLAUSE_DECL (*sc);
poly_offset_int offsetn;
poly_int64 bitposn;
tree tree_offsetn;
tree base
= extract_base_bit_offset (sc_decl, NULL,
&bitposn, &offsetn);
&bitposn, &offsetn,
&tree_offsetn);
if (base != decl)
break;
if (scp)
@ -9809,16 +9956,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
continue;
}
}
skip_map_struct:
;
}
else if ((code == OACC_ENTER_DATA
|| code == OACC_EXIT_DATA
|| code == OACC_DATA
|| code == OACC_PARALLEL
|| code == OACC_KERNELS
|| code == OACC_SERIAL)
|| code == OACC_SERIAL
|| code == OMP_TARGET_ENTER_DATA
|| code == OMP_TARGET_EXIT_DATA)
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
gomp_map_kind k = (code == OACC_EXIT_DATA
gomp_map_kind k = ((code == OACC_EXIT_DATA
|| code == OMP_TARGET_EXIT_DATA)
? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
OMP_CLAUSE_SET_MAP_KIND (c, k);
}
@ -10650,6 +10802,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
ctx->clauses = *orig_list_p;
gimplify_omp_ctxp = ctx;
if (struct_seen_clause)
delete struct_seen_clause;
if (struct_map_to_clause)
delete struct_map_to_clause;
if (struct_deref_set)

5
gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
View File

@ -37,13 +37,12 @@ int main(int argc, char* argv[])
{
int j, k;
for (k = 0; k < S; k++)
#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */
#pragma acc parallel loop copy(m[k].a[0:N])
for (j = 0; j < N; j++)
m[k].a[j]++;
for (k = 0; k < S; k++)
#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected .\\\). before .\\\.. token" } */
/* { dg-error ".m. appears more than once in data clauses" "" { target c++ } .-1 } */
#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10])
for (j = 0; j < N; j++)
{
m[k].b[j]++;

24
gcc/testsuite/c-c++-common/gomp/target-enter-data-1.c Normal file
View File

@ -0,0 +1,24 @@
/* { dg-do compile } */
/* { dg-additional-options "-fopenmp -fdump-tree-gimple" } */
struct bar
{
int num_vectors;
double *vectors;
};
struct foo
{
int num_vectors;
struct bar *bars;
double **vectors;
};
void func (struct foo *f, int n, int m)
{
#pragma omp target enter data map (to: f->vectors[m][:n])
#pragma omp target enter data map (to: f->bars[n].vectors[:m])
#pragma omp target enter data map (to: f->bars[n].vectors[:f->bars[n].num_vectors])
}
/* { dg-final { scan-tree-dump-times "map\\(to:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:\[^-\]+->vectors \\\[bias: \[^\]\]+\\\]\\)" 3 "gimple" } } */

52
gcc/testsuite/c-c++-common/gomp/target-implicit-map-2.c Normal file
View File

@ -0,0 +1,52 @@
/* { dg-do compile } */
/* { dg-additional-options "-fdump-tree-gimple" } */
#include <stdlib.h>
#define N 10
struct S
{
int a, b;
int *ptr;
int c, d;
};
int
main (void)
{
struct S a;
a.ptr = (int *) malloc (sizeof (int) * N);
for (int i = 0; i < N; i++)
a.ptr[i] = 0;
#pragma omp target enter data map(to: a.ptr, a.ptr[:N])
#pragma omp target
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target update from(a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 1)
abort ();
#pragma omp target map(a.ptr[:N])
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target update from(a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 2)
abort ();
#pragma omp target exit data map(from:a.ptr, a.ptr[:N])
return 0;
}
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:a \[len: [0-9]+\]\[implicit\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(attach:a\.ptr \[bias: 0\]\)} "gimple" } } */

107
libgomp/target.c
View File

@ -581,11 +581,30 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
address/length adjustment is a TODO. */
assert (!implicit_subset);
gomp_copy_host2dev (devicep, aq,
(void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ newn->host_start - oldn->host_start),
(void *) newn->host_start,
newn->host_end - newn->host_start, false, cbuf);
if (oldn->aux && oldn->aux->attach_count)
{
/* We have to be careful not to overwrite still attached pointers
during the copyback to host. */
uintptr_t addr = newn->host_start;
while (addr < newn->host_end)
{
size_t i = (addr - oldn->host_start) / sizeof (void *);
if (oldn->aux->attach_count[i] == 0)
gomp_copy_host2dev (devicep, aq,
(void *) (oldn->tgt->tgt_start
+ oldn->tgt_offset
+ addr - oldn->host_start),
(void *) addr,
sizeof (void *), false, cbuf);
addr += sizeof (void *);
}
}
else
gomp_copy_host2dev (devicep, aq,
(void *) (oldn->tgt->tgt_start + oldn->tgt_offset
+ newn->host_start - oldn->host_start),
(void *) newn->host_start,
newn->host_end - newn->host_start, false, cbuf);
}
gomp_increment_refcount (oldn, refcount_set);
@ -2009,17 +2028,45 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
(void *) n->host_end);
}
if (n->aux && n->aux->attach_count)
{
uintptr_t addr = cur_node.host_start;
while (addr < cur_node.host_end)
{
/* We have to be careful not to overwrite still attached
pointers during host<->device updates. */
size_t i = (addr - cur_node.host_start) / sizeof (void *);
if (n->aux->attach_count[i] == 0)
{
void *devaddr = (void *) (n->tgt->tgt_start
+ n->tgt_offset
+ addr - n->host_start);
if (GOMP_MAP_COPY_TO_P (kind & typemask))
gomp_copy_host2dev (devicep, NULL,
devaddr, (void *) addr,
sizeof (void *), false, NULL);
if (GOMP_MAP_COPY_FROM_P (kind & typemask))
gomp_copy_dev2host (devicep, NULL,
(void *) addr, devaddr,
sizeof (void *));
}
addr += sizeof (void *);
}
}
else
{
void *hostaddr = (void *) cur_node.host_start;
void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start
- n->host_start);
size_t size = cur_node.host_end - cur_node.host_start;
void *hostaddr = (void *) cur_node.host_start;
void *devaddr = (void *) (n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start);
size_t size = cur_node.host_end - cur_node.host_start;
if (GOMP_MAP_COPY_TO_P (kind & typemask))
gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
false, NULL);
if (GOMP_MAP_COPY_FROM_P (kind & typemask))
gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
if (GOMP_MAP_COPY_TO_P (kind & typemask))
gomp_copy_host2dev (devicep, NULL, devaddr, hostaddr, size,
false, NULL);
if (GOMP_MAP_COPY_FROM_P (kind & typemask))
gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
}
}
}
gomp_mutex_unlock (&devicep->lock);
@ -2932,11 +2979,31 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
if ((kind == GOMP_MAP_FROM && do_copy)
|| kind == GOMP_MAP_ALWAYS_FROM)
gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
(void *) (k->tgt->tgt_start + k->tgt_offset
+ cur_node.host_start
- k->host_start),
cur_node.host_end - cur_node.host_start);
{
if (k->aux && k->aux->attach_count)
{
/* We have to be careful not to overwrite still attached
pointers during the copyback to host. */
uintptr_t addr = k->host_start;
while (addr < k->host_end)
{
size_t i = (addr - k->host_start) / sizeof (void *);
if (k->aux->attach_count[i] == 0)
gomp_copy_dev2host (devicep, NULL, (void *) addr,
(void *) (k->tgt->tgt_start
+ k->tgt_offset
+ addr - k->host_start),
sizeof (void *));
addr += sizeof (void *);
}
}
else
gomp_copy_dev2host (devicep, NULL, (void *) cur_node.host_start,
(void *) (k->tgt->tgt_start + k->tgt_offset
+ cur_node.host_start
- k->host_start),
cur_node.host_end - cur_node.host_start);
}
/* Structure elements lists are removed altogether at once, which
may cause immediate deallocation of the target_mem_desc, causing

14
libgomp/testsuite/libgomp.c++/target-11.C
View File

@ -23,9 +23,11 @@ foo ()
e = c + 18;
D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
int err = 0;
#pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \
map (tofrom:s.s[3:3], s. template v. template d[z + 1:z + 3]) \
map (from: s.w[z:4], s.x[1:3], err) private (i)
#pragma omp target map (to: s.v.b, s.v.b[0:z + 7]) \
map (s.template u, s.template u[z + 1:z + 4]) \
map (tofrom: s.s, s.s[3:3]) \
map (tofrom: s. template v. template d[z + 1:z + 3])\
map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
{
err = 0;
for (i = 0; i < 7; i++)
@ -80,9 +82,9 @@ main ()
e = c + 18;
S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
int err = 0;
#pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
map (from: s.w[z:4], s.x[1:3], err) private (i)
#pragma omp target map (to: s.v.b, s.v.b[0:z + 7], s.u, s.u[z + 1:z + 4]) \
map (tofrom: s.s, s.s[3:3], s.v.d[z + 1:z + 3]) \
map (from: s.w, s.w[z:4], s.x, s.x[1:3], err) private (i)
{
err = 0;
for (i = 0; i < 7; i++)

2
libgomp/testsuite/libgomp.c++/target-12.C
View File

@ -53,7 +53,7 @@ main ()
int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
int *v = u + 4;
#pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
#pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
s.s++;
u[3]++;
s.v[1]++;

20
libgomp/testsuite/libgomp.c++/target-15.C
View File

@ -14,7 +14,7 @@ foo (S s)
d = id;
int err;
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@ -48,7 +48,7 @@ foo (S s)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@ -61,8 +61,8 @@ foo (S s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@ -73,7 +73,7 @@ foo (S s)
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@ -97,7 +97,7 @@ foo (S s)
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@ -109,8 +109,8 @@ foo (S s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@ -121,7 +121,7 @@ foo (S s)
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@ -133,7 +133,7 @@ foo (S s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)

20
libgomp/testsuite/libgomp.c++/target-16.C
View File

@ -16,7 +16,7 @@ foo (S<C, I, L, UC, SH> s)
d = id;
int err;
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@ -50,7 +50,7 @@ foo (S<C, I, L, UC, SH> s)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@ -63,8 +63,8 @@ foo (S<C, I, L, UC, SH> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@ -75,7 +75,7 @@ foo (S<C, I, L, UC, SH> s)
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@ -99,7 +99,7 @@ foo (S<C, I, L, UC, SH> s)
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@ -111,8 +111,8 @@ foo (S<C, I, L, UC, SH> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@ -123,7 +123,7 @@ foo (S<C, I, L, UC, SH> s)
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@ -135,7 +135,7 @@ foo (S<C, I, L, UC, SH> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)

20
libgomp/testsuite/libgomp.c++/target-17.C
View File

@ -16,7 +16,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
d = id;
int err;
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@ -50,7 +50,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| omp_target_is_present (&s.h, d)
|| omp_target_is_present (&s.h[2], d)))
abort ();
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@ -63,8 +63,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
@ -75,7 +75,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
}
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@ -99,7 +99,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@ -111,8 +111,8 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
@ -123,7 +123,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
}
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
@ -135,7 +135,7 @@ foo (S<C, I, L, UCR, CAR, SH, IPR> s)
|| !omp_target_is_present (&s.h, d)
|| !omp_target_is_present (&s.h[2], d))
abort ();
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3], s.e, s.f, s.g[1:2], s.h, s.h[2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)

8
libgomp/testsuite/libgomp.c++/target-21.C
View File

@ -7,7 +7,7 @@ void
foo (S s)
{
int err;
#pragma omp target map (s.x[0:N], s.y[0:N]) map (s.t.t[16:3]) map (from: err)
#pragma omp target map (s.x[0:N], s.y, s.y[0:N]) map (s.t.t[16:3]) map (from: err)
{
err = s.x[2] != 28 || s.y[2] != 37 || s.t.t[17] != 81;
s.x[2]++;
@ -38,7 +38,7 @@ void
foo2 (S &s)
{
int err;
#pragma omp target map (s.x[N:10], s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
#pragma omp target map (s.x[N:10], s.y, s.y[N:10]) map (from: err) map (s.t.t[N+16:N+3])
{
err = s.x[2] != 30 || s.y[2] != 38 || s.t.t[17] != 81;
s.x[2]++;
@ -69,7 +69,7 @@ void
foo3 (U s)
{
int err;
#pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
#pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
{
err = s.x[2] != 32 || s.y[2] != 39 || s.t.t[17] != 82;
s.x[2]++;
@ -100,7 +100,7 @@ void
foo4 (U &s)
{
int err;
#pragma omp target map (s.x[0:10], s.y[0:10]) map (from: err) map (s.t.t[16:3])
#pragma omp target map (s.x[0:10], s.y, s.y[0:10]) map (from: err) map (s.t.t[16:3])
{
err = s.x[2] != 34 || s.y[2] != 40 || s.t.t[17] != 82;
s.x[2]++;

4
libgomp/testsuite/libgomp.c++/target-23.C
View File

@ -16,13 +16,13 @@ main (void)
s->data[i] = 0;
#pragma omp target enter data map(to: s)
#pragma omp target enter data map(to: s->data[:SZ])
#pragma omp target enter data map(to: s->data, s->data[:SZ])
#pragma omp target
{
for (int i = 0; i < SZ; i++)
s->data[i] = i;
}
#pragma omp target exit data map(from: s->data[:SZ])
#pragma omp target exit data map(from: s->data, s->data[:SZ])
#pragma omp target exit data map(from: s)
for (int i = 0; i < SZ; i++)

46
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-2.c Normal file
View File

@ -0,0 +1,46 @@
#include <stdlib.h>
#define N 10
struct S
{
int a, b;
int *ptr;
int c, d;
};
int
main (void)
{
struct S a;
a.ptr = (int *) malloc (sizeof (int) * N);
for (int i = 0; i < N; i++)
a.ptr[i] = 0;
#pragma omp target enter data map(to: a.ptr, a.ptr[:N])
#pragma omp target
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target update from(a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 1)
abort ();
#pragma omp target map(a.ptr[:N])
for (int i = 0; i < N; i++)
a.ptr[i] += 1;
#pragma omp target update from(a.ptr[:N])
for (int i = 0; i < N; i++)
if (a.ptr[i] != 2)
abort ();
#pragma omp target exit data map(from:a.ptr, a.ptr[:N])
return 0;
}

2
libgomp/testsuite/libgomp.c/target-23.c
View File

@ -8,7 +8,7 @@ main ()
int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
int *v = u + 4;
#pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
#pragma omp target enter data map (to: s.s, s.u, s.u[0:5]) map (alloc: s.v[1:3])
s.s++;
u[3]++;
s.v[1]++;

20
libgomp/testsuite/libgomp.c/target-29.c
View File

@ -14,7 +14,7 @@ foo (struct S s)
d = id;
int err;
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err)
#pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(to: sep) map(from: err)
{
err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
@ -35,7 +35,7 @@ foo (struct S s)
|| omp_target_is_present (s.d, d)
|| omp_target_is_present (&s.d[-2], d)))
abort ();
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
{
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
@ -43,15 +43,15 @@ foo (struct S s)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d))
abort ();
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
#pragma omp target update to(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
{
err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
s.a = 17; s.b[0] = 18; s.b[1] = 19;
s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
}
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target update from(s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
}
if (sep
&& (omp_target_is_present (&s.a, d)
@ -66,29 +66,29 @@ foo (struct S s)
if (err) abort ();
s.a = 33; s.b[0] = 34; s.b[1] = 35;
s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d))
abort ();
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
#pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
#pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d, s.d[-2:3]) map(from: err)
{
err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
s.a = 49; s.b[0] = 48; s.b[1] = 47;
s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
}
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
if (!omp_target_is_present (&s.a, d)
|| !omp_target_is_present (s.b, d)
|| !omp_target_is_present (&s.c[1], d)
|| !omp_target_is_present (s.d, d)
|| !omp_target_is_present (&s.d[-2], d))
abort ();
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3])
#pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d, s.d[-2:3])
if (sep
&& (omp_target_is_present (&s.a, d)
|| omp_target_is_present (s.b, d)