Patchwork [gomp-4_0-branch] acc nested function support

login
register
mail settings
Submitter Cesar Philippidis
Date July 18, 2014, 9:07 p.m.
Message ID <53C98C74.9030508@codesourcery.com>
Download mbox | patch
Permalink /patch/371721/
State New
Headers show

Comments

Cesar Philippidis - July 18, 2014, 9:07 p.m.
Hi Thomas,

This patch enables acc constructs to be used inside nested functions. I
doubt nested functions will be used much in c, but some of the openacc
fortran tutorials I've seen online make use of internal subroutines in
fortran. Those internal subroutines are implemented as nested functions.

Does this look OK to commit to gomp-4_0-branch?

Thanks,
Cesar
Thomas Schwinge - July 29, 2014, 8:16 a.m.
Hi Cesar!

On Fri, 18 Jul 2014 14:07:00 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch enables acc constructs to be used inside nested functions.

Thanks!

> I
> doubt nested functions will be used much in c, but some of the openacc
> fortran tutorials I've seen online make use of internal subroutines in
> fortran. Those internal subroutines are implemented as nested functions.

> Does this look OK to commit to gomp-4_0-branch?

I think we first need to resolve the following issue:

> --- a/gcc/gimple.h
> +++ b/gcc/gimple.h
> @@ -576,22 +576,6 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
>    tree data_arg;
>  };
>  
> -/* GIMPLE_OACC_KERNELS */
> -struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
> -  gimple_statement_oacc_kernels : public gimple_statement_omp_parallel_layout
> -{
> -    /* No extra fields; adds invariant:
> -         stmt->code == GIMPLE_OACC_KERNELS.  */
> -};
> -
> -/* GIMPLE_OACC_PARALLEL */
> -struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
> -  gimple_statement_oacc_parallel : public gimple_statement_omp_parallel_layout
> -{
> -    /* No extra fields; adds invariant:
> -         stmt->code == GIMPLE_OACC_PARALLEL.  */
> -};
> -
>  /* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */
>  struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
>    gimple_statement_omp_taskreg : public gimple_statement_omp_parallel_layout
> @@ -617,6 +601,22 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
>           stmt->code == GIMPLE_OMP_TARGET.  */
>  };
>  
> +/* GIMPLE_OACC_KERNELS */
> +struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
> +  gimple_statement_oacc_kernels : public gimple_statement_omp_taskreg
> +{
> +    /* No extra fields; adds invariant:
> +         stmt->code == GIMPLE_OACC_KERNELS.  */
> +};
> +
> +/* GIMPLE_OACC_PARALLEL */
> +struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
> +  gimple_statement_oacc_parallel : public gimple_statement_omp_taskreg
> +{
> +    /* No extra fields; adds invariant:
> +         stmt->code == GIMPLE_OACC_PARALLEL.  */
> +};
> +
>  /* GIMPLE_OMP_TASK */
>  
>  struct GTY((tag("GSS_OMP_TASK")))
> @@ -927,7 +927,8 @@ template <>
>  inline bool
>  is_a_helper <gimple_statement_omp_taskreg *>::test (gimple gs)
>  {
> -  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
> +  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK
> +    || gs->code == GIMPLE_OACC_PARALLEL || gs->code == GIMPLE_OACC_KERNELS;
>  }
>  
>  template <>
> @@ -1135,7 +1136,8 @@ template <>
>  inline bool
>  is_a_helper <const gimple_statement_omp_taskreg *>::test (const_gimple gs)
>  {
> -  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
> +  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK
> +    || gs->code == GIMPLE_OACC_PARALLEL || gs->code == GIMPLE_OACC_KERNELS;
>  }

Doing it this way has been "disapproved" before, and I raised the issue
in
<http://news.gmane.org/find-root.php?message_id=%3C87egxcjkmc.fsf%40kepler.schwinge.homeip.net%3E>.
If that makes sense to you, please continue this discussion until I'm
back next week.

> --- a/gcc/tree-nested.c
> +++ b/gcc/tree-nested.c
> @@ -1322,7 +1320,22 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
>  
>      case GIMPLE_OACC_KERNELS:
>      case GIMPLE_OACC_PARALLEL:
> -      gcc_unreachable ();
> +      save_suppress = info->suppress_expansion;
> +      convert_nonlocal_omp_clauses (gimple_omp_taskreg_clauses_ptr (stmt),
> +				    wi);
> +      save_local_var_chain = info->new_local_var_chain;
> +      info->new_local_var_chain = NULL;
> +
> +      walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
> +	         info, gimple_omp_body_ptr (stmt));
> +
> +      if (info->new_local_var_chain)
> +	declare_vars (info->new_local_var_chain,
> +	              gimple_seq_first_stmt (gimple_omp_body (stmt)),
> +		      false);
> +      info->new_local_var_chain = save_local_var_chain;
> +      info->suppress_expansion = save_suppress;
> +      break;
>  
>      case GIMPLE_OMP_PARALLEL:
>      case GIMPLE_OMP_TASK:

Wouldn't we rather group the GIMPLE_OACC_* with GIMPLE_OMP_TARGET
(handling all these together)?  (Either way, please avoid the code
duplication for GIMPLE_OACC_KERNELS/GIMPLE_OACC_PARALLEL even if grouping
with GIMPLE_OMP_PARALLEL/GIMPLE_OMP_TASK; I think the latter's existing
code can basically be re-used as is?


Grüße,
 Thomas

Patch

2014-07-17  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* gcc/gimple.h (gimple_statement_oacc_kernels,
	gimple_statment_oacc_parallel): Derive from
	gimple_statement_omp_taskreg instestead of
	gimple_statement_omp_parallel_layout.
	(is_a_helper <gimple_statement_omp_taskreg *>::test): Permit gimple
	codes GIMPLE_OACC_PARALLEL and GIMPLE_OACC_KERNELS.
	(is_a_helper <const gimple_statement_omp_taskreg *>::test): Likewise.
	*tree-nested.c (walk_gimple_omp_for): Handle openacc kernels and
	parallel constructs.
	(convert_nonlocal_reference_stmt): Likewise.
	(convert_local_reference_stmt): Likewise.
	(convert_tramp_reference_stmt): Likewise.
	(convert_gimple_call): Likewise.

	gcc/testsuite/
	* c-c++-common/goacc/nested-function-1.c: New test.
	* gfortran.dg/goacc/cray-2.f95: New test.
	* gfortran.dg/goacc/loop-4.f95: New test.
	* gfortran.dg/goacc/loop-5.f95: New test.


diff --git a/gcc/gimple.h b/gcc/gimple.h
index 68d1745..d45010f 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -576,22 +576,6 @@  struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   tree data_arg;
 };
 
-/* GIMPLE_OACC_KERNELS */
-struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
-  gimple_statement_oacc_kernels : public gimple_statement_omp_parallel_layout
-{
-    /* No extra fields; adds invariant:
-         stmt->code == GIMPLE_OACC_KERNELS.  */
-};
-
-/* GIMPLE_OACC_PARALLEL */
-struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
-  gimple_statement_oacc_parallel : public gimple_statement_omp_parallel_layout
-{
-    /* No extra fields; adds invariant:
-         stmt->code == GIMPLE_OACC_PARALLEL.  */
-};
-
 /* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */
 struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   gimple_statement_omp_taskreg : public gimple_statement_omp_parallel_layout
@@ -617,6 +601,22 @@  struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
          stmt->code == GIMPLE_OMP_TARGET.  */
 };
 
+/* GIMPLE_OACC_KERNELS */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+  gimple_statement_oacc_kernels : public gimple_statement_omp_taskreg
+{
+    /* No extra fields; adds invariant:
+         stmt->code == GIMPLE_OACC_KERNELS.  */
+};
+
+/* GIMPLE_OACC_PARALLEL */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+  gimple_statement_oacc_parallel : public gimple_statement_omp_taskreg
+{
+    /* No extra fields; adds invariant:
+         stmt->code == GIMPLE_OACC_PARALLEL.  */
+};
+
 /* GIMPLE_OMP_TASK */
 
 struct GTY((tag("GSS_OMP_TASK")))
@@ -927,7 +927,8 @@  template <>
 inline bool
 is_a_helper <gimple_statement_omp_taskreg *>::test (gimple gs)
 {
-  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
+  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK
+    || gs->code == GIMPLE_OACC_PARALLEL || gs->code == GIMPLE_OACC_KERNELS;
 }
 
 template <>
@@ -1135,7 +1136,8 @@  template <>
 inline bool
 is_a_helper <const gimple_statement_omp_taskreg *>::test (const_gimple gs)
 {
-  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK;
+  return gs->code == GIMPLE_OMP_PARALLEL || gs->code == GIMPLE_OMP_TASK
+    || gs->code == GIMPLE_OACC_PARALLEL || gs->code == GIMPLE_OACC_KERNELS;
 }
 
 template <>
diff --git a/gcc/testsuite/c-c++-common/goacc/nested-function-1.c b/gcc/testsuite/c-c++-common/goacc/nested-function-1.c
new file mode 100644
index 0000000..51a0e9f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/nested-function-1.c
@@ -0,0 +1,47 @@ 
+/* { dg-do compile } */
+
+extern void abort (void);
+
+int
+main (void)
+{
+  int j = 0, k = 6, l = 7, m = 8;
+  void simple (void)
+  {
+    int i;
+#pragma acc parallel
+    {
+#pragma acc loop
+      for (i = 0; i < m; i+= k)
+	j = (m + i - j) * l;
+    }
+  }
+  void collapse (void)
+  {
+    int x, y, z;
+#pragma acc parallel
+    {
+#pragma acc loop collapse (3)
+      for (x = 0; x < k; x++)
+	for (y = -5; y < l; y++)
+	  for (z = 0; z < m; z++)
+	    j += x + y + z;
+    }
+  }
+  void reduction (void)
+  {
+    int x, y, z;
+#pragma acc parallel
+    {
+#pragma acc loop collapse (3) reduction (+:j)
+      for (x = 0; x < k; x++)
+	for (y = -5; y < l; y++)
+	  for (z = 0; z < m; z++)
+	    j += x + y + z;
+    }
+  }
+  simple();
+  collapse();
+  reduction();
+  return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/cray-2.f95 b/gcc/testsuite/gfortran.dg/goacc/cray-2.f95
new file mode 100644
index 0000000..70f7cf6
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/cray-2.f95
@@ -0,0 +1,55 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fcray-pointer" }
+
+program test
+  call oacc1
+contains
+  subroutine oacc1
+    implicit none
+    integer :: i
+    real :: pointee
+    pointer (ptr, pointee)
+    !$acc declare device_resident (pointee)
+    !$acc declare device_resident (ptr)
+    !$acc data copy (pointee) ! { dg-error "Cray pointee" }
+    !$acc end data
+    !$acc data deviceptr (pointee) ! { dg-error "Cray pointee" }
+    !$acc end data
+    !$acc parallel private (pointee) ! { dg-error "Cray pointee" }
+    !$acc end parallel
+    !$acc host_data use_device (pointee) ! { dg-error "Cray pointee" }
+    !$acc end host_data
+    !$acc parallel loop reduction(+:pointee) ! { dg-error "Cray pointee" }
+    do i = 1,5
+    enddo
+    !$acc end parallel loop
+    !$acc parallel loop
+    do i = 1,5
+      ! Subarrays are not implemented yet
+      !$acc cache (pointee) ! TODO: This must fail, as in openacc-1_0-branch
+    enddo
+    !$acc end parallel loop
+    !$acc update host (pointee) ! { dg-error "Cray pointee" }
+    !$acc update device (pointee) ! { dg-error "Cray pointee" }
+    !$acc data copy (ptr)
+    !$acc end data
+    !$acc data deviceptr (ptr) ! { dg-error "Cray pointer" }
+    !$acc end data
+    !$acc parallel private (ptr)
+    !$acc end parallel
+    !$acc host_data use_device (ptr) ! { dg-error "Cray pointer" }
+    !$acc end host_data
+    !$acc parallel loop reduction(+:ptr) ! { dg-error "Cray pointer" }
+    do i = 1,5
+    enddo
+    !$acc end parallel loop
+    !$acc parallel loop
+    do i = 1,5
+      !$acc cache (ptr) ! TODO: This must fail, as in openacc-1_0-branch
+    enddo
+    !$acc end parallel loop
+    !$acc update host (ptr)
+    !$acc update device (ptr)
+  end subroutine oacc1
+end program test
+! { dg-prune-output "unimplemented" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-4.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
new file mode 100644
index 0000000..f876106
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-4.f95
@@ -0,0 +1,170 @@ 
+! { dg-do compile }
+! { dg-additional-options "-fmax-errors=100" }
+program test
+  call test1
+contains
+subroutine test1
+  integer :: i, j, k, b(10)
+  integer, dimension (30) :: a
+  double precision :: d
+  real :: r
+  i = 0
+  !$acc loop
+  do 100 ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+    if (i .gt. 0) exit ! { dg-error "EXIT statement" }
+  100 i = i + 1
+  i = 0
+  !$acc loop
+  do ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+      if (i .gt. 0) exit ! { dg-error "EXIT statement" }
+       i = i + 1
+  end do
+  i = 0
+  !$acc loop
+  do 200 while (i .lt. 4) ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+  200 i = i + 1
+  !$acc loop
+  do while (i .lt. 8) ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+       i = i + 1
+  end do
+  !$acc loop
+  do 300 d = 1, 30, 6 ! { dg-error "integer" }
+      i = d
+  300 a(i) = 1
+  !$acc loop
+  do d = 1, 30, 5 ! { dg-error "integer" }
+       i = d
+      a(i) = 2
+  end do
+  !$acc loop
+  do i = 1, 30
+      if (i .eq. 16) exit ! { dg-error "EXIT statement" }
+  end do
+  !$acc loop
+  outer: do i = 1, 30
+      do j = 5, 10
+          if (i .eq. 6 .and. j .eq. 7) exit outer ! { dg-error "EXIT statement" }
+      end do
+  end do outer
+  last: do i = 1, 30
+   end do last
+
+  ! different types of loop are allowed
+  !$acc loop
+  do i = 1,10
+  end do
+  !$acc loop
+  do 400, i = 1,10
+400   a(i) = i
+
+  ! after loop directive must be loop
+  !$acc loop
+  a(1) = 1 ! { dg-error "Expected DO loop" }
+  do i = 1,10
+  enddo
+
+  ! combined directives may be used with/without end
+  !$acc parallel loop
+  do i = 1,10
+  enddo
+  !$acc parallel loop
+  do i = 1,10
+  enddo
+  !$acc end parallel loop
+  !$acc kernels loop
+  do i = 1,10
+  enddo
+  !$acc kernels loop
+  do i = 1,10
+  enddo
+  !$acc end kernels loop
+
+  !$acc kernels loop reduction(max:i)
+  do i = 1,10
+  enddo
+  !$acc kernels
+  !$acc loop reduction(max:i)
+  do i = 1,10
+  enddo
+  !$acc end kernels
+
+  !$acc parallel loop collapse(0) ! { dg-error "constant positive integer" }
+  do i = 1,10
+  enddo
+
+  !$acc parallel loop collapse(-1) ! { dg-error "constant positive integer" }
+  do i = 1,10
+  enddo
+
+  !$acc parallel loop collapse(i) ! { dg-error "Constant expression required" }
+  do i = 1,10
+  enddo
+
+  !$acc parallel loop collapse(4) ! { dg-error "not enough DO loops for collapsed" }
+    do i = 1, 3
+        do j = 4, 6
+          do k = 5, 7
+              a(i+j-k) = i + j + k
+          end do
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 5, 2
+        do j = i + 1, 7, i  ! { dg-error "collapsed loops don.t form rectangular iteration space" }
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+        k = 4
+    end do
+    !$acc parallel loop collapse(3-1)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+        k = 4
+    end do
+    !$acc parallel loop collapse(1+1)
+    do i = 1, 3
+        do j = 4, 6
+        end do
+        k = 4
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do      ! { dg-error "cannot be a DO WHILE or DO without loop control" }
+        end do
+    end do
+    !$acc parallel loop collapse(2)
+    do i = 1, 3
+        do r = 4, 6    ! { dg-error "integer" }
+        end do
+    end do
+
+    ! Both seq and independent are not allowed
+  !$acc loop independent seq ! { dg-error "SEQ conflicts with INDEPENDENT" }
+  do i = 1,10
+  enddo
+
+
+  !$acc cache (a) ! { dg-error "inside of loop" }
+
+  do i = 1,10
+    !$acc cache(a)
+  enddo
+
+  do i = 1,10
+    a(i) = i
+    !$acc cache(a)
+  enddo
+
+end subroutine test1
+end program test
+! { dg-prune-output "Deleted" }
+! { dg-prune-output "ACC cache unimplemented" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-5.f95 b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
new file mode 100644
index 0000000..448d2f5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-5.f95
@@ -0,0 +1,58 @@ 
+! { dg-do compile }
+! { dg-additional-options "-std=f2008" }
+
+program test
+  call test1
+contains
+subroutine test1
+  implicit none
+  integer :: i, j
+
+  ! !$acc end loop not required by spec
+  !$acc loop
+  do i = 1,5
+  enddo
+  !$acc end loop ! { dg-warning "Redundant" }
+
+  !$acc loop
+  do i = 1,5
+  enddo
+  j = 1
+  !$acc end loop ! { dg-error "Unexpected" }
+
+  !$acc parallel
+  !$acc loop
+  do i = 1,5
+  enddo
+  !$acc end parallel
+  !$acc end loop ! { dg-error "Unexpected" }
+
+  ! OpenACC supports Fortran 2008 do concurrent statement
+  !$acc loop
+  do concurrent (i = 1:5)
+  end do
+
+  !$acc loop
+  outer_loop: do i = 1, 5
+    inner_loop: do j = 1,5
+      if (i .eq. j) cycle outer_loop
+      if (i .ne. j) exit outer_loop ! { dg-error "EXIT statement" }
+    end do inner_loop
+  end do outer_loop
+
+  outer_loop1: do i = 1, 5
+    !$acc loop
+    inner_loop1: do j = 1,5
+      if (i .eq. j) cycle outer_loop1 ! { dg-error "CYCLE statement" }
+    end do inner_loop1
+  end do outer_loop1
+
+  !$acc loop collapse(2)
+  outer_loop2: do i = 1, 5
+    inner_loop2: do j = 1,5
+      if (i .eq. j) cycle outer_loop2 ! { dg-error "CYCLE statement" }
+      if (i .ne. j) exit outer_loop2 ! { dg-error "EXIT statement" }
+    end do inner_loop2
+  end do outer_loop2
+end subroutine test1
+end program test
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index 06f3589..3d0af52 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -622,8 +622,6 @@  walk_gimple_omp_for (gimple for_stmt,
     		     walk_stmt_fn callback_stmt, walk_tree_fn callback_op,
     		     struct nesting_info *info)
 {
-  gcc_assert (!is_gimple_omp_oacc_specifically (for_stmt));
-
   struct walk_stmt_info wi;
   gimple_seq seq;
   tree t;
@@ -1322,7 +1320,22 @@  convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 
     case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
+      save_suppress = info->suppress_expansion;
+      convert_nonlocal_omp_clauses (gimple_omp_taskreg_clauses_ptr (stmt),
+				    wi);
+      save_local_var_chain = info->new_local_var_chain;
+      info->new_local_var_chain = NULL;
+
+      walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
+	         info, gimple_omp_body_ptr (stmt));
+
+      if (info->new_local_var_chain)
+	declare_vars (info->new_local_var_chain,
+	              gimple_seq_first_stmt (gimple_omp_body (stmt)),
+		      false);
+      info->new_local_var_chain = save_local_var_chain;
+      info->suppress_expansion = save_suppress;
+      break;
 
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
@@ -1354,7 +1367,6 @@  convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_FOR:
-      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       save_suppress = info->suppress_expansion;
       convert_nonlocal_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi);
       walk_gimple_omp_for (stmt, convert_nonlocal_reference_stmt,
@@ -1895,8 +1907,6 @@  convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
     {
     case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
-
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_suppress = info->suppress_expansion;
@@ -1926,7 +1936,6 @@  convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_FOR:
-      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       save_suppress = info->suppress_expansion;
       convert_local_omp_clauses (gimple_omp_for_clauses_ptr (stmt), wi);
       walk_gimple_omp_for (stmt, convert_local_reference_stmt,
@@ -2286,18 +2295,15 @@  convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 	break;
       }
 
-    case GIMPLE_OACC_KERNELS:
-    case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
-
     case GIMPLE_OMP_TARGET:
-      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
 	{
 	  *handled_ops_p = false;
 	  return NULL_TREE;
 	}
       /* FALLTHRU */
+    case GIMPLE_OACC_KERNELS:
+    case GIMPLE_OACC_PARALLEL:
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       {
@@ -2359,8 +2365,6 @@  convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
 
     case GIMPLE_OACC_KERNELS:
     case GIMPLE_OACC_PARALLEL:
-      gcc_unreachable ();
-
     case GIMPLE_OMP_PARALLEL:
     case GIMPLE_OMP_TASK:
       save_static_chain_added = info->static_chain_added;
@@ -2431,7 +2435,6 @@  convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_FOR:
-      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       walk_body (convert_gimple_call, NULL, info,
 	  	 gimple_omp_for_pre_body_ptr (stmt));
       /* FALLTHRU */
@@ -2443,7 +2446,6 @@  convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
     case GIMPLE_OMP_TASKGROUP:
     case GIMPLE_OMP_ORDERED:
     case GIMPLE_OMP_CRITICAL:
-      gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
       walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
       break;