From: Nathan Sidwell <nathan@acm.org>
To: GCC Patches <gcc-patches@gcc.gnu.org>
Cc: james norris <James_Norris@mentor.com>
Subject: [gomp4] openacc default handling
Date: Thu, 30 Jul 2015 02:30:00 -0000 [thread overview]
Message-ID: <55B95DFD.6090703@acm.org> (raw)
[-- Attachment #1: Type: text/plain, Size: 444 bytes --]
I've committed this to gomp4. When I broke out the oacc_default_clause
function from omp_notice_variable, I was puzzled by the ordering of the lookups,
but could quite figure out what was wrong. further investigation and standard
reading showed that for openacc, we look in outer lexical scopes to find a data
clause for the object, before applying the default behaviour. This patch
reorganizes things to make that the case.
nathan
[-- Attachment #2: gomp4-dflt-2.patch --]
[-- Type: text/x-patch, Size: 4142 bytes --]
2015-07-29 Nathan Sidwell <nathan@codesourcery.com>
gcc/
* gimplify.c (oacc_default_clause): Outer scope searching moved to
omp_notice_variable.
(omp_notice_variable): For OpenACC search enclosing scopes before
applying default.
gcc/testsuite/
* c-c++-common/goacc/default-2.c: New.
Index: gcc/testsuite/c-c++-common/goacc/default-2.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/default-2.c (revision 0)
+++ gcc/testsuite/c-c++-common/goacc/default-2.c (revision 0)
@@ -0,0 +1,21 @@
+void Foo ()
+{
+ int ary[10];
+
+#pragma acc parallel default(none) /* { dg-error "enclosing" } */
+ {
+ ary[0] = 5; /* { dg-error "not specified" } */
+ }
+}
+
+void Baz ()
+{
+ int ary[10];
+#pragma acc data copy (ary)
+ {
+#pragma acc parallel default(none)
+ {
+ ary[0] = 5;
+ }
+ }
+}
Index: gcc/gimplify.c
===================================================================
--- gcc/gimplify.c (revision 226334)
+++ gcc/gimplify.c (working copy)
@@ -5934,30 +5934,10 @@ oacc_default_clause (struct gimplify_omp
case OMP_CLAUSE_DEFAULT_UNSPECIFIED:
{
- if (struct gimplify_omp_ctx *octx = ctx->outer_context)
- {
- omp_notice_variable (octx, decl, in_code);
-
- for (; octx; octx = octx->outer_context)
- {
- if (octx->region_type & ORT_HOST_DATA)
- continue;
- if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
- break;
- splay_tree_node n2
- = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
- if (n2)
- {
- flags |= GOVD_MAP;
- goto found_outer;
- }
- }
- }
-
if (is_global_var (decl) && device_resident_p (decl))
flags |= GOVD_MAP_TO_ONLY | GOVD_MAP;
- /* Scalars under kernels are default 'copy'. */
else if (ctx->acc_region_kind == ARK_KERNELS)
+ /* Scalars under kernels are default 'copy'. */
flags |= GOVD_FORCE_MAP | GOVD_MAP;
else if (ctx->acc_region_kind == ARK_PARALLEL)
{
@@ -5968,16 +5948,14 @@ oacc_default_clause (struct gimplify_omp
type = TREE_TYPE (type);
if (AGGREGATE_TYPE_P (type))
- /* Aggregates default to 'copy'. This should really
- include GOVD_FORCE_MAP. */
+ /* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
else
- /* Scalars default tp 'firstprivate'. */
+ /* Scalars default to 'firstprivate'. */
flags |= GOVD_GANGLOCAL | GOVD_MAP_TO_ONLY | GOVD_MAP;
}
else
gcc_unreachable ();
- found_outer:;
}
break;
}
@@ -6020,21 +5998,49 @@ omp_notice_variable (struct gimplify_omp
if (ctx->region_type == ORT_TARGET)
{
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
- if (n == NULL)
+ bool is_oacc = ctx->region_kind == ORK_OACC;
+
+ if (!n)
{
- bool is_oacc = ctx->region_kind == ORK_OACC;
+ struct gimplify_omp_ctx *octx = ctx->outer_context;
- if (is_oacc)
- flags = oacc_default_clause (ctx, decl, in_code, flags);
- else
- flags |= GOVD_MAP;
+ /* OpenMP doesn't look in outer contexts to find an
+ enclosing data clause. */
+ if (is_oacc && octx)
+ {
+ omp_notice_variable (octx, decl, in_code);
+
+ for (; octx; octx = octx->outer_context)
+ {
+ if (octx->region_type & ORT_HOST_DATA)
+ continue;
+ if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
+ break;
+ splay_tree_node n2
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ if (n2)
+ {
+ flags |= GOVD_MAP;
+ goto found_outer;
+ }
+ }
+ }
- if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl), is_oacc))
+ if (!lang_hooks.types.omp_mappable_type
+ (TREE_TYPE (decl), ctx->region_kind == ORK_OACC))
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
flags |= GOVD_EXPLICIT;
}
+
+ if (is_oacc)
+ flags = oacc_default_clause (ctx, decl, in_code, flags);
+ else
+ flags |= GOVD_MAP;
+
+ found_outer:;
omp_add_variable (ctx, decl, flags);
}
else
next reply other threads:[~2015-07-29 23:13 UTC|newest]
Thread overview: 2+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-07-30 2:30 Nathan Sidwell [this message]
2015-08-03 11:07 ` Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=55B95DFD.6090703@acm.org \
--to=nathan@acm.org \
--cc=James_Norris@mentor.com \
--cc=gcc-patches@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).