diff options
author | Nathan Sidwell <nathan@codesourcery.com> | 2015-11-12 13:51:13 +0000 |
---|---|---|
committer | Nathan Sidwell <nathan@gcc.gnu.org> | 2015-11-12 13:51:13 +0000 |
commit | fffeedeb5aefd79d90fd9a2b331fe095965ffbd2 (patch) | |
tree | 86db058605686d31d1e68578774a9ac54758d69e /gcc | |
parent | 8339a33e54150c0d643593c9358aa35f4c3ffd5e (diff) | |
download | gcc-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/ChangeLog | 16 | ||||
-rw-r--r-- | gcc/gimplify.c | 61 | ||||
-rw-r--r-- | gcc/testsuite/ChangeLog | 8 | ||||
-rw-r--r-- | gcc/testsuite/c-c++-common/goacc/data-default-1.c | 37 |
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; +} |