From db0f1c7a906bd8798a4016c1622f5999897e7469 Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Wed, 2 Dec 2015 15:48:45 +0000 Subject: [PATCH] Fix oacc kernels default mapping for scalars 2015-12-02 Tom de Vries * gimplify.c (enum gimplify_omp_var_data): Add enum value GOVD_MAP_FORCE. (oacc_default_clause): Fix default for scalars in oacc kernels. (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE. * c-c++-common/goacc/kernels-default-2.c: New test. * c-c++-common/goacc/kernels-default.c: New test. From-SVN: r231183 --- gcc/ChangeLog | 7 +++++++ gcc/gimplify.c | 19 ++++++++++++++----- gcc/testsuite/ChangeLog | 5 +++++ .../c-c++-common/goacc/kernels-default-2.c | 17 +++++++++++++++++ .../c-c++-common/goacc/kernels-default.c | 14 ++++++++++++++ 5 files changed, 57 insertions(+), 5 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-default-2.c create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-default.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 0e4a3dd82da..ef5fbefd7b4 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,10 @@ +2015-12-02 Tom de Vries + + * gimplify.c (enum gimplify_omp_var_data): Add enum value + GOVD_MAP_FORCE. + (oacc_default_clause): Fix default for scalars in oacc kernels. + (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE. + 2015-12-02 Tom de Vries * omp-low.c (install_var_field, scan_sharing_clauses): Add and handle diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 4229e2d52dd..7146a01a5af 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -90,6 +90,9 @@ enum gimplify_omp_var_data /* Flag for shared vars that are or might be stored to in the region. */ GOVD_WRITTEN = 131072, + /* Flag for GOVD_MAP, if it is a forced mapping. */ + GOVD_MAP_FORCE = 262144, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -5980,8 +5983,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) gcc_unreachable (); case ORT_ACC_KERNELS: - /* Everything under kernels are default 'present_or_copy'. */ + /* Scalars are default 'copy' under kernels, non-scalars are default + 'present_or_copy'. */ flags |= GOVD_MAP; + if (!AGGREGATE_TYPE_P (TREE_TYPE (decl))) + flags |= GOVD_MAP_FORCE; + rkind = "kernels"; break; @@ -7640,10 +7647,12 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) } else if (code == OMP_CLAUSE_MAP) { - OMP_CLAUSE_SET_MAP_KIND (clause, - flags & GOVD_MAP_TO_ONLY - ? GOMP_MAP_TO - : GOMP_MAP_TOFROM); + int kind = (flags & GOVD_MAP_TO_ONLY + ? GOMP_MAP_TO + : GOMP_MAP_TOFROM); + if (flags & GOVD_MAP_FORCE) + kind |= GOMP_MAP_FLAG_FORCE; + OMP_CLAUSE_SET_MAP_KIND (clause, kind); if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog index bb523db54b1..5fe26bb9aa0 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,8 @@ +2015-12-02 Tom de Vries + + * c-c++-common/goacc/kernels-default-2.c: New test. + * c-c++-common/goacc/kernels-default.c: New test. + 2015-12-02 Tom de Vries * c-c++-common/goacc/kernels-alias-2.c: New test. diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c new file mode 100644 index 00000000000..232b1236f38 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default-2.c @@ -0,0 +1,17 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 2 + +void +foo (void) +{ + unsigned int a[N]; + +#pragma acc kernels + { + a[0]++; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(tofrom" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-default.c b/gcc/testsuite/c-c++-common/goacc/kernels-default.c new file mode 100644 index 00000000000..58cd5e10a5b --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-default.c @@ -0,0 +1,14 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void +foo (void) +{ + unsigned int i; +#pragma acc kernels + { + i++; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(force_tofrom" 1 "gimple" } } */ -- 2.30.2