Hi!
On 01/11/2016 05:55 AM, Thomas Schwinge wrote:
Hi!
On Thu, 7 Jan 2016 12:57:32 -0600, James Norris <jnor...@codesourcery.com>
wrote:
The checking of variables declared by OpenACC declare directives
used within an OpenACC routine procedure was not being done correctly.
This fix adds the checking required and also adds to the testing.
This fix resolves the issue pointed out by Cesar with declare-4.c
(https://gcc.gnu.org/ml/gcc-patches/2016-01/msg00339.html).
This patch has been applied to gomp-4_0-branch.
Such a patch needs to go into trunk; see my report in
<http://news.gmane.org/find-root.php?message_id=%3C87mvtapdp0.fsf%40kepler.schwinge.homeip.net%3E>.
--- gcc/c/c-parser.c (revision 232138)
+++ gcc/c/c-parser.c (working copy)
@@ -14115,6 +14115,10 @@
/* Also add an "omp declare target" attribute, with clauses. */
DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"),
clauses, DECL_ATTRIBUTES (fndecl));
+
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("oacc routine"),
+ clauses, DECL_ATTRIBUTES (fndecl));
}
Yuck, another attribute... (..., and it's not listed/documented in
gcc/c-family/c-common.c:c_common_attribute_table.)
Yuck is right. In the interim I've found a function: get_oacc_fn_attrib,
in omp-low.c, that seems to suit my needs. So I'll be revising to use
that one instead.
You store clauses in the "oacc routine" here, but it's unused as far as I
can tell.
Given that we have the clauses available from the "omp declare target"
attribute (subject to change to switching to a generic "omp clauses"
attribute as suggested in
<http://news.gmane.org/find-root.php?message_id=%3C87twns3ebs.fsf%40hertz.schwinge.homeip.net%3E>),
could we maybe just look up some specific clause instead of detecting the
presence of this extra attribute? (Jakub, any preference?) Anyway, have
we verified that the desired behavior:
--- gcc/c/c-typeck.c (revision 232138)
+++ gcc/c/c-typeck.c (working copy)
@@ -2664,6 +2664,26 @@
tree ref;
tree decl = lookup_name (id);
+ if (decl
+ && decl != error_mark_node
+ && current_function_decl
+ && TREE_CODE (decl) == VAR_DECL
+ && is_global_var (decl)
+ && lookup_attribute ("oacc routine",
+ DECL_ATTRIBUTES (current_function_decl)))
+ {
+ if (lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl))
+ || ((!lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl))
+ && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
+ || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl))))))
+ {
+ error_at (loc, "invalid use in %<routine%> function");
+ return error_mark_node;
+ }
+ }
..., isn't applicable to OpenMP as well (thus, no "oacc routine"
attribute conditional is needed here)?
This appears to be a situation where OpenMP and OpenACC diverge,
so I need to discriminate between OpenACC routine and OpenMP
declare target.
--- gcc/cp/parser.c (revision 232138)
+++ gcc/cp/parser.c (working copy)
@@ -36732,6 +36732,10 @@
DECL_ATTRIBUTES (fndecl)
= tree_cons (get_identifier ("omp declare target"),
clauses, DECL_ATTRIBUTES (fndecl));
+
+ DECL_ATTRIBUTES (fndecl)
+ = tree_cons (get_identifier ("oacc routine"),
+ NULL_TREE, DECL_ATTRIBUTES (fndecl));
}
You don't store clauses in the "oacc routine" here.
--- gcc/cp/semantics.c (revision 232138)
+++ gcc/cp/semantics.c (working copy)
@@ -3700,6 +3700,25 @@
decl = convert_from_reference (decl);
}
+
+ if (decl != error_mark_node
+ && current_function_decl
+ && TREE_CODE (decl) == VAR_DECL
+ && is_global_var (decl)
+ && lookup_attribute ("oacc routine",
+ DECL_ATTRIBUTES (current_function_decl)))
+ {
+ if (lookup_attribute ("omp declare target link",
+ DECL_ATTRIBUTES (decl))
+ || ((!lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl))
+ && ((TREE_STATIC (decl) && !DECL_EXTERNAL (decl))
+ || (!TREE_STATIC (decl) && DECL_EXTERNAL (decl))))))
+ {
+ *error_msg = "invalid use in %<routine%> function";
+ return error_mark_node;
+ }
+ }
Likewise.
No equivalent change is needed for Fortran?
No. C/C++ statics and extern's don't appear in Fortran.
--- libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c (revision
232138)
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c (working copy)
@@ -4,7 +4,7 @@
#include <openacc.h>
float b;
-#pragma acc declare link (b)
+#pragma acc declare create (b)
#pragma acc routine
int
I have not verified the details, but a very similar fix was required to
get rid of a number of regressions:
@@ -2637,18 +2637,18 @@ PASS: c-c++-common/goacc-gomp/nesting-fail-1.c
(test for errors, line 350)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 356)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 358)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 360)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 371)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 378)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 386)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 395)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 402)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 409)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 416)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 42)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 423)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 430)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 432)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
errors, line 434)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 47)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 52)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 57)
@@ -2667,7 +2667,7 @@ PASS: c-c++-common/goacc-gomp/nesting-fail-1.c
(test for errors, line 93)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 95)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 97)
PASS: c-c++-common/goacc-gomp/nesting-fail-1.c (test for errors, line 99)
[-PASS:-]{+FAIL:+} c-c++-common/goacc-gomp/nesting-fail-1.c (test for
excess errors)
Same for C++.
Committed to gomp-4_0-branch in r232219:
commit 1ac87f2b59dd03cb305ec94a7c6b5657dbb54e66
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon Jan 11 11:51:57 2016 +0000
Resolve c-c++-common/goacc-gomp/nesting-fail-1.c regressions
gcc/testsuite/
* c-c++-common/goacc-gomp/nesting-fail-1.c: Add OpenACC declare
directive for "i".
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@232219
138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/testsuite/ChangeLog.gomp | 5 +++++
gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c | 1 +
2 files changed, 6 insertions(+)
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 607ca8e..2db11df 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2016-01-11 Thomas Schwinge <tho...@codesourcery.com>
+
+ * c-c++-common/goacc-gomp/nesting-fail-1.c: Add OpenACC declare
+ directive for "i".
+
2016-01-07 James Norris <jnor...@codesourcery.com>
* c-c++-common/goacc/routine-5.c: Additional tests.
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 8e0f217..9011fcf 100644
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -1,4 +1,5 @@
extern int i;
+#pragma acc declare create(i)
void
f_omp (void)
Grüße
Thomas
Thank you, thank you,
Jim
