Re: [gomp4] acc enter/exit data

[Thomas Schwinge]

Hi!

On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis <ce...@codesourcery.com> 
wrote:
> This patch add support for OpenACC's enter/exit data directive. [...]

>       gcc/
>       * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.

In r218567, I committed the following to gomp-4_0-branch:

commit 86724db93ad780106102573f2cfadd6f884e8650
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed Dec 10 09:52:14 2014 +0000

    Fix OpenACC enter/exit data ICE.
    
        [...]: In function 'f_acc_data':
        [...]:4:1: internal compiler error: in expand_gimple_stmt_1, at 
cfgexpand.c:3413
         f_acc_data (void)
         ^
        0x70cad3 expand_gimple_stmt_1
                [...]/source-gcc/gcc/cfgexpand.c:3413
        0x70cad3 expand_gimple_stmt
                [...]/source-gcc/gcc/cfgexpand.c:3440
        0x712b3d expand_gimple_basic_block
                [...]/source-gcc/gcc/cfgexpand.c:5273
        0x71479e execute
                [...]/source-gcc/gcc/cfgexpand.c:5882
    
        gcc/
        * omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
        <GIMPLE_OMP_TARGET>: Handle
        GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA the same as
        GF_OMP_TARGET_KIND_OACC_UPDATE.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218567 
138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                           |  7 +++++++
 gcc/omp-low.c                                |  8 ++++++--
 gcc/testsuite/c-c++-common/goacc/nesting-2.c | 11 +++++++++++
 3 files changed, 24 insertions(+), 2 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index af59ada..bece7c1 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2014-12-10  Thomas Schwinge  <tho...@codesourcery.com>
+
+       * omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
+       <GIMPLE_OMP_TARGET>: Handle
+       GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA the same as
+       GF_OMP_TARGET_KIND_OACC_UPDATE.
+
 2014-11-13  Cesar Philippidis  <ce...@codesourcery.com>
 
        * omp-low.c (oacc_get_reduction_array_id): Fix whitespace.
diff --git gcc/omp-low.c gcc/omp-low.c
index 9af3b8a..6fed38f 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -9404,7 +9404,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region 
*parent,
       else if (code == GIMPLE_OMP_TARGET
               && (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE
                   || (gimple_omp_target_kind (stmt)
-                      == GF_OMP_TARGET_KIND_OACC_UPDATE)))
+                      == GF_OMP_TARGET_KIND_OACC_UPDATE)
+                  || (gimple_omp_target_kind (stmt)
+                      == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)))
        new_omp_region (bb, code, parent);
       else
        {
@@ -12270,7 +12272,9 @@ make_gimple_omp_edges (basic_block bb, struct 
omp_region **region,
       cur_region = new_omp_region (bb, code, cur_region);
       fallthru = true;
       if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
-         || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE)
+         || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE
+         || (gimple_omp_target_kind (last)
+             == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA))
        cur_region = cur_region->outer;
       break;
 
diff --git gcc/testsuite/c-c++-common/goacc/nesting-2.c 
gcc/testsuite/c-c++-common/goacc/nesting-2.c
new file mode 100644
index 0000000..0d350c6
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-2.c
@@ -0,0 +1,11 @@
+int i;
+
+void
+f_acc_data (void)
+{
+#pragma acc data
+  {
+#pragma acc update host(i)
+#pragma acc enter data copyin(i)
+  }
+}


Grüße,
 Thomas

signature.asc
Description: PGP signature