[gomp4] acc enter/exit data

[Thomas Schwinge]

Hi!

On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis <cesar@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  <thomas@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  <cesar@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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 472 bytes
Desc: not available
URL: <http://gcc.gnu.org/pipermail/gcc-patches/attachments/20141210/f35ce0dc/attachment.sig>