backport: "Fix oacc kernels default mapping for scalars"
2015-12-02 Tom de Vries <tom@codesourcery.com>
backport from trunk:
* 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.
---
gcc/gimplify.c | 21 +++++++++++++++------
.../c-c++-common/goacc/kernels-default-2.c | 17 +++++++++++++++++
gcc/testsuite/c-c++-common/goacc/kernels-default.c | 14 ++++++++++++++
gcc/testsuite/gfortran.dg/goacc/reduction-2.f95 | 2 +-
4 files changed, 47 insertions(+), 7 deletions(-)
@@ -90,8 +90,11 @@ 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,
+
/* OpenACC deviceptr clause. */
- GOVD_USE_DEVPTR = 1 << 18,
+ GOVD_USE_DEVPTR = 1 << 19,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
@@ -5988,8 +5991,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;
@@ -7700,10 +7707,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)
{
new file mode 100644
@@ -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" } } */
new file mode 100644
@@ -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" } } */
@@ -17,6 +17,6 @@ end subroutine
! { dg-final { scan-tree-dump-times "target oacc_parallel firstprivate.a." 1 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop reduction..:a. private.p." 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "target oacc_kernels map.tofrom:a .len: 4.." 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "target oacc_kernels map.force_tofrom:a .len: 4.." 1 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop reduction..:a. private.k." 1 "gimple" } }