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
2014-07-17 Cesar Philippidis <ce...@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;
