aboutsummaryrefslogtreecommitdiff
path: root/gcc
diff options
context:
space:
mode:
authorNathan Sidwell <nathan@codesourcery.com>2015-11-12 13:51:13 +0000
committerNathan Sidwell <nathan@gcc.gnu.org>2015-11-12 13:51:13 +0000
commitfffeedeb5aefd79d90fd9a2b331fe095965ffbd2 (patch)
tree86db058605686d31d1e68578774a9ac54758d69e /gcc
parent8339a33e54150c0d643593c9358aa35f4c3ffd5e (diff)
downloadgcc-fffeedeb5aefd79d90fd9a2b331fe095965ffbd2.zip
gcc-fffeedeb5aefd79d90fd9a2b331fe095965ffbd2.tar.gz
gcc-fffeedeb5aefd79d90fd9a2b331fe095965ffbd2.tar.bz2
gimplify.c (oacc_default_clause): New.
gcc/ * gimplify.c (oacc_default_clause): New. (omp_notice_variable): Call it. gcc/testsuite/ * c-c++-common/goacc/data-default-1.c: New. libgomp/ * testsuite/libgomp.oacc-c-c++-common/default-1.c: New. Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com> From-SVN: r230256
Diffstat (limited to 'gcc')
-rw-r--r--gcc/ChangeLog16
-rw-r--r--gcc/gimplify.c61
-rw-r--r--gcc/testsuite/ChangeLog8
-rw-r--r--gcc/testsuite/c-c++-common/goacc/data-default-1.c37
4 files changed, 115 insertions, 7 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b7c3df7..7089985 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,15 @@
+2015-11-12 Nathan Sidwell <nathan@codesourcery.com>
+
+ gcc/
+ * gimplify.c (oacc_default_clause): New.
+ (omp_notice_variable): Call it.
+
+ gcc/testsuite/
+ * c-c++-common/goacc/data-default-1.c: New.
+
+ libgomp/
+ * testsuite/libgomp.oacc-c-c++-common/default-1.c: New.
+
2015-11-12 Ilya Enkovich <enkovich.gnu@gmail.com>
PR tree-optimization/68305
@@ -311,7 +323,7 @@
2015-11-11 Nathan Sidwell <nathan@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
- * gcc/gimplify.c (enum omp_region_type): Add ORT_ACC,
+ * gimplify.c (enum omp_region_type): Add ORT_ACC,
ORT_ACC_DATA, ORT_ACC_PARALLEL, ORT_ACC_KERNELS. Adjust ORT_NONE.
(gimple_add_tmp_var): Add ORT_ACC checks.
(gimplify_var_or_parm_decl): Likewise.
@@ -327,7 +339,7 @@
(gimplify_oacc_cache): Specify ORT_ACC.
(gimplify_omp_workshare): Adjust OpenACC region types.
(gimplify_omp_target_update): Likewise.
- * gcc/omp-low.c (scan_sharing_clauses): Remove Openacc
+ * omp-low.c (scan_sharing_clauses): Remove Openacc
firstprivate sorry.
(lower-rec_input_clauses): Don't handle openacc firstprivate
references here.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 66e5168..74d8765 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -5900,6 +5900,60 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
return flags;
}
+
+/* Determine outer default flags for DECL mentioned in an OACC region
+ but not declared in an enclosing clause. */
+
+static unsigned
+oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
+{
+ const char *rkind;
+
+ switch (ctx->region_type)
+ {
+ default:
+ gcc_unreachable ();
+
+ case ORT_ACC_KERNELS:
+ /* Everything under kernels are default 'present_or_copy'. */
+ flags |= GOVD_MAP;
+ rkind = "kernels";
+ break;
+
+ case ORT_ACC_PARALLEL:
+ {
+ tree type = TREE_TYPE (decl);
+
+ if (TREE_CODE (type) == REFERENCE_TYPE
+ || POINTER_TYPE_P (type))
+ type = TREE_TYPE (type);
+
+ if (AGGREGATE_TYPE_P (type))
+ /* Aggregates default to 'present_or_copy'. */
+ flags |= GOVD_MAP;
+ else
+ /* Scalars default to 'firstprivate'. */
+ flags |= GOVD_FIRSTPRIVATE;
+ rkind = "parallel";
+ }
+ break;
+ }
+
+ if (DECL_ARTIFICIAL (decl))
+ ; /* We can get compiler-generated decls, and should not complain
+ about them. */
+ else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
+ {
+ error ("%qE not specified in enclosing OpenACC %s construct",
+ DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
+ error_at (ctx->location, "enclosing OpenACC %s construct", rkind);
+ }
+ else
+ gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
+
+ return flags;
+}
+
/* Record the fact that DECL was used within the OMP context CTX.
IN_CODE is true when real code uses DECL, and false when we should
merely emit default(none) errors. Return true if DECL is going to
@@ -6023,7 +6077,12 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
nflags |= GOVD_MAP | GOVD_EXPLICIT;
}
else if (nflags == flags)
- nflags |= GOVD_MAP;
+ {
+ if ((ctx->region_type & ORT_ACC) != 0)
+ nflags = oacc_default_clause (ctx, decl, flags);
+ else
+ nflags |= GOVD_MAP;
+ }
}
found_outer:
omp_add_variable (ctx, decl, nflags);
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 7754469..a7395cf 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,7 @@
+2015-11-12 Nathan Sidwell <nathan@codesourcery.com>
+
+ * c-c++-common/goacc/data-default-1.c: New.
+
2015-11-12 David Edelsohn <dje.gcc@gmail.com>
* gcc.target/powerpc/pr67789.c: Skip on AIX and Darwin.
@@ -82,10 +86,6 @@
* gcc.dg/tree-ssa/pr68234.c: New testcase.
-2015-11-10 Nathan Sidwell <nathan@codesourcery.com>
-
- * gcc.dg/goacc/nvptx-opt-1.c: New test.
-
2015-11-10 Ilya Enkovich <enkovich.gnu@gmail.com>
* gcc.target/i386/mask-pack.c: New test.
diff --git a/gcc/testsuite/c-c++-common/goacc/data-default-1.c b/gcc/testsuite/c-c++-common/goacc/data-default-1.c
new file mode 100644
index 0000000..0741abc
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/data-default-1.c
@@ -0,0 +1,37 @@
+/* { dg-do compile } */
+
+
+int main ()
+{
+ int n = 2;
+ int ary[2];
+
+#pragma acc parallel default (none) /* { dg-message "parallel construct" 2 } */
+ {
+ ary[0] /* { dg-error "not specified in enclosing" } */
+ = n; /* { dg-error "not specified in enclosing" } */
+ }
+
+#pragma acc kernels default (none) /* { dg-message "kernels construct" 2 } */
+ {
+ ary[0] /* { dg-error "not specified in enclosing" } */
+ = n; /* { dg-error "not specified in enclosing" } */
+ }
+
+#pragma acc data copy (ary, n)
+ {
+#pragma acc parallel default (none)
+ {
+ ary[0]
+ = n;
+ }
+
+#pragma acc kernels default (none)
+ {
+ ary[0]
+ = n;
+ }
+ }
+
+ return 0;
+}