From 7e9d0d7fec32a74c91f5d6ba65780ce682ec2b1a Mon Sep 17 00:00:00 2001 From: Tobias Burnus Date: Mon, 16 Mar 2020 16:22:57 +0100 Subject: [PATCH] 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. --- 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 409c0f41f5b1..bd5d7de4cd7b 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,13 @@ +2020-03-16 Tobias Burnus + + * 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 * 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 bceb03c422e0..6acb6367a7f1 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 88f0e426c553..6488f3bb60d6 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,5 +1,10 @@ +2020-03-17 Tobias Burnus + + * 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 - Julian Brown + Julian Brown * 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 5defe2ea85de..c21dfaeece2a 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 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.47.3