diff mbox

[gomp4] acc enter/exit data

Message ID 871to796ey.fsf@kepler.schwinge.homeip.net
State New
Headers show

Commit Message

Thomas Schwinge Dec. 10, 2014, 9:54 a.m. UTC 
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(-)



Grüße,
 Thomas

Comments

Jakub Jelinek Dec. 10, 2014, 9:59 a.m. UTC | #1 
On Wed, Dec 10, 2014 at 10:54:13AM +0100, Thomas Schwinge wrote:
> --- 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))

I'd say that at this point a
  switch (gimple_omp_target_kind (last))
    {
    case GF_OMP_TARGET_KIND_UPDATE:
    case GF_OMP_TARGET_KIND_OACC_UPDATE:
    case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
      ...
    default:
      ...
    }
would be cleaner.  The first hunk is more questionable, because there is
else and it would require duplicating of the else body in default:, goto
or similar, but perhaps it would be better that way too.

	Jakub
diff mbox

Patch

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)
+  }
+}