[37/40] Fix for is_gimple_reg vars to 'data kernels'
Commit Message
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(-)
--
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
@@ -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. */
@@ -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 } }