procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

OpenACC 2.6 deep copy: C and C++ front-end parts

Browse source 
gcc/c-family/
	* c-common.h (c_omp_map_clause_name): Add prototype.
	* c-omp.c (c_omp_map_clause_name): New function.
	* c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and
	PRAGMA_OACC_CLAUSE_DETACH.

	gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
	detach clauses.
	(c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter.
	Allow deref (->) in variable lists if true.
	(c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter.
	Pass to c_parser_omp_variable_list.
	(c_parser_oacc_data_clause): Support attach and detach clauses.  Update
	call to c_parser_omp_variable_list.
	(c_parser_oacc_all_clauses): Support attach and detach clauses.
	(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK,
	OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
	OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH.
	(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
	* c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach
	and detach.  Support deref.
	(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
	GOMP_MAP_ALWAYS_POINTER for OpenACC.
	(c_oacc_check_attachments): New function.
	(c_finish_omp_clauses): Check attach/detach arguments for being
	pointers using above.  Support deref.

	gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Support attach and detach
	clauses.
	(cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter.
	Parse deref if true.
	(cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter.  Pass to
	cp_parser_omp_var_list_no_open.
	(cp_parser_oacc_data_clause): Support attach and detach clauses.
	Update call to cp_parser_omp_var_list_no_open.
	(cp_parser_oacc_all_clauses): Support attach and detach.
	(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK,
	OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
	OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH.
	(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
	* semantics.c (handle_omp_array_sections_1): Reject subarrays for
	attach and detach.
	(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
	GOMP_MAP_ALWAYS_POINTER for OpenACC.
	(cp_oacc_check_attachments): New function.
	(finish_omp_clauses): Use above function.  Allow structure fields and
	class members to appear in OpenACC data clauses.  Support
	GOMP_MAP_ATTACH_DETACH.  Support deref.

	gcc/testsuite/
	* c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
	* c-c++-common/goacc/mdc-1.c: New test.
	* c-c++-common/goacc/mdc-2.c: New test.
	* gcc.dg/goacc/mdc.C: New test.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r279627
This commit is contained in:
Julian Brown 2019-12-20 01:20:38 +00:00 • committed by Julian Brown
parent 4fd872bc0c
commit 519d7496be
15 changed files with 618 additions and 34 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

7
gcc/c-family/ChangeLog
View file

@ -1,3 +1,10 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
* c-common.h (c_omp_map_clause_name): Add prototype.
* c-omp.c (c_omp_map_clause_name): New function.
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and
PRAGMA_OACC_CLAUSE_DETACH.
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>

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

@ -1205,6 +1205,7 @@ extern bool c_omp_predefined_variable (tree);
extern enum omp_clause_default_kind c_omp_predetermined_sharing (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);
/* Return next tree in the chain for chain_next walking of tree nodes. */
static inline tree

33
gcc/c-family/c-omp.c
View file

@ -2259,3 +2259,36 @@ c_omp_mark_declare_variant (location_t loc, tree variant, tree construct)
error_at (loc, "%qD used as a variant with incompatible %<construct%> "
"selector sets", variant);
}
/* For OpenACC, the OMP_CLAUSE_MAP_KIND of an OMP_CLAUSE_MAP is used internally
to distinguish clauses as seen by the user. Return the "friendly" clause
name for error messages etc., where possible. See also
c/c-parser.c:c_parser_oacc_data_clause and
cp/parser.c:cp_parser_oacc_data_clause. */
const char *
c_omp_map_clause_name (tree clause, bool oacc)
{
if (oacc && OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (clause))
{
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_ALLOC: return "create";
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_TO: return "copyin";
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FROM: return "copyout";
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_TOFROM: return "copy";
case GOMP_MAP_RELEASE: return "delete";
case GOMP_MAP_FORCE_PRESENT: return "present";
case GOMP_MAP_ATTACH: return "attach";
case GOMP_MAP_FORCE_DETACH:
case GOMP_MAP_DETACH: return "detach";
case GOMP_MAP_DEVICE_RESIDENT: return "device_resident";
case GOMP_MAP_LINK: return "link";
case GOMP_MAP_FORCE_DEVICEPTR: return "deviceptr";
default: break;
}
return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
}

2
gcc/c-family/c-pragma.h
View file

@ -143,11 +143,13 @@ enum pragma_omp_clause {
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
PRAGMA_OACC_CLAUSE_ATTACH,
PRAGMA_OACC_CLAUSE_AUTO,
PRAGMA_OACC_CLAUSE_COPY,
PRAGMA_OACC_CLAUSE_COPYOUT,
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
PRAGMA_OACC_CLAUSE_DETACH,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_FINALIZE,

24
gcc/c/ChangeLog
View file

@ -1,3 +1,27 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
* c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
detach clauses.
(c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter.
Allow deref (->) in variable lists if true.
(c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter.
Pass to c_parser_omp_variable_list.
(c_parser_oacc_data_clause): Support attach and detach clauses. Update
call to c_parser_omp_variable_list.
(c_parser_oacc_all_clauses): Support attach and detach clauses.
(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK,
OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH.
(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
* c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach
and detach. Support deref.
(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
GOMP_MAP_ALWAYS_POINTER for OpenACC.
(c_oacc_check_attachments): New function.
(c_finish_omp_clauses): Check attach/detach arguments for being
pointers using above. Support deref.
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>

53
gcc/c/c-parser.c
View file

@ -12564,6 +12564,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_ALIGNED;
else if (!strcmp ("async", p))
result = PRAGMA_OACC_CLAUSE_ASYNC;
else if (!strcmp ("attach", p))
result = PRAGMA_OACC_CLAUSE_ATTACH;
break;
case 'b':
if (!strcmp ("bind", p))
@ -12590,6 +12592,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OACC_CLAUSE_DELETE;
else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
else if (!strcmp ("detach", p))
result = PRAGMA_OACC_CLAUSE_DETACH;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
@ -12835,12 +12839,16 @@ c_parser_oacc_wait_list (c_parser *parser, location_t clause_loc, tree list)
If KIND is nonzero, CLAUSE_LOC is the location of the clause.
If KIND is zero, create a TREE_LIST with the decl in TREE_PURPOSE;
return the list created. */
return the list created.
The optional ALLOW_DEREF argument is true if list items can use the deref
(->) operator. */
static tree
c_parser_omp_variable_list (c_parser *parser,
location_t clause_loc,
enum omp_clause_code kind, tree list)
enum omp_clause_code kind, tree list,
bool allow_deref = false)
{
auto_vec<c_token> tokens;
unsigned int tokens_avail = 0;
@ -12967,9 +12975,13 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
while (c_parser_next_token_is (parser, CPP_DOT))
while (c_parser_next_token_is (parser, CPP_DOT)
|| (allow_deref
&& c_parser_next_token_is (parser, CPP_DEREF)))
{
location_t op_loc = c_parser_peek_token (parser)->location;
if (c_parser_next_token_is (parser, CPP_DEREF))
t = build_simple_mem_ref (t);
c_parser_consume_token (parser);
if (!c_parser_next_token_is (parser, CPP_NAME))
{
@ -13091,11 +13103,12 @@ c_parser_omp_variable_list (c_parser *parser,
}
/* Similarly, but expect leading and trailing parenthesis. This is a very
common case for OpenACC and OpenMP clauses. */
common case for OpenACC and OpenMP clauses. The optional ALLOW_DEREF
argument is true if list items can use the deref (->) operator. */
static tree
c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
tree list)
tree list, bool allow_deref = false)
{
/* The clauses location. */
location_t loc = c_parser_peek_token (parser)->location;
@ -13103,7 +13116,7 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
matching_parens parens;
if (parens.require_open (parser))
{
list = c_parser_omp_variable_list (parser, loc, kind, list);
list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
parens.skip_until_found_close (parser);
}
return list;
@ -13118,7 +13131,9 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
present ( variable-list )
OpenACC 2.6:
no_create ( variable-list ) */
no_create ( variable-list )
attach ( variable-list )
detach ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
@ -13127,6 +13142,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
enum gomp_map_kind kind;
switch (c_kind)
{
case PRAGMA_OACC_CLAUSE_ATTACH:
kind = GOMP_MAP_ATTACH;
break;
case PRAGMA_OACC_CLAUSE_COPY:
kind = GOMP_MAP_TOFROM;
break;
@ -13142,6 +13160,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_DELETE:
kind = GOMP_MAP_RELEASE;
break;
case PRAGMA_OACC_CLAUSE_DETACH:
kind = GOMP_MAP_DETACH;
break;
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
break;
@ -13164,7 +13185,7 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
gcc_unreachable ();
}
tree nl, c;
nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@ -15879,6 +15900,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses);
c_name = "auto";
break;
case PRAGMA_OACC_CLAUSE_ATTACH:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "attach";
break;
case PRAGMA_OACC_CLAUSE_COLLAPSE:
clauses = c_parser_omp_clause_collapse (parser, clauses);
c_name = "collapse";
@ -15907,6 +15932,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_default (parser, clauses, true);
c_name = "default";
break;
case PRAGMA_OACC_CLAUSE_DETACH:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "detach";
break;
case PRAGMA_OACC_CLAUSE_DEVICE:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device";
@ -16421,7 +16450,8 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
*/
#define OACC_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
@ -16605,6 +16635,7 @@ c_parser_oacc_declare (c_parser *parser)
#define OACC_ENTER_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@ -16614,6 +16645,7 @@ c_parser_oacc_declare (c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@ -16753,6 +16785,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
#define OACC_KERNELS_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@ -16769,6 +16802,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
#define OACC_PARALLEL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@ -16788,6 +16822,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
#define OACC_SERIAL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \

76
gcc/c/c-typeck.c
View file

@ -12897,7 +12897,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
if (TREE_CODE (t) == COMPONENT_REF
&& ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
@ -12918,6 +12917,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
if (ort == C_ORT_ACC && 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);
}
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@ -13003,7 +13011,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
length = fold_convert (sizetype, length);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
{
if (length != integer_one_node)
{
error_at (OMP_CLAUSE_LOCATION (c),
"expected single pointer in %qs clause",
c_omp_map_clause_name (c, ort == C_ORT_ACC));
return error_mark_node;
}
}
if (length != NULL_TREE)
{
if (!integer_nonzerop (length))
@ -13444,7 +13463,11 @@ 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)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
{
gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
@ -13681,6 +13704,35 @@ c_omp_finish_iterators (tree iter)
return ret;
}
/* Ensure that pointers are used in OpenACC attach and detach clauses.
Return true if an error has been detected. */
static bool
c_oacc_check_attachments (tree c)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
/* OpenACC attach / detach clauses must be pointers. */
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
{
tree t = OMP_CLAUSE_DECL (c);
while (TREE_CODE (t) == TREE_LIST)
t = TREE_CHAIN (t);
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause",
c_omp_map_clause_name (c, true));
return true;
}
}
return false;
}
/* For all elements of CLAUSES, validate them against their constraints.
Remove any elements from the list that are invalid. */
@ -14434,6 +14486,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
}
}
if (c_oacc_check_attachments (c))
remove = true;
break;
}
if (t == error_mark_node)
@ -14441,8 +14495,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
/* OpenACC attach / detach clauses must be pointers. */
if (c_oacc_check_attachments (c))
{
remove = true;
break;
}
if (TREE_CODE (t) == COMPONENT_REF
&& (ort & C_ORT_OMP)
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
@ -14477,6 +14536,15 @@ 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 (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);
}
}
if (remove)
break;

25
gcc/cp/ChangeLog
View file

@ -1,3 +1,28 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
* parser.c (cp_parser_omp_clause_name): Support attach and detach
clauses.
(cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter.
Parse deref if true.
(cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter. Pass to
cp_parser_omp_var_list_no_open.
(cp_parser_oacc_data_clause): Support attach and detach clauses.
Update call to cp_parser_omp_var_list_no_open.
(cp_parser_oacc_all_clauses): Support attach and detach.
(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK,
OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH.
(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
* semantics.c (handle_omp_array_sections_1): Reject subarrays for
attach and detach.
(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
GOMP_MAP_ALWAYS_POINTER for OpenACC.
(cp_oacc_check_attachments): New function.
(finish_omp_clauses): Use above function. Allow structure fields and
class members to appear in OpenACC data clauses. Support
GOMP_MAP_ATTACH_DETACH. Support deref.
2019-12-19 Jason Merrill <jason@redhat.com>
PR c++/52320 - EH cleanups for partially constructed arrays.

56
gcc/cp/parser.c
View file

@ -33538,6 +33538,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_ALIGNED;
else if (!strcmp ("async", p))
result = PRAGMA_OACC_CLAUSE_ASYNC;
else if (!strcmp ("attach", p))
result = PRAGMA_OACC_CLAUSE_ATTACH;
break;
case 'b':
if (!strcmp ("bind", p))
@ -33562,6 +33564,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_DEFAULTMAP;
else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
else if (!strcmp ("detach", p))
result = PRAGMA_OACC_CLAUSE_DETACH;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
@ -33766,11 +33770,15 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
COLON can be NULL if only closing parenthesis should end the list,
or pointer to bool which will receive false if the list is terminated
by closing parenthesis or true if the list is terminated by colon. */
by closing parenthesis or true if the list is terminated by colon.
The optional ALLOW_DEREF argument is true if list items can use the deref
(->) operator. */
static tree
cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
tree list, bool *colon)
tree list, bool *colon,
bool allow_deref = false)
{
cp_token *token;
bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
@ -33851,15 +33859,20 @@ 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:
while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
|| (allow_deref
&& cp_lexer_next_token_is (parser->lexer, CPP_DEREF)))
{
cpp_ttype ttype
= cp_lexer_next_token_is (parser->lexer, CPP_DOT)
? CPP_DOT : CPP_DEREF;
location_t loc
= cp_lexer_peek_token (parser->lexer)->location;
cp_id_kind idk = CP_ID_KIND_NONE;
cp_lexer_consume_token (parser->lexer);
decl = convert_from_reference (decl);
decl
= cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
= cp_parser_postfix_dot_deref_expression (parser, ttype,
decl, false,
&idk, loc);
}
@ -33977,10 +33990,12 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
common case for omp clauses. */
static tree
cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
bool allow_deref = false)
{
if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return cp_parser_omp_var_list_no_open (parser, kind, list, NULL);
return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
allow_deref);
return list;
}
@ -33993,7 +34008,9 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
present ( variable-list )
OpenACC 2.6:
no_create ( variable-list ) */
no_create ( variable-list )
attach ( variable-list )
detach ( variable-list ) */
static tree
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@ -34002,6 +34019,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
enum gomp_map_kind kind;
switch (c_kind)
{
case PRAGMA_OACC_CLAUSE_ATTACH:
kind = GOMP_MAP_ATTACH;
break;
case PRAGMA_OACC_CLAUSE_COPY:
kind = GOMP_MAP_TOFROM;
break;
@ -34017,6 +34037,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_DELETE:
kind = GOMP_MAP_RELEASE;
break;
case PRAGMA_OACC_CLAUSE_DETACH:
kind = GOMP_MAP_DETACH;
break;
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
break;
@ -34039,7 +34062,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
gcc_unreachable ();
}
tree nl, c;
nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list);
nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@ -36517,6 +36540,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses);
c_name = "auto";
break;
case PRAGMA_OACC_CLAUSE_ATTACH:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "attach";
break;
case PRAGMA_OACC_CLAUSE_COLLAPSE:
clauses = cp_parser_omp_clause_collapse (parser, clauses, here);
c_name = "collapse";
@ -36545,6 +36572,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_clause_default (parser, clauses, here, true);
c_name = "default";
break;
case PRAGMA_OACC_CLAUSE_DETACH:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "detach";
break;
case PRAGMA_OACC_CLAUSE_DEVICE:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device";
@ -40397,10 +40428,12 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
structured-block */
#define OACC_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
@ -40601,6 +40634,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
#define OACC_ENTER_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
@ -40611,6 +40645,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@ -40718,6 +40753,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
#define OACC_KERNELS_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@ -40734,6 +40770,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
#define OACC_PARALLEL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@ -40753,6 +40790,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
#define OACC_SERIAL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \

98
gcc/cp/semantics.c
View file

@ -4744,7 +4744,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t = TREE_OPERAND (t, 0);
ret = t;
if (TREE_CODE (t) == COMPONENT_REF
&& ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
@ -4768,6 +4767,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
t = TREE_OPERAND (t, 0);
}
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
@ -4867,6 +4868,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
{
if (length != integer_one_node)
{
error_at (OMP_CLAUSE_LOCATION (c),
"expected single pointer in %qs clause",
c_omp_map_clause_name (c, ort == C_ORT_ACC));
return error_mark_node;
}
}
if (length != NULL_TREE)
{
if (!integer_nonzerop (length))
@ -5315,12 +5328,18 @@ 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)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
{
gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
@ -6243,6 +6262,41 @@ cp_omp_finish_iterators (tree iter)
return ret;
}
/* Ensure that pointers are used in OpenACC attach and detach clauses.
Return true if an error has been detected. */
static bool
cp_oacc_check_attachments (tree c)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
/* OpenACC attach / detach clauses must be pointers. */
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
{
tree t = OMP_CLAUSE_DECL (c);
tree type;
while (TREE_CODE (t) == TREE_LIST)
t = TREE_CHAIN (t);
type = TREE_TYPE (t);
if (TREE_CODE (type) == REFERENCE_TYPE)
type = TREE_TYPE (type);
if (TREE_CODE (type) != POINTER_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause",
c_omp_map_clause_name (c, true));
return true;
}
}
return false;
}
/* For all elements of CLAUSES, validate them vs OpenMP constraints.
Remove any elements from the list that are invalid. */
@ -6507,7 +6561,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
check_dup_generic_t:
if (t == current_class_ptr
&& (ort != C_ORT_OMP_DECLARE_SIMD
&& ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
|| (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
{
@ -6577,8 +6631,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
handle_field_decl:
if (!remove
&& TREE_CODE (t) == FIELD_DECL
&& t == OMP_CLAUSE_DECL (c)
&& ort != C_ORT_ACC)
&& t == OMP_CLAUSE_DECL (c))
{
OMP_CLAUSE_DECL (c)
= omp_privatize_field (t, (OMP_CLAUSE_CODE (c)
@ -6645,7 +6698,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
if (t == current_class_ptr)
if (ort != C_ORT_ACC && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@ -7134,7 +7187,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (t == error_mark_node)
remove = true;
else if (t == current_class_ptr)
else if (ort != C_ORT_ACC && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@ -7266,6 +7319,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
}
}
if (cp_oacc_check_attachments (c))
remove = true;
break;
}
if (t == error_mark_node)
@ -7273,14 +7328,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
/* OpenACC attach / detach clauses must be pointers. */
if (cp_oacc_check_attachments (c))
{
remove = true;
break;
}
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
if (ort == C_ORT_ACC
&& TREE_CODE (t) == COMPONENT_REF
&& TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
if (TREE_CODE (t) == COMPONENT_REF
&& (ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
|| ort == C_ORT_ACC)
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
if (type_dependent_expression_p (t))
@ -7330,7 +7396,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER))
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH))
break;
if (DECL_P (t))
error_at (OMP_CLAUSE_LOCATION (c),
@ -7412,7 +7479,9 @@ 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_ACC
|| !bitmap_bit_p (&map_field_head, DECL_UID (t))))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
@ -7467,7 +7536,12 @@ 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_ALWAYS_POINTER);
{
gomp_map_kind k
= (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER;
OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);

8
gcc/testsuite/ChangeLog
View file

@ -1,3 +1,11 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
* c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
* c-c++-common/goacc/mdc-1.c: New test.
* c-c++-common/goacc/mdc-2.c: New test.
* gcc.dg/goacc/mdc.C: New test.
2019-12-19 Vladimir Makarov <vmakarov@redhat.com>
PR target/92905

84
gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c Normal file
View file

@ -0,0 +1,84 @@
/* { dg-do compile } */
#include <stdlib.h>
#include <stdio.h>
typedef struct {
int *a;
int *b;
int *c;
} mystruct;
int main(int argc, char* argv[])
{
const int N = 1024;
const int S = 32;
mystruct *m = (mystruct *) calloc (S, sizeof (*m));
int i, j;
for (i = 0; i < S; i++)
{
m[i].a = (int *) malloc (N * sizeof (int));
m[i].b = (int *) malloc (N * sizeof (int));
m[i].c = (int *) malloc (N * sizeof (int));
}
for (j = 0; j < S; j++)
for (i = 0; i < N; i++)
{
m[j].a[i] = 0;
m[j].b[i] = 0;
m[j].c[i] = 0;
}
#pragma acc enter data copyin(m[0:1])
for (int i = 0; i < 99; i++)
{
int j, k;
for (k = 0; k < S; k++)
#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before .\\\.. token" } */
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 } */
for (j = 0; j < N; j++)
{
m[k].b[j]++;
if (j > 5 && j < N - 5)
m[k].c[j]++;
}
}
#pragma acc exit data copyout(m[0:1])
for (j = 0; j < S; j++)
{
for (i = 0; i < N; i++)
{
if (m[j].a[i] != 99)
abort ();
if (m[j].b[i] != 99)
abort ();
if (i > 5 && i < N-5)
{
if (m[j].c[i] != 99)
abort ();
}
else
{
if (m[j].c[i] != 0)
abort ();
}
}
free (m[j].a);
free (m[j].b);
free (m[j].c);
}
free (m);
return 0;
}

55
gcc/testsuite/c-c++-common/goacc/mdc-1.c Normal file
View file

@ -0,0 +1,55 @@
/* Test OpenACC's support for manual deep copy, including the attach
and detach clauses. */
/* { dg-do compile { target int32 } } */
/* { dg-additional-options "-fdump-tree-omplower" } */
void
t1 ()
{
struct foo {
int *a, *b, c, d, *e;
} s;
int *a, *z;
#pragma acc enter data copyin(s)
{
#pragma acc data copy(s.a[0:10]) copy(z[0:10])
{
s.e = z;
#pragma acc parallel loop attach(s.e)
for (int i = 0; i < 10; i++)
s.a[i] = s.e[i];
a = s.e;
#pragma acc enter data attach(a)
#pragma acc exit data detach(a)
}
#pragma acc enter data copyin(a)
#pragma acc acc enter data attach(s.e)
#pragma acc exit data detach(s.e)
#pragma acc data attach(s.e)
{
}
#pragma acc exit data delete(a)
#pragma acc exit data detach(a) finalize
#pragma acc exit data detach(s.a) finalize
}
}
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 8.." 1 "omplower" } } */

62
gcc/testsuite/c-c++-common/goacc/mdc-2.c Normal file
View file

@ -0,0 +1,62 @@
/* Test OpenACC's support for manual deep copy, including the attach
and detach clauses. */
void
t1 ()
{
struct foo {
int *a, *b, c, d, *e;
} s;
int *a, *z, scalar, **y;
#pragma acc enter data copyin(s) detach(z) /* { dg-error ".detach. is not valid for" } */
{
#pragma acc data copy(s.a[0:10]) copy(z[0:10])
{
s.e = z;
#pragma acc parallel loop attach(s.e) detach(s.b) /* { dg-error ".detach. is not valid for" } */
for (int i = 0; i < 10; i++)
s.a[i] = s.e[i];
a = s.e;
#pragma acc enter data attach(a) detach(s.c) /* { dg-error ".detach. is not valid for" } */
#pragma acc exit data detach(a)
}
#pragma acc enter data attach(z[:5]) /* { dg-error "expected single pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(z[:5]) /* { dg-error "expected single pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(z[1:]) /* { dg-error "expected single pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(z[1:]) /* { dg-error "expected single pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(z[:]) /* { dg-error "expected single pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(z[:]) /* { dg-error "expected single pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(z[3]) /* { dg-error "expected pointer in .attach. clause" } */
#pragma acc exit data detach(z[3]) /* { dg-error "expected pointer in .detach. clause" } */
#pragma acc acc enter data attach(s.e)
#pragma acc exit data detach(s.e) attach(z) /* { dg-error ".attach. is not valid for" } */
#pragma acc data attach(s.e)
{
}
#pragma acc exit data delete(a) attach(s.a) /* { dg-error ".attach. is not valid for" } */
#pragma acc enter data attach(scalar) /* { dg-error "expected pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(scalar) /* { dg-error "expected pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(s) /* { dg-error "expected pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(s) /* { dg-error "expected pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
}
#pragma acc enter data attach(y[10])
#pragma acc exit data detach(y[10])
}

68
gcc/testsuite/g++.dg/goacc/mdc.C Normal file
View file

@ -0,0 +1,68 @@
/* Test OpenACC's support for manual deep copy, including the attach
and detach clauses. */
void
t1 ()
{
struct foo {
int *a, *b, c, d, *e;
} s;
struct foo& rs = s;
int *a, *z, scalar, **y;
int* const &ra = a;
int* const &rz = z;
int& rscalar = scalar;
int** const &ry = y;
#pragma acc enter data copyin(rs) detach(rz) /* { dg-error ".detach. is not valid for" } */
{
#pragma acc data copy(rs.a[0:10]) copy(rz[0:10])
{
s.e = z;
#pragma acc parallel loop attach(rs.e) detach(rs.b) /* { dg-error ".detach. is not valid for" } */
for (int i = 0; i < 10; i++)
s.a[i] = s.e[i];
a = s.e;
#pragma acc enter data attach(ra) detach(rs.c) /* { dg-error ".detach. is not valid for" } */
#pragma acc exit data detach(ra)
}
#pragma acc enter data attach(rz[:5]) /* { dg-error "expected single pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(rz[:5]) /* { dg-error "expected single pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(rz[1:]) /* { dg-error "expected single pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(rz[1:]) /* { dg-error "expected single pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(rz[:]) /* { dg-error "expected single pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(rz[:]) /* { dg-error "expected single pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(rz[3]) /* { dg-error "expected pointer in .attach. clause" } */
#pragma acc exit data detach(rz[3]) /* { dg-error "expected pointer in .detach. clause" } */
#pragma acc acc enter data attach(rs.e)
#pragma acc exit data detach(rs.e) attach(rz) /* { dg-error ".attach. is not valid for" } */
#pragma acc data attach(rs.e)
{
}
#pragma acc exit data delete(ra) attach(rs.a) /* { dg-error ".attach. is not valid for" } */
#pragma acc enter data attach(rscalar) /* { dg-error "expected pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(rscalar) /* { dg-error "expected pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc enter data attach(rs) /* { dg-error "expected pointer in .attach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
#pragma acc exit data detach(rs) /* { dg-error "expected pointer in .detach. clause" } */
/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */
}
#pragma acc enter data attach(ry[10])
#pragma acc exit data detach(ry[10])
}