This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |
| Other format: | [Raw text] | |
Hi Jim!
A few things I noticed when working on merging your trunk r230275 into
gomp-4_0-branch. Please fix these (on trunk).
| --- gcc/c-family/c-pragma.h
| +++ gcc/c-family/c-pragma.h
| @@ -176,7 +178,8 @@ enum pragma_omp_clause {
| PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
| 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_REDUCTION = PRAGMA_OMP_CLAUSE_REDUCTION,
| + PRAGMA_OACC_CLAUSE_LINK = PRAGMA_OMP_CLAUSE_LINK
| };
Maintain alphabetic sorting (as it had been present on gomp-4_0-branch, I
think?).
| --- gcc/c/c-parser.c
| +++ gcc/c/c-parser.c
| @@ -10018,6 +10023,8 @@ c_parser_omp_clause_name (c_parser *parser)
| result = PRAGMA_OMP_CLAUSE_DEVICE;
| else if (!strcmp ("deviceptr", p))
| result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
| + else if (!strcmp ("device_resident", p))
| + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
| else if (!strcmp ("dist_schedule", p))
| result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
Lower-case "device_resident" sorts before "deviceptr".
| @@ -10454,10 +10461,16 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
| case PRAGMA_OACC_CLAUSE_DEVICE:
| kind = GOMP_MAP_FORCE_TO;
| break;
| + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
| + kind = GOMP_MAP_DEVICE_RESIDENT;
| + break;
| case PRAGMA_OACC_CLAUSE_HOST:
| case PRAGMA_OACC_CLAUSE_SELF:
| kind = GOMP_MAP_FORCE_FROM;
| break;
| + case PRAGMA_OACC_CLAUSE_LINK:
| + kind = GOMP_MAP_LINK;
| + break;
| case PRAGMA_OACC_CLAUSE_PRESENT:
| kind = GOMP_MAP_FORCE_PRESENT;
| break;
Update accepted syntax comment for c_parser_oacc_data_clause function (as
present on gomp-4_0-branch).
| --- gcc/cp/parser.c
| +++ gcc/cp/parser.c
| @@ -29128,6 +29128,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
| result = PRAGMA_OMP_CLAUSE_DEVICE;
| else if (!strcmp ("deviceptr", p))
| result = PRAGMA_OACC_CLAUSE_DEVICEPTR;
| + else if (!strcmp ("device_resident", p))
| + result = PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT;
| else if (!strcmp ("dist_schedule", p))
| result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
As in gcc/c/c-parser.c.
| @@ -29541,10 +29543,16 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
| case PRAGMA_OACC_CLAUSE_DEVICE:
| kind = GOMP_MAP_FORCE_TO;
| break;
| + case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT:
| + kind = GOMP_MAP_DEVICE_RESIDENT;
| + break;
| case PRAGMA_OACC_CLAUSE_HOST:
| case PRAGMA_OACC_CLAUSE_SELF:
| kind = GOMP_MAP_FORCE_FROM;
| break;
| + case PRAGMA_OACC_CLAUSE_LINK:
| + kind = GOMP_MAP_LINK;
| + break;
| case PRAGMA_OACC_CLAUSE_PRESENT:
| kind = GOMP_MAP_FORCE_PRESENT;
| break;
Likewise.
| +static tree
| +cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
| +{
| + tree clauses, stmt, t;
| + bool error = false;
| +
| + clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
| + "#pragma acc declare", pragma_tok, true);
| +
| +
| + if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
| + {
| + error_at (pragma_tok->location,
| + "no valid clauses specified in %<#pragma acc declare%>");
The last parameter to cp_parser_oacc_all_clauses, true, is the default
anyway, and is not specified at other call sites, so remove that. Also
remove the following double blank lines.
In the C front end, instead of find_omp_clause OMP_CLAUSE_MAP, you just
check for "!clauses" -- unless there is a reason to be different here,
for uniformity settle on one variant.
That said, even if it doesn't make sense, is it actually a hard error to
not specify any clauses with the declare directive?
| --- gcc/gimplify.c
| +++ gcc/gimplify.c
| +/* Return true if global var DECL is device resident. */
| +
| +static bool
| +device_resident_p (tree decl)
I suggest to improve that function's very generic name, and its
descriptive comment. Without more context, the casual reader will not
understand what "device resident" means, for example. At least note that
this relates to the OpenACC declare directive, or something like that.
| +{
| + tree attr = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (decl));
| +
| + if (!attr)
| + return false;
As discussed in
<http://news.gmane.org/find-root.php?message_id=%3C87611h1zi7.fsf%40kepler.schwinge.homeip.net%3E>
already, for "oacc declare" used in an earlier version of this patch:
there is no "oacc declare target" attribute defined/ever set, so I
suspect device_resident_p will always return false, and thus isn't doing
what it's intended to be doing?
| +
| + for (tree t = TREE_VALUE (attr); t; t = TREE_PURPOSE (t))
| + {
| + tree c = TREE_VALUE (t);
| + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DEVICE_RESIDENT)
| + return true;
| + }
| +
| + return false;
| +}
And thus here:
| @@ -5908,6 +5960,15 @@ static unsigned
| oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
| {
| const char *rkind;
| + bool on_device = false;
| +
| + if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
| + && is_global_var (decl)
| + && device_resident_p (decl))
| + {
| + on_device = true;
| + flags |= GOVD_MAP_TO_ONLY;
| + }
|
| switch (ctx->region_type)
| {
| @@ -5928,7 +5989,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
| || POINTER_TYPE_P (type))
| type = TREE_TYPE (type);
|
| - if (AGGREGATE_TYPE_P (type))
| + if (on_device || AGGREGATE_TYPE_P (type))
| /* Aggregates default to 'present_or_copy'. */
| flags |= GOVD_MAP;
| else
... on_device will always be set to false? Missing testsuite coverage?
| +static void
| +gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
| +{
| + tree expr = *expr_p;
| + gomp_target *stmt;
| + tree clauses, t;
| +
| + clauses = OACC_DECLARE_CLAUSES (expr);
| +
| + gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE);
Add a comment why ORT_TARGET_DATA.
It eventually follows ever other gimplify_scan_omp_clauses call, but no
gimplify_adjust_omp_clauses call is needed here?
| --- gcc/omp-builtins.def
| +++ gcc/omp-builtins.def
| @@ -353,3 +353,5 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA,
| BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST)
| DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams",
| BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST)
| +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DECLARE, "GOACC_declare",
| + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
Sort alphabetically with the other OpenACC builtins.
| --- include/gomp-constants.h
| +++ include/gomp-constants.h
| @@ -72,6 +72,11 @@ enum gomp_map_kind
| POINTER_SIZE_UNITS. */
| GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0),
| /* Do not map, copy bits for firstprivate instead. */
| + /* OpenACC device_resident. */
| + GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
| + /* OpenACC link. */
| + GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
| + /* Allocate. */
| GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
| /* Similarly, but store the value in the pointer rather than
| pointed by the pointer. */
I suspect the "Do not map, copy bits for firstprivate instead" comment
still applies to GOMP_MAP_FIRSTPRIVATE only, which here (unintended?) got
an "Allocate" comment added?
To answer your question from
<http://news.gmane.org/find-root.php?message_id=%3C563A3C70.6060500%40codesourcery.com%3E>:
On Wed, 4 Nov 2015 11:12:16 -0600, James Norris <jnorris@codesourcery.com> wrote:
> On 11/04/2015 10:49 AM, Thomas Schwinge wrote:
> > On Tue, 3 Nov 2015 10:31:32 -0600, James Norris <jnorris@codesourcery.com> wrote:
> >> --- a/include/gomp-constants.h
> >> +++ b/include/gomp-constants.h
> >> @@ -73,6 +73,11 @@ enum gomp_map_kind
> >> POINTER_SIZE_UNITS. */
> >> GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0),
> >> /* Do not map, copy bits for firstprivate instead. */
> >> + /* OpenACC device_resident. */
> >> + GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
> >> + /* OpenACC link. */
> >> + GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
> >> + /* Allocate. */
> >> GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
> >> /* Similarly, but store the value in the pointer rather than
> >> pointed by the pointer. */
> >
> > Confused -- I don't see these two getting handled in libgomp?
>
> These won't be 'seen' by libgomp. So should these
> be defined by some other means?
In this case, I suspect they should be moved later in enum gomp_map_kind,
into the "Internal to GCC, not used in libgomp" section.
GrÃÃe
Thomas
Attachment:
signature.asc
Description: PGP signature
| Index Nav: | [Date Index] [Subject Index] [Author Index] [Thread Index] | |
|---|---|---|
| Message Nav: | [Date Prev] [Date Next] | [Thread Prev] [Thread Next] |