Handle OpenACC 'self' clause for compute constructs in OpenACC 'kernels' decomposition

... to fix up recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a
"OpenACC 2.7: Implement self clause for compute constructs" for that case.

	gcc/
	* omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1):
	Handle 'OMP_CLAUSE_SELF' like 'OMP_CLAUSE_IF'.
	* omp-expand.cc (expand_omp_target): Handle 'OMP_CLAUSE_SELF' for
	'GF_OMP_TARGET_KIND_OACC_DATA_KERNELS'.
	gcc/testsuite/
	* c-c++-common/goacc/self-clause-2.c: Verify
	'--param=openacc-kernels=decompose'.
	* gfortran.dg/goacc/kernels-tree.f95: Adjust.
	libgomp/
	* oacc-parallel.c (GOACC_data_start): Handle
	'GOACC_FLAG_LOCAL_DEVICE'.
	(GOACC_parallel_keyed): Simplify accordingly.
	* testsuite/libgomp.oacc-fortran/self-1.f90: Adjust.

4 files changed, 27 insertions, 10 deletions

diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 8576b93..5c6a7f2 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -10334,9 +10334,19 @@ expand_omp_target (struct omp_region *region)
 
   if ((c = omp_find_clause (clauses, OMP_CLAUSE_SELF)) != NULL_TREE)
     {
-      gcc_assert (is_gimple_omp_oacc (entry_stmt) && offloaded);
+      gcc_assert ((is_gimple_omp_oacc (entry_stmt) && offloaded)
+		  || (gimple_omp_target_kind (entry_stmt)
+		      == GF_OMP_TARGET_KIND_OACC_DATA_KERNELS));
 
-      edge e = split_block_after_labels (new_bb);
+      edge e;
+      if (offloaded)
+	e = split_block_after_labels (new_bb);
+      else
+	{
+	  gsi = gsi_last_nondebug_bb (new_bb);
+	  gsi_prev (&gsi);
+	  e = split_block (new_bb, gsi_stmt (gsi));
+	}
       basic_block cond_bb = e->src;
       new_bb = e->dest;
       remove_edge (e);
diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc
index ffc0a8f..dfbb349 100644
--- a/gcc/omp-oacc-kernels-decompose.cc
+++ b/gcc/omp-oacc-kernels-decompose.cc
@@ -1519,17 +1519,18 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt)
 	      break;
 	    }
 	}
-      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF)
+      else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF
+	       || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF)
 	{
-	  /* If there is an 'if' clause, it must be duplicated to the
-	     enclosing data region.  Temporarily remove the if clause's
-	     chain to avoid copying it.  */
+	  /* If there is an 'if' or 'self' clause, it must be duplicated to the
+	     enclosing data region.  Temporarily remove its chain to avoid
+	     copying it.  */
 	  tree saved_chain = OMP_CLAUSE_CHAIN (c);
 	  OMP_CLAUSE_CHAIN (c) = NULL;
-	  tree new_if_clause = unshare_expr (c);
+	  tree new_clause = unshare_expr (c);
 	  OMP_CLAUSE_CHAIN (c) = saved_chain;
-	  OMP_CLAUSE_CHAIN (new_if_clause) = data_clauses;
-	  data_clauses = new_if_clause;
+	  OMP_CLAUSE_CHAIN (new_clause) = data_clauses;
+	  data_clauses = new_clause;
 	}
     }
   /* Restore the original order of the clauses.  */
diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
index 769694b..3ac29a0 100644
--- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c
@@ -1,6 +1,8 @@
 /* See also 'if-clause-2.c'.  */
 
 /* { dg-additional-options "-fdump-tree-gimple" } */
+/* { dg-additional-options "--param=openacc-kernels=decompose" }
+   { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */
 
 void
 f (short c)
@@ -11,6 +13,8 @@ f (short c)
 
 #pragma acc kernels self(c) copy(c)
   /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */
+  /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } }
+     { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single async\(-1\) num_gangs\(1\) map\(force_present:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } */
   ++c;
 
 #pragma acc serial self(c) copy(c)
@@ -29,6 +33,8 @@ g (short d)
 
 #pragma acc kernels self copy(d)
   /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */
+  /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "omp_oacc_kernels_decompose" } }
+     { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single async\(-1\) num_gangs\(1\) map\(force_present:d \[len: [0-9]+\]\) self\(1+\)$} 1 "omp_oacc_kernels_decompose" } } */
   ++d;
 
 #pragma acc serial self copy(d)
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index 1ba04a8..2ee578f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -42,5 +42,5 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } 
 
-! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } }
+! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\) self\(1\)$} 1 "omp_oacc_kernels_decompose" } }
 ! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } }