From: James Norris <jnorris@codesourcery.com>
To: Jakub Jelinek <jakub@redhat.com>
Cc: GCC Patches <gcc-patches@gcc.gnu.org>
Subject: Re: [PATCH] Fix use of declare'd vars by routine procedures.
Date: Mon, 01 Feb 2016 16:00:00 -0000 [thread overview]
Message-ID: <56AF8133.7000004@codesourcery.com> (raw)
In-Reply-To: <20160129083151.GI3017@tucnak.redhat.com>
[-- Attachment #1: Type: text/plain, Size: 2158 bytes --]
Hi!
On 01/29/2016 02:31 AM, Jakub Jelinek wrote:
> On Thu, Jan 28, 2016 at 12:26:38PM -0600, James Norris wrote:
>> I think the attached change is what you had in mind with
>> regard to doing the check at gimplification time.
>
> Nope, this is still a wrong location for that.
> If you look at the next line after the block you've added, you'll see
> if (gimplify_omp_ctxp && omp_notice_variable (gimplify_omp_ctxp, decl, true))
> And that function fairly early calls is_global_var (decl). So you know
> already gimplify_omp_ctxp && is_global_var (decl), just put the rest into
> that block.
When said you said "rest into that block", I assumed you meant
the block in omp_notice_variable () that has the function call
to is_global_var () as the if condition.
> The TREE_CODE (decl) == VAR_DECL check is VAR_P (decl).
Fixed.
> What do you want to achieve with
>
>> + && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
>> + || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl))))
>
> ? is_global_var already guarantees you that it is either TREE_STATIC
> or DECL_EXTERNAL, why is that not good enough?
As you pointed out above, the check is not necessary.
>
>> [snip snip]
>>
>> + /* Validate variable for use within routine function. */
>> + if (gimplify_omp_ctxp && gimplify_omp_ctxp->region_type == ORT_TARGET
>> + && get_oacc_fn_attrib (current_function_decl)
>
> If you only care about the implicit target region of acc routine, I think
> you also want to check that gimplify_omp_ctxp->outer_context == NULL.
Check added.
>
>> [snip snip]
>
> And I'm really confused by this error message. If you are complaining that
> the variable is not listed in acc declare clauses, why don't you say that?
> What does the error have to do with its storage class?
> Also, splitting one error into two is weird, usually there would be one
> error message and perhaps inform after it.
Error messages re-written to make better sense.
Thanks!
Jim
===== ChangeLog entries....
gcc/
* gimplify.c (omp_notice_variable): Add variable usage check.
gcc/testsuite/
* c-c++-common/goacc/routine-5.c: Add tests.
[-- Attachment #2: declare.patch --]
[-- Type: text/x-patch, Size: 4452 bytes --]
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9f99842..41cd7a7 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,7 @@
+2016-01-XX James Norris <jnorris@codesourcery.com>
+
+ * gimplify.c (omp_notice_variable): Add usage check.
+
2016-01-29 Jakub Jelinek <jakub@redhat.com>
PR target/69551
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 32bc1fd..5833605 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6087,9 +6087,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if (ctx->region_type == ORT_NONE)
return lang_hooks.decls.omp_disregard_value_expr (decl, false);
- /* Threadprivate variables are predetermined. */
if (is_global_var (decl))
{
+ /* Threadprivate variables are predetermined. */
if (DECL_THREAD_LOCAL_P (decl))
return omp_notice_threadprivate_variable (ctx, decl, NULL_TREE);
@@ -6100,6 +6100,30 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
return omp_notice_threadprivate_variable (ctx, decl, value);
}
+
+ if (gimplify_omp_ctxp->outer_context == NULL
+ && VAR_P (decl)
+ && get_oacc_fn_attrib (current_function_decl))
+ {
+ location_t loc = DECL_SOURCE_LOCATION (decl);
+
+ if (lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl)))
+ {
+ error_at (loc,
+ "%qE with %<link%> clause used in %<routine%> function",
+ DECL_NAME (decl));
+ return false;
+ }
+ else if (!lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl)))
+ {
+ error_at (loc,
+ "%qE requires a %<declare%> directive for use "
+ "in a %<routine%> function", DECL_NAME (decl));
+ return false;
+ }
+ }
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
@@ -8223,7 +8247,7 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
if (oacc_declare_returns == NULL)
oacc_declare_returns = new hash_map<tree, tree>;
- oacc_declare_returns->put (decl, c);
+ oacc_declare_returns->put (decl, c);
}
}
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 150ebc8..61003c2 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,7 @@
+2016-01-XX James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/routine-5.c: Add tests.
+
2016-01-29 Jakub Jelinek <jakub@redhat.com>
PR target/69551
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index ccda097..c34838f 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -45,3 +45,97 @@ using namespace g;
#pragma acc routine (a) /* { dg-error "does not refer to" } */
#pragma acc routine (c) /* { dg-error "does not refer to" } */
+
+int vb1; /* { dg-error "directive for use" } */
+extern int vb2; /* { dg-error "directive for use" } */
+static int vb3; /* { dg-error "directive for use" } */
+
+#pragma acc routine
+int
+func1 (int a)
+{
+ vb1 = a + 1;
+ vb2 = vb1 + 1;
+ vb3 = vb2 + 1;
+
+ return vb3;
+}
+
+#pragma acc routine
+int
+func2 (int a)
+{
+ extern int vb4; /* { dg-error "directive for use" } */
+ static int vb5; /* { dg-error "directive for use" } */
+
+ vb4 = a + 1;
+ vb5 = vb4 + 1;
+
+ return vb5;
+}
+
+extern int vb6; /* { dg-error "clause used in" } */
+#pragma acc declare link (vb6)
+static int vb7; /* { dg-error "clause used in" } */
+#pragma acc declare link (vb7)
+
+#pragma acc routine
+int
+func3 (int a)
+{
+ vb6 = a + 1;
+ vb7 = vb6 + 1;
+
+ return vb7;
+}
+
+int vb8;
+#pragma acc declare create (vb8)
+extern int vb9;
+#pragma acc declare create (vb9)
+static int vb10;
+#pragma acc declare create (vb10)
+
+#pragma acc routine
+int
+func4 (int a)
+{
+ vb8 = a + 1;
+ vb9 = vb8 + 1;
+ vb10 = vb9 + 1;
+
+ return vb10;
+}
+
+int vb11;
+#pragma acc declare device_resident (vb11)
+extern int vb12;
+#pragma acc declare device_resident (vb12)
+extern int vb13;
+#pragma acc declare device_resident (vb13)
+
+#pragma acc routine
+int
+func5 (int a)
+{
+ vb11 = a + 1;
+ vb12 = vb11 + 1;
+ vb13 = vb12 + 1;
+
+ return vb13;
+}
+
+#pragma acc routine
+int
+func6 (int a)
+{
+ extern int vb14;
+#pragma acc declare create (vb14)
+ static int vb15;
+#pragma acc declare create (vb15)
+
+ vb14 = a + 1;
+ vb15 = vb14 + 1;
+
+ return vb15;
+}
next prev parent reply other threads:[~2016-02-01 16:00 UTC|newest]
Thread overview: 6+ messages / expand[flat|nested] mbox.gz Atom feed top
2016-01-17 23:01 James Norris
2016-01-22 10:39 ` Jakub Jelinek
2016-01-28 18:26 ` James Norris
2016-01-29 8:32 ` Jakub Jelinek
2016-02-01 16:00 ` James Norris [this message]
2016-02-01 16:06 ` Jakub Jelinek
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=56AF8133.7000004@codesourcery.com \
--to=jnorris@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
/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).