From fffeedeb5aefd79d90fd9a2b331fe095965ffbd2 Mon Sep 17 00:00:00 2001 From: Nathan Sidwell Date: Thu, 12 Nov 2015 13:51:13 +0000 Subject: [PATCH] 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 From-SVN: r230256 --- gcc/ChangeLog | 16 +++- gcc/gimplify.c | 61 ++++++++++++- gcc/testsuite/ChangeLog | 8 +- .../c-c++-common/goacc/data-default-1.c | 37 ++++++++ libgomp/ChangeLog | 4 + .../libgomp.oacc-c-c++-common/default-1.c | 87 +++++++++++++++++++ 6 files changed, 206 insertions(+), 7 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/data-default-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index b7c3df7830c..7089985c980 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,15 @@ +2015-11-12 Nathan Sidwell + + 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 PR tree-optimization/68305 @@ -311,7 +323,7 @@ 2015-11-11 Nathan Sidwell Cesar Philippidis - * 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 66e5168746f..74d8765950f 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 77544697fb7..a7395cf4f37 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,7 @@ +2015-11-12 Nathan Sidwell + + * c-c++-common/goacc/data-default-1.c: New. + 2015-11-12 David Edelsohn * 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 - - * gcc.dg/goacc/nvptx-opt-1.c: New test. - 2015-11-10 Ilya Enkovich * 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 00000000000..0741abca4ea --- /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; +} diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 406c5720325..b4f708e5e46 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,7 @@ +2015-11-12 Nathan Sidwell + + * testsuite/libgomp.oacc-c-c++-common/default-1.c: New. + 2015-11-1 Nathan Sidwell * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: New. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c new file mode 100644 index 00000000000..1ac0b9587b9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/default-1.c @@ -0,0 +1,87 @@ +/* { dg-do run } */ + +#include + +int test_parallel () +{ + int ok = 1; + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + + /* val defaults to firstprivate, ary defaults to copy. */ +#pragma acc parallel num_gangs (32) copy (ok) copy(ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop gang(static:1) + for (unsigned i = 0; i < 32; i++) + { + if (val != 2) + ok = 0; + val += i; + ary[i] = val; + } + } + + if (ondev) + { + if (!ok) + return 1; + if (val != 2) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} + +int test_kernels () +{ + int val = 2; + int ary[32]; + int ondev = 0; + + for (int i = 0; i < 32; i++) + ary[i] = ~0; + + /* val defaults to copy, ary defaults to copy. */ +#pragma acc kernels copy(ondev) + { + ondev = acc_on_device (acc_device_not_host); +#pragma acc loop + for (unsigned i = 0; i < 32; i++) + { + ary[i] = val; + val++; + } + } + + if (ondev) + { + if (val != 2 + 32) + return 1; + + for (int i = 0; i < 32; i++) + if (ary[i] != 2 + i) + return 1; + } + + return 0; +} + +int main () +{ + if (test_parallel ()) + return 1; + + if (test_kernels ()) + return 1; + + return 0; +} -- 2.30.2