[PATCH 37/40] Fix for is_gimple_reg vars to 'data kernels'

[Frederik Harwath <frederik@codesourcery.com>]

From: Tobias Burnus <tobias@codesourcery.com>

Nearly all variable mapping is moved from 'kernels' to a surrounding
'data kernels' and then 'force_present' mapped for the 'kernels'. However, as
libgomp.oacc-c-c++-common/declare-vla.c shows, moving 'int i, N' will fail as
there is a special case for is_gimple_reg in mapping and that fails badly if
outside a target region (e.g. offloading = false). As those are transferred by
value and not as a pointer, it makes more sense to only map them at
'kernels' and ignore them for 'data kernels'.
Additionally, as e.g. libgomp.oacc-c-c++-common/kernels-decompose-1.c shows,
one still additionally to handle 'kernels'-declared variables which now are
declared in 'kernels data' and and can be handled as is_gimple_reg.

        gcc/
        * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region):
        is_gimple_reg vars are not yet mapped, fall through to map is as
        before the transformation.
        (omp_oacc_kernels_decompose_1): Don't map is_gimple_reg vars.
        (decompose_kernels_region_body): Use tofrom for is_gimple_reg vars.
        (omp_oacc_kernels_decompose_1): Handle is_gimple_reg vars as without
        data kernels.

        gcc/testsuite/
        * gfortran.dg/goacc/declare-3.f95: Update scan-tree-dump-times.
---
 gcc/omp-oacc-kernels-decompose.cc             | 9 +++++++--
 gcc/testsuite/gfortran.dg/goacc/declare-3.f95 | 2 +-
 2 files changed, 8 insertions(+), 3 deletions(-)

diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index c96207d96250..a6be1f1ed238 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -873,7 +873,7 @@ maybe_build_inner_data_region (location_t loc, gimple *body,
          else
            inner_bind_vars = next;
        }
-      else
+      else if (!is_gimple_reg (v))
        {
          /* Otherwise, build the map clause.  */
          tree new_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
@@ -1222,7 +1222,9 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses)
       if (!DECL_ARTIFICIAL (var) && TREE_CODE (var) != CONST_DECL)
        {
          tree present_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
-         OMP_CLAUSE_SET_MAP_KIND (present_clause, GOMP_MAP_FORCE_PRESENT);
+         OMP_CLAUSE_SET_MAP_KIND (present_clause,
+                                  is_gimple_reg (var)
+                                  ? GOMP_MAP_TOFROM : GOMP_MAP_FORCE_PRESENT);
          OMP_CLAUSE_DECL (present_clause) = var;
          OMP_CLAUSE_SIZE (present_clause) = DECL_SIZE_UNIT (var);
          OMP_CLAUSE_CHAIN (present_clause) = present_clauses;
@@ -1437,6 +1439,9 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt)
                   region causes runtime errors.  */
                break;

+             if (is_gimple_reg (decl))
+               break;
+
              /* For non-artificial variables, and for non-declaration
                 expressions like A[0:n], copy the clause to the data
                 region.  */
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
index 9127cba6600d..2a1fe0a68465 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
@@ -39,7 +39,7 @@ program test
   use mod_d
   use mod_e

-  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(force_to:b\) map\(force_alloc:a\)$} original } }
+  ! { dg-final { scan-tree-dump {(?n)#pragma acc data map\(force_alloc:d\) map\(to:b\) map\(alloc:a\)$} original } }
 end program test

 ! { dg-final { scan-tree-dump-times {#pragma acc data} 1 original } }
--
2.33.0

-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955