From: Cesar Philippidis <cesar@codesourcery.com>
To: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>,
Jakub Jelinek <jakub@redhat.com>
Subject: [PATCH] Fix PR80029
Date: Wed, 22 Mar 2017 13:40:00 -0000 [thread overview]
Message-ID: <278ce4b7-cc02-453b-10fc-8094e9078d9a@codesourcery.com> (raw)
[-- Attachment #1: Type: text/plain, Size: 309 bytes --]
In addition to resolving the memory leak involving OpenACC delare
clauses, this patch also corrects an ICE involving VLA variables as data
clause arguments used in acc declare constructs. More details can be
found here <https://gcc.gnu.org/ml/gcc-patches/2016-09/msg01630.html>.
Is this OK for trunk?
Cesar
[-- Attachment #2: pr80029.diff --]
[-- Type: text/x-patch, Size: 3846 bytes --]
2017-03-22 Cesar Philippidis <cesar@codesourcery.com>
PR c++/80029
gcc/
* gimplify.c (is_oacc_declared): New function.
(oacc_default_clause): Use it to set default flags for acc declared
variables inside parallel regions.
(gimplify_scan_omp_clauses): Strip firstprivate pointers for acc
declared variables.
(gimplify_oacc_declare): Gimplify the declare clauses. Add the
declare attribute to any decl as necessary.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/declare-vla.c: New test.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index fbf136f..26d35d1 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6787,6 +6787,16 @@ device_resident_p (tree decl)
return false;
}
+/* Return true if DECL has an ACC DECLARE attribute. */
+
+static bool
+is_oacc_declared (tree decl)
+{
+ tree t = TREE_CODE (decl) == MEM_REF ? TREE_OPERAND (decl, 0) : decl;
+ tree declared = lookup_attribute ("oacc declare target", DECL_ATTRIBUTES (t));
+ return declared != NULL_TREE;
+}
+
/* Determine outer default flags for DECL mentioned in an OMP region
but not declared in an enclosing clause.
@@ -6887,6 +6897,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
const char *rkind;
bool on_device = false;
+ bool declared = is_oacc_declared (decl);
tree type = TREE_TYPE (decl);
if (lang_hooks.decls.omp_privatize_by_reference (decl))
@@ -6917,7 +6928,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
case ORT_ACC_PARALLEL:
{
- if (on_device || AGGREGATE_TYPE_P (type))
+ if (on_device || AGGREGATE_TYPE_P (type) || declared)
/* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
else
@@ -7346,6 +7357,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_DECLARE:
case OACC_HOST_DATA:
ctx->target_firstprivatize_array_bases = true;
default:
@@ -9231,18 +9243,26 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
{
tree expr = *expr_p;
gomp_target *stmt;
- tree clauses, t;
+ tree clauses, t, decl;
clauses = OACC_DECLARE_CLAUSES (expr);
gimplify_scan_omp_clauses (&clauses, pre_p, ORT_TARGET_DATA, OACC_DECLARE);
+ gimplify_adjust_omp_clauses (pre_p, NULL, &clauses, OACC_DECLARE);
for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
{
- tree decl = OMP_CLAUSE_DECL (t);
+ decl = OMP_CLAUSE_DECL (t);
if (TREE_CODE (decl) == MEM_REF)
- continue;
+ decl = TREE_OPERAND (decl, 0);
+
+ if (VAR_P (decl) && !is_oacc_declared (decl))
+ {
+ tree attr = get_identifier ("oacc declare target");
+ DECL_ATTRIBUTES (decl) = tree_cons (attr, NULL_TREE,
+ DECL_ATTRIBUTES (decl));
+ }
if (VAR_P (decl)
&& !is_global_var (decl)
@@ -9258,7 +9278,8 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
}
}
- omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
+ if (gimplify_omp_ctxp)
+ omp_add_variable (gimplify_omp_ctxp, decl, GOVD_SEEN);
}
stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_OACC_DECLARE,
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
new file mode 100644
index 0000000..3ea148e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-vla.c
@@ -0,0 +1,25 @@
+/* Verify that acc declare accept VLA variables. */
+
+#include <assert.h>
+
+int
+main ()
+{
+ int N = 1000;
+ int i, A[N];
+#pragma acc declare copy(A)
+
+ for (i = 0; i < N; i++)
+ A[i] = -i;
+
+#pragma acc kernels
+ for (i = 0; i < N; i++)
+ A[i] = i;
+
+#pragma acc update host(A)
+
+ for (i = 0; i < N; i++)
+ assert (A[i] == i);
+
+ return 0;
+}
next prev reply other threads:[~2017-03-22 13:40 UTC|newest]
Thread overview: 7+ messages / expand[flat|nested] mbox.gz Atom feed top
2016-09-22 23:39 [openacc] fix an ICE with acc declared VLAs Cesar Philippidis
2016-09-23 15:19 ` [gomp4] " Cesar Philippidis
2017-03-22 13:40 ` Cesar Philippidis [this message]
2017-03-22 13:43 ` [PATCH] Fix PR80029 Jakub Jelinek
2017-03-22 13:48 ` Cesar Philippidis
2017-03-22 13:50 ` Jakub Jelinek
2017-04-07 9:53 ` 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=278ce4b7-cc02-453b-10fc-8094e9078d9a@codesourcery.com \
--to=cesar@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).