mirrors/gcc
0
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2025-01-25 13:53:56 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR OpenMP description.

Browse Source 
* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
	OpenMP description.  Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
	* tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
	for OMP_CLAUSE_USE_DEVICE_ADDR clause.
	(walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree-nested.c (convert_nonlocal_omp_clauses,
	convert_local_omp_clauses): Likewise.
	* gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
	Likewise.
	* omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
	Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
	clause with array or reference to array types, no matter what type
	except for reference it has.
gcc/c-family/
	* c-pragma.h (enum pragma_omp_clause): Add
	PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.  Set PRAGMA_OACC_CLAUSE_USE_DEVICE
	equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
	enumeration value.
gcc/c/
	* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
	(c_parser_omp_clause_use_device_addr): New function.
	(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer type rather than pointer or array type.
	Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
	(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
	(cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
	like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
	map or use_device_* clauses.
	* semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
	in OpenMP, require pointer or reference to pointer type rather than
	pointer or array or reference to pointer or array type. Handle
	OMP_CLAUSE_USE_DEVICE_ADDR.
	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
gcc/testsuite/
	* c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
	instead of use_device_ptr clause where required by OpenMP 5.0, add
	further tests for both use_device_ptr and use_device_addr clauses.
libgomp/
	* testsuite/libgomp.c/target-18.c (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.
	* testsuite/libgomp.c++/target-9.C (struct S): New type.
	(foo): Use use_device_addr clause instead of use_device_ptr clause
	where required by OpenMP 5.0, add further tests for both use_device_ptr
	and use_device_addr clauses.  Add t and u arguments.
	(main): Adjust caller.

From-SVN: r274159
This commit is contained in:
Jakub Jelinek 2019-08-07 09:27:10 +02:00 committed by Jakub Jelinek
parent 4e708f5ebd
commit 398e3feb8a
21 changed files with 291 additions and 46 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

17
gcc/ChangeLog
View File

@ -1,3 +1,20 @@
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* tree-core.h (enum omp_clause_code): Adjust OMP_CLAUSE_USE_DEVICE_PTR
OpenMP description. Add OMP_CLAUSE_USE_DEVICE_ADDR clause.
* tree.c (omp_clause_num_ops, omp_clause_code_name): Add entries
for OMP_CLAUSE_USE_DEVICE_ADDR clause.
(walk_tree_1): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
* tree-pretty-print.c (dump_omp_clause): Likewise.
* tree-nested.c (convert_nonlocal_omp_clauses,
convert_local_omp_clauses): Likewise.
* gimplify.c (gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses):
Likewise.
* omp-low.c (scan_sharing_clauses, lower_omp_target): Likewise.
Treat OMP_CLAUSE_USE_DEVICE_ADDR like OMP_CLAUSE_USE_DEVICE_PTR
clause with array or reference to array types, no matter what type
except for reference it has.
2019-08-07 Kewen Lin <linkw@gcc.gnu.org>
* config/rs6000/vector.md (vrotr<mode>3): New define_expand.

7
gcc/c-family/ChangeLog
View File

@ -1,3 +1,10 @@
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* c-pragma.h (enum pragma_omp_clause): Add
PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR. Set PRAGMA_OACC_CLAUSE_USE_DEVICE
equal to PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR instead of being a separate
enumeration value.
2019-08-05 Marek Polacek <polacek@redhat.com>
PR c++/91338 - Implement P1161R3: Deprecate a[b,c].

5
gcc/c-family/c-pragma.h
View File

@ -137,6 +137,7 @@ enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_UNIFORM,
PRAGMA_OMP_CLAUSE_UNTIED,
PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR,
PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR,
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
@ -157,7 +158,6 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_SELF,
PRAGMA_OACC_CLAUSE_SEQ,
PRAGMA_OACC_CLAUSE_TILE,
PRAGMA_OACC_CLAUSE_USE_DEVICE,
PRAGMA_OACC_CLAUSE_VECTOR,
PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
PRAGMA_OACC_CLAUSE_WAIT,
@ -171,7 +171,8 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE,
PRAGMA_OACC_CLAUSE_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION,
PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK
PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK,
PRAGMA_OACC_CLAUSE_USE_DEVICE = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR
};
extern struct cpp_reader* parse_in;

13
gcc/c/ChangeLog
View File

@ -1,3 +1,16 @@
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
(c_parser_omp_clause_use_device_addr): New function.
(c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
(c_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
map or use_device_* clauses.
* c-typeck.c (c_finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
in OpenMP, require pointer type rather than pointer or array type.
Handle OMP_CLAUSE_USE_DEVICE_ADDR.
2019-07-31 Jakub Jelinek <jakub@redhat.com>
PR c/91192

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

@ -11866,6 +11866,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
else if (!strcmp ("use_device", p))
result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
else if (!strcmp ("use_device_addr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
break;
@ -13121,6 +13123,16 @@ c_parser_omp_clause_use_device_ptr (c_parser *parser, tree list)
list);
}
/* OpenMP 5.0:
use_device_addr ( variable-list ) */
static tree
c_parser_omp_clause_use_device_addr (c_parser *parser, tree list)
{
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE_ADDR,
list);
}
/* OpenMP 4.5:
is_device_ptr ( variable-list ) */
@ -15321,6 +15333,10 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_use_device_ptr (parser, clauses);
c_name = "use_device_ptr";
break;
case PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR:
clauses = c_parser_omp_clause_use_device_addr (parser, clauses);
c_name = "use_device_addr";
break;
case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
clauses = c_parser_omp_clause_is_device_ptr (parser, clauses);
c_name = "is_device_ptr";
@ -18288,7 +18304,8 @@ c_parser_omp_teams (location_t loc, c_parser *parser,
( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR))
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR))
static tree
c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
@ -18323,7 +18340,8 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
*pc = OMP_CLAUSE_CHAIN (*pc);
continue;
}
else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR)
else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
map_seen = 3;
pc = &OMP_CLAUSE_CHAIN (*pc);
}
@ -18333,7 +18351,8 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
if (map_seen == 0)
error_at (loc,
"%<#pragma omp target data%> must contain at least "
"one %<map%> or %<use_device_ptr%> clause");
"one %<map%>, %<use_device_ptr%> or %<use_device_addr%> "
"clause");
return NULL_TREE;
}

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

@ -14609,16 +14609,32 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
t = OMP_CLAUSE_DECL (c);
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE
&& TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is neither a pointer nor an array",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
&& ort == C_ORT_OMP)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is not a pointer",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is neither a pointer nor an array",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
}
goto check_dup_generic;
case OMP_CLAUSE_USE_DEVICE_ADDR:
t = OMP_CLAUSE_DECL (c);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
c_mark_addressable (t);
goto check_dup_generic;
case OMP_CLAUSE_NOWAIT:
if (copyprivate_seen)
{

14
gcc/cp/ChangeLog
View File

@ -1,3 +1,17 @@
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* parser.c (cp_parser_omp_clause_name): Parse use_device_addr clause.
(cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
(OMP_TARGET_DATA_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR.
(cp_parser_omp_target_data): Handle PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR
like PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR, adjust diagnostics about no
map or use_device_* clauses.
* semantics.c (finish_omp_clauses): For OMP_CLAUSE_USE_DEVICE_PTR
in OpenMP, require pointer or reference to pointer type rather than
pointer or array or reference to pointer or array type. Handle
OMP_CLAUSE_USE_DEVICE_ADDR.
* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_USE_DEVICE_ADDR.
2019-08-06 Jason Merrill <jason@redhat.com>
PR c++/91378 - ICE with noexcept and auto return type.

16
gcc/cp/parser.c
View File

@ -32648,6 +32648,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_UNTIED;
else if (!strcmp ("use_device", p))
result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
else if (!strcmp ("use_device_addr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
break;
@ -35637,6 +35639,11 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses);
c_name = "use_device_ptr";
break;
case PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_ADDR,
clauses);
c_name = "use_device_addr";
break;
case PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_IS_DEVICE_PTR,
clauses);
@ -38715,7 +38722,8 @@ cp_parser_omp_teams (cp_parser *parser, cp_token *pragma_tok,
( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR))
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_USE_DEVICE_ADDR))
static tree
cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
@ -38751,7 +38759,8 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
*pc = OMP_CLAUSE_CHAIN (*pc);
continue;
}
else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR)
else if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
|| OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
map_seen = 3;
pc = &OMP_CLAUSE_CHAIN (*pc);
}
@ -38761,7 +38770,8 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
if (map_seen == 0)
error_at (pragma_tok->location,
"%<#pragma omp target data%> must contain at least "
"one %<map%> or %<use_device_ptr%> clause");
"one %<map%>, %<use_device_ptr%> or %<use_device_addr%> "
"clause");
return NULL_TREE;
}

2
gcc/cp/pt.c
View File

@ -16303,6 +16303,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_NONTEMPORAL:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:
@ -16427,6 +16428,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
case OMP_CLAUSE_IN_REDUCTION:
case OMP_CLAUSE_TASK_REDUCTION:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:

39
gcc/cp/semantics.c
View File

@ -7524,20 +7524,41 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
{
tree type = TREE_TYPE (t);
if (!TYPE_PTR_P (type)
&& TREE_CODE (type) != ARRAY_TYPE
&& (!TYPE_REF_P (type)
|| (!TYPE_PTR_P (TREE_TYPE (type))
&& TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE)))
&& (!TYPE_REF_P (type) || !TYPE_PTR_P (TREE_TYPE (type))))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is neither a pointer, nor an array "
"nor reference to pointer or array",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
&& ort == C_ORT_OMP)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is neither a pointer "
"nor reference to pointer",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
else if (TREE_CODE (type) != ARRAY_TYPE
&& (!TYPE_REF_P (type)
|| TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qs variable is neither a pointer, nor an "
"array nor reference to pointer or array",
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
}
}
goto check_dup_generic;
case OMP_CLAUSE_USE_DEVICE_ADDR:
field_ok = true;
t = OMP_CLAUSE_DECL (c);
if (!processing_template_decl
&& (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
&& !TYPE_REF_P (TREE_TYPE (t))
&& !cxx_mark_addressable (t))
remove = true;
goto check_dup_generic;
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_DEFAULT:
case OMP_CLAUSE_UNTIED:

4
gcc/gimplify.c
View File

@ -9015,8 +9015,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_notice;
case OMP_CLAUSE_USE_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
@ -10264,6 +10263,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_ORDER:
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:

31
gcc/omp-low.c
View File

@ -1238,8 +1238,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (decl))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 3, ctx);
else
install_var_field (decl, false, 3, ctx);
@ -1635,6 +1638,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_ORDER:
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_NONTEMPORAL:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
@ -11465,6 +11469,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
map_cnt++;
@ -11481,7 +11486,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (var))
|| TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
{
tree new_var = lookup_decl (var, ctx);
tree type = build_pointer_type (TREE_TYPE (var));
@ -11846,23 +11853,27 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
tkind = GOMP_MAP_USE_DEVICE_PTR;
else
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
type = TREE_TYPE (ovar);
if (TREE_CODE (type) == ARRAY_TYPE)
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (ovar))
|| TREE_CODE (type) == ARRAY_TYPE)
var = build_fold_addr_expr (var);
else
{
if (omp_is_reference (ovar))
{
type = TREE_TYPE (type);
if (TREE_CODE (type) != ARRAY_TYPE)
if (TREE_CODE (type) != ARRAY_TYPE
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR)
var = build_simple_mem_ref (var);
var = fold_convert (TREE_TYPE (x), var);
}
@ -12017,9 +12028,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
break;
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
x = build_sender_ref (var, ctx);
else
x = build_receiver_ref (var, false, ctx);
@ -12034,7 +12046,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq_add_stmt (&new_body,
gimple_build_assign (new_var, x));
}
else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_reference (var))
|| TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
{
tree new_var = lookup_decl (var, ctx);
new_var = DECL_VALUE_EXPR (new_var);
@ -12052,7 +12066,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (omp_is_reference (var))
{
type = TREE_TYPE (type);
if (TREE_CODE (type) != ARRAY_TYPE)
if (TREE_CODE (type) != ARRAY_TYPE
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR)
{
tree v = create_tmp_var_raw (type, get_name (var));
gimple_add_tmp_var (v);

6
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,9 @@
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* c-c++-common/gomp/target-data-1.c (foo): Use use_device_addr clause
instead of use_device_ptr clause where required by OpenMP 5.0, add
further tests for both use_device_ptr and use_device_addr clauses.
2019-08-07 Kewen Lin <linkw@gcc.gnu.org>
* gcc.target/powerpc/vec_rotate-1.c: New test.

32
gcc/testsuite/c-c++-common/gomp/target-data-1.c
View File

@ -4,15 +4,39 @@ void
foo (void)
{
int a[4] = { 1, 2, 3, 4 };
int *p = &a[0];
int x = 5;
#pragma omp target data map(to:p[:4])
#pragma omp target data use_device_ptr(p)
#pragma omp target is_device_ptr(p)
{
p[0]++;
}
#pragma omp target data map(to:a)
#pragma omp target data use_device_ptr(a)
#pragma omp target data use_device_addr(a)
#pragma omp target is_device_ptr(a)
{
a[0]++;
p[0]++;
}
#pragma omp target data map(to:x)
#pragma omp target data use_device_addr(x)
{
int *q = &x;
#pragma omp target is_device_ptr(q)
{
q[0]++;
}
}
#pragma omp target data /* { dg-error "must contain at least one" } */
a[0]++;
#pragma omp target data map(to:a)
#pragma omp target data use_device_ptr(a) use_device_ptr(a) /* { dg-error "appears more than once in data clauses" } */
#pragma omp target data map(to:p)
#pragma omp target data use_device_ptr(p) use_device_ptr(p) /* { dg-error "appears more than once in data clauses" } */
a[0]++;
#pragma omp target data map(to:a)
#pragma omp target data use_device_addr(a) use_device_addr(a) /* { dg-error "appears more than once in data clauses" } */
a[0]++;
#pragma omp target data map(to:a)
#pragma omp target data use_device_ptr(a) /* { dg-error "'use_device_ptr' variable is not a pointer" "" { target c } } */
/* { dg-error "'use_device_ptr' variable is neither a pointer nor reference to pointer" "" { target c++ } .-1 } */
a[0]++; /* { dg-error "must contain at least one" "" { target *-*-* } .-2 } */
}

5
gcc/tree-core.h
View File

@ -307,9 +307,12 @@ enum omp_clause_code {
OMP_CLAUSE_MAP,
/* OpenACC clause: use_device (variable-list).
OpenMP clause: use_device_ptr (variable-list). */
OpenMP clause: use_device_ptr (ptr-list). */
OMP_CLAUSE_USE_DEVICE_PTR,
/* OpenMP clause: use_device_addr (variable-list). */
OMP_CLAUSE_USE_DEVICE_ADDR,
/* OpenMP clause: is_device_ptr (variable-list). */
OMP_CLAUSE_IS_DEVICE_PTR,

2
gcc/tree-nested.c
View File

@ -1227,6 +1227,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
do_decl_clause:
decl = OMP_CLAUSE_DECL (clause);
@ -1947,6 +1948,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
do_decl_clause:
decl = OMP_CLAUSE_DECL (clause);

3
gcc/tree-pretty-print.c
View File

@ -465,6 +465,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_USE_DEVICE_PTR:
name = "use_device_ptr";
goto print_remap;
case OMP_CLAUSE_USE_DEVICE_ADDR:
name = "use_device_addr";
goto print_remap;
case OMP_CLAUSE_IS_DEVICE_PTR:
name = "is_device_ptr";
goto print_remap;

3
gcc/tree.c
View File

@ -299,6 +299,7 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_USE_DEVICE_PTR */
1, /* OMP_CLAUSE_USE_DEVICE_ADDR */
1, /* OMP_CLAUSE_IS_DEVICE_PTR */
1, /* OMP_CLAUSE_INCLUSIVE */
1, /* OMP_CLAUSE_EXCLUSIVE */
@ -382,6 +383,7 @@ const char * const omp_clause_code_name[] =
"to",
"map",
"use_device_ptr",
"use_device_addr",
"is_device_ptr",
"inclusive",
"exclusive",
@ -12384,6 +12386,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_ADDR:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_INCLUSIVE:
case OMP_CLAUSE_EXCLUSIVE:

12
libgomp/ChangeLog
View File

@ -1,3 +1,15 @@
2019-08-07 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c/target-18.c (struct S): New type.
(foo): Use use_device_addr clause instead of use_device_ptr clause
where required by OpenMP 5.0, add further tests for both use_device_ptr
and use_device_addr clauses.
* testsuite/libgomp.c++/target-9.C (struct S): New type.
(foo): Use use_device_addr clause instead of use_device_ptr clause
where required by OpenMP 5.0, add further tests for both use_device_ptr
and use_device_addr clauses. Add t and u arguments.
(main): Adjust caller.
2019-08-06 Jakub Jelinek <jakub@redhat.com>
* testsuite/libgomp.c++/loop-13.C: New test.

43
libgomp/testsuite/libgomp.c++/target-9.C
View File

@ -1,10 +1,13 @@
extern "C" void abort (void);
struct S { int e, f; };
void
foo (int *&p, int (&s)[5], int n)
foo (int *&p, int (&s)[5], int &t, S &u, int n)
{
int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
int *r = a + 1, *q = p - 1, i, err;
int v = 27;
S w = { 28, 29 };
for (i = 0; i < n; i++)
b[i] = 9 + i;
#pragma omp target data map(to:a)
@ -30,7 +33,7 @@ foo (int *&p, int (&s)[5], int n)
if (err)
abort ();
#pragma omp target data map(to:b)
#pragma omp target data use_device_ptr(b) map(from:err)
#pragma omp target data use_device_addr(b) map(from:err)
#pragma omp target is_device_ptr(b) private(i) map(from:err)
{
err = 0;
@ -41,7 +44,7 @@ foo (int *&p, int (&s)[5], int n)
if (err)
abort ();
#pragma omp target data map(to:c)
#pragma omp target data use_device_ptr(c) map(from:err)
#pragma omp target data use_device_addr(c) map(from:err)
#pragma omp target is_device_ptr(c) private(i) map(from:err)
{
err = 0;
@ -52,7 +55,7 @@ foo (int *&p, int (&s)[5], int n)
if (err)
abort ();
#pragma omp target data map(to:s[:5])
#pragma omp target data use_device_ptr(s) map(from:err)
#pragma omp target data use_device_addr(s) map(from:err)
#pragma omp target is_device_ptr(s) private(i) map(from:err)
{
err = 0;
@ -62,6 +65,34 @@ foo (int *&p, int (&s)[5], int n)
}
if (err)
abort ();
#pragma omp target data map(to: v) map(to:u)
#pragma omp target data use_device_addr (v) use_device_addr (u) map(from:err)
{
int *z = &v;
S *x = &u;
#pragma omp target is_device_ptr (z, x) map(from:err)
{
err = 0;
if (*z != 27 || x->e != 25 || x->f != 26)
err = 1;
}
}
if (err)
abort ();
#pragma omp target data map(to: t, w)
#pragma omp target data use_device_addr (t, w) map(from:err)
{
int *z = &t;
S *x = &w;
#pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err)
{
err = 0;
if (*z != 24 || x->e != 28 || x->f != 29)
err = 1;
}
}
if (err)
abort ();
}
int
@ -69,5 +100,7 @@ main ()
{
int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
int *p = a + 1;
foo (p, b, 9);
int t = 24;
S u = { 25, 26 };
foo (p, b, t, u, 9);
}

30
libgomp/testsuite/libgomp.c/target-18.c
View File

@ -1,9 +1,11 @@
extern void abort (void);
struct S { int e, f; };
void
foo (int n)
{
int a[4] = { 0, 1, 2, 3 }, b[n];
int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
struct S d = { 5, 6 };
int *p = a + 1, i, err;
for (i = 0; i < n; i++)
b[i] = 9 + i;
@ -21,7 +23,7 @@ foo (int n)
for (i = 0; i < 4; i++)
a[i] = 23 + i;
#pragma omp target data map(to:a)
#pragma omp target data use_device_ptr(a) map(from:err)
#pragma omp target data use_device_addr(a) map(from:err)
#pragma omp target is_device_ptr(a) private(i) map(from:err)
{
err = 0;
@ -32,7 +34,7 @@ foo (int n)
if (err)
abort ();
#pragma omp target data map(to:b)
#pragma omp target data use_device_ptr(b) map(from:err)
#pragma omp target data use_device_addr(b) map(from:err)
#pragma omp target is_device_ptr(b) private(i) map(from:err)
{
err = 0;
@ -42,6 +44,28 @@ foo (int n)
}
if (err)
abort ();
#pragma omp target data map(to:c)
#pragma omp target data use_device_addr(c) map(from:err)
{
int *q = &c;
#pragma omp target is_device_ptr(q) map(from:err)
{
err = *q != 4;
}
}
if (err)
abort ();
#pragma omp target data map(to:d)
#pragma omp target data use_device_addr(d) map(from:err)
{
struct S *r = &d;
#pragma omp target is_device_ptr(r) map(from:err)
{
err = r->e != 5 || r->f != 6;
}
}
if (err)
abort ();
}
int