public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Fix for is_gimple_reg vars to 'data kernels'
@ 2021-05-13 16:14 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:14 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:7e9d0d7fec32a74c91f5d6ba65780ce682ec2b1a
commit 7e9d0d7fec32a74c91f5d6ba65780ce682ec2b1a
Author: Tobias Burnus <tobias@codesourcery.com>
Date: Mon Mar 16 16:22:57 2020 +0100
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-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/common-block-3.f90: Update scan-tree-dump-times.
* gfortran.dg/goacc/declare-3.f95: Update scan-tree-dump-times.
Diff:
---
gcc/ChangeLog.omp | 10 ++++++++++
gcc/omp-oacc-kernels-decompose.cc | 9 +++++++--
gcc/testsuite/ChangeLog.omp | 7 ++++++-
gcc/testsuite/gfortran.dg/goacc/common-block-3.f90 | 6 ++++--
gcc/testsuite/gfortran.dg/goacc/declare-3.f95 | 2 +-
5 files changed, 28 insertions(+), 6 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 409c0f41f5b..bd5d7de4cd7 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,13 @@
+2020-03-16 Tobias Burnus <tobias@codesourcery.com>
+
+ * 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.
+
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-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index bceb03c422e..6acb6367a7f 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -870,7 +870,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);
@@ -1231,7 +1231,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;
@@ -1501,6 +1503,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 88f0e426c55..6488f3bb60d 100644
--- 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 5defe2ea85d..c21dfaeece2 100644
--- 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 9127cba6600..2a1fe0a6846 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 } }
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:14 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:14 [gcc/devel/omp/gcc-11] Fix for is_gimple_reg vars to 'data kernels' Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).