Fix for is_gimple_reg vars to 'data kernels'

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.c (maybe_build_inner_data_region): is_gimple_reg vars
	are not yet mapped, fall through to map is as before the transformation.
	(transform_kernels_region): Don't map is_gimple_reg vars.
	(decompose_kernels_region_body): Use tofrom for is_gimple_reg vars.
	(transform_kernels_region): Handle is_gimple_reg vars as without
	data kernels.

	gcc/testsuite/
	* gfortran.dg/goacc/common-block-3.f90: Update scan-tree-dump-times.
	* gfortran.dg/goacc/declare-3.f95: Update scan-tree-dump-times.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index a9689b77fbe9c64ca08ed38c7cfcfebd5c044eb4..5db79be770030477aa68b83ec034c2095ea75089 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2020-03-16  Tobias Burnus  <tobias@codesourcery.com>
+
+       * omp-oacc-kernels.c (maybe_build_inner_data_region): is_gimple_reg vars
+       are not yet mapped, fall through to map is as before the transformation.
+       (transform_kernels_region): Don't map is_gimple_reg vars.
+       (decompose_kernels_region_body): Use tofrom for is_gimple_reg vars.
+       (transform_kernels_region): Handle is_gimple_reg vars as without
+       data kernels.
+
 2020-03-12  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * omp-sese.c (install_var_field): Generate a field name for a VAR_DECL

diff --git a/gcc/omp-oacc-kernels.c b/gcc/omp-oacc-kernels.c
index a08a29d5ec190b60a8e0c26cee6dcfa98da1bcd1..fb36b50bbba7d96825fa225b6fdc153c9e01bba1 100644 (file)
--- a/gcc/omp-oacc-kernels.c
+++ b/gcc/omp-oacc-kernels.c
@@ -856,7 +856,7 @@ maybe_build_inner_data_region (location_t loc, gimple *body,
          else
            inner_bind_vars = next;
        }
-      else if (!idx_vars->contains (v))
+      else if (!idx_vars->contains (v) && !is_gimple_reg (v))
        {
          /* Otherwise, build the map clause.  */
          tree new_clause = build_omp_clause (loc, OMP_CLAUSE_MAP);
@@ -1217,7 +1217,9 @@ decompose_kernels_region_body (gimple *kernels_region, tree kernels_clauses,
          && !idx_vars.contains (var))
        {
          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;
@@ -1487,6 +1489,9 @@ transform_kernels_region (gimple *kernels_region)
                   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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index a38aac1d4f3b625bab547b56cb7e4c141c15e736..a21bf07eb78148715d5a66e2418fa6194c745589 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,5 +1,10 @@
+2020-03-17  Tobias Burnus  <tobias@codesourcery.com>
+
+       * gfortran.dg/goacc/common-block-3.f90: Update scan-tree-dump-times.
+       * gfortran.dg/goacc/declare-3.f95: Update scan-tree-dump-times.
+
 2018-10-04  Cesar Philippidis  <cesar@codesourcery.com>
-            Julian Brown  <julian@codesourcery.com>
+           Julian Brown  <julian@codesourcery.com>
 
        * gfortran.dg/goacc/declare-allocatable-1.f90: New test.
 

diff --git a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
index 5defe2ea85deb953e771b91652433a7cd73b7122..c21dfaeece2adc543e9510381ee4c7be3060f12c 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/common-block-3.f90
@@ -30,9 +30,11 @@ end program main
 ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:b \\\[len: 400\\\]\\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_parallel .*map\\(tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
 
+! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:x \\\[len: 400\\\]\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_present:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:i \\\[len: 4\\\]\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:x \\\[len: 400\\\]\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(tofrom:y \\\[len: 400\\\]\\\)" 1 "omplower" } }
 ! { dg-final { scan-tree-dump-times "omp target oacc_kernels .*map\\(force_tofrom:c \\\[len: 4\\\]\\)" 1 "omplower" } }
 
 ! Expecting no mapping of un-referenced common-blocks variables

diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-3.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-3.f95
index 9127cba6600d7acb18bda5db29bf57692545f033..2a1fe0a68465d751a8aa1cb88d4feb5bc04bded2 100644 (file)
--- 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 } }