Fix oacc kernels default mapping for scalars

2015-12-02  Tom de Vries  <tom@codesourcery.com>

	* 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.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@231183 138bc75d-0d04-0410-961f-82ee72b054a4

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 0e4a3dd82daae73a2931c94380635dc10027bc03..ef5fbefd7b4830ed8b9283bedd3ac3c5009c1cd9 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,10 @@
+2015-12-02  Tom de Vries  <tom@codesourcery.com>
+
+       * 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  <tom@codesourcery.com>
 
        * omp-low.c (install_var_field, scan_sharing_clauses): Add and handle

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 4229e2d52dd049088db1787519dcbfb6a462f45d..7146a01a5afe0656f1fca3beb5c9a4129a1641ac 100644 (file)
--- 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 bb523db54b17f10d38de6baed66a128a965d6208..5fe26bb9aa00dd9fe9e7c0ddfe9a465e35e5c045 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2015-12-02  Tom de Vries  <tom@codesourcery.com>
+
+       * 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  <tom@codesourcery.com>
 
        * 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 (file)
index 0000000..232b123
--- /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 (file)
index 0000000..58cd5e1
--- /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" } } */