This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [patch,gomp-4_0-branch] acc nested function support
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: Thomas Schwinge <thomas at codesourcery dot com>, "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, <dmalcolm at redhat dot com>
- Date: Tue, 4 Nov 2014 16:45:31 -0800
- Subject: Re: [patch,gomp-4_0-branch] acc nested function support
- Authentication-results: sourceware.org; auth=none
- References: <53C98C74 dot 9030508 at codesourcery dot com>
Here's an updated version of my nested function patch.
David, I tweaked the gimple class hierarchy a little bit. Here's what
the updated class diagram looks like:
+ gimple_statement_omp
| | layout: GSS_OMP. Used for code GIMPLE_OMP_SECTION
| |
| + gimple_statement_omp_parallel_layout
| | | layout: GSS_OMP_PARALLEL_LAYOUT
| | |
| | + gimple_statement_omp_targetreg
| | |
| | + gimple_statement_oacc_kernels
| | | code: GIMPLE_OACC_KERNELS
| | |
| | + gimple_statement_oacc_parallel
| | | code: GIMPLE_OACC_PARALLEL
| | |
| | + gimple_statement_omp_target
| | code: GIMPLE_OMP_TARGET
Basically, I've introduced gimple_statement_omp_targetreg and made
GIMPLE_OACC_{PARALLEL,KERNELS} and GIMPLE_OMP_TARGET inherit it. This
seems to work out pretty good. It cleans up both
{lower,expand}_oacc_offload in omp-low.c and allows OpenACC kernel and
parallel regions to be treated as OpenMP target regions in
tree-nested.c. Are these changes to gimple.h OK?
Thomas, assuming these gimple changes are OK, should I commit this
change to gomp-4_0-branch, or do you want to include this patch with
your middle end trunk submission?
Thanks,
Cesar
2014-11-04 Cesar Philippidis <cesar@codesourcery.com>
gcc/
* doc/gimple.texi (gimple class hierarchy): Add
gimple_statement_omp_targetreg, gimple_statement_oacc_kernels and
gimple_statement_oacc_parallel. Make gimple_statement_omp_target
inherit gimple_statement_omp_targetreg.
* gcc/gimple.h (gimple_statement_omp_targetreg): Declare.
(gimple_statement_oacc_kernels): Derive from
gimple_statement_omp_targetreg.
(gimple_statement_oacc_parallel): Likewise.
(gimple_statement_oacc_target): Likewise.
(is_a_helper <gimple_statement_omp_targetreg *>): Define.
(is_a_helper <const gimple_statement_omp_targetreg *>): Define.
(gimple_omp_subcode): Use GIMPLE_OACC_KERNELS as the starting point
for OpenACC/OpenMP subcodes.
(gimple_omp_targetreg_clauses): Declare.
(gimple_omp_targetreg_clauses_ptr): Declare.
(gimple_omp_targetreg_set_clauses): Declare.
(gimple_omp_targetreg_child_fn): Declare.
(gimple_omp_targetreg_child_fn_ptr): Declare.
(gimple_omp_targetreg_set_child_fn): Declare.
(gimple_omp_targetreg_data_arg): Declare.
(gimple_omp_targetreg_data_arg_ptr): Declare.
(gimple_omp_targetreg_set_data_arg): Declare.
(gimple_omp_targetreg_kind): Declare.
(gimple_omp_targetreg_set_kind): Declare.
* gcc/omp-low.c (expand_oacc_offload): Use
gimple_omp_targetreg_child_fn and gimple_omp_target_reg_data_arg
instead of the specific functions for OpenACC kernels and parallel
regions.
(lower_oacc_offload): Use gimple_omp_targetreg_clauses and
gimple_omp_targetreg_set_data_arg for similar reasons.
* tree-nested.c (walk_gimple_omp_for): Remove OpenACC assert.
(convert_nonlocal_reference_stmt): Handle GIMPLE_OACC_KERNELS
and GIMPLE_OACC_PARALLEL.
(convert_local_reference_stmt): Remove OpenACC asserts.
(convert_tramp_reference_stmt): Handle GIMPLE_OACC_KERNELS and
GIMPLE_OACC_PARALLEL.
(convert_gimple_call): Remove OpenACC asserts.
gcc/testsuite/
* gcc.dg/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.
libgomp/
* testsuite/libgomp.oacc-c/sub-collapse-1.c: New test.
* testsuite/libgomp.oacc-c/sub-collapse-2.c: New test.
* testsuite/libgomp.oacc-fortran/sub-collapse-1.f90: New test.
* testsuite/libgomp.oacc-fortran/sub-collapse-2.f90: New test.
* testsuite/libgomp.oacc-fortran/sub-collapse-3.f90: New test.
diff --git a/gcc/doc/gimple.texi b/gcc/doc/gimple.texi
index 4c59748..860cb2c 100644
--- a/gcc/doc/gimple.texi
+++ b/gcc/doc/gimple.texi
@@ -354,8 +354,16 @@ kinds, along with their relationships to @code{GSS_} values (layouts) and
| | | + gimple_statement_omp_task
| | | code: GIMPLE_OMP_TASK
| | |
- | | + gimple_statement_omp_target
- | | code: GIMPLE_OMP_TARGET
+ | | + gimple_statement_omp_targetreg
+ | | |
+ | | + gimple_statement_oacc_kernels
+ | | | code: GIMPLE_OACC_KERNELS
+ | | |
+ | | + gimple_statement_oacc_parallel
+ | | | code: GIMPLE_OACC_PARALLEL
+ | | |
+ | | + gimple_statement_omp_target
+ | | code: GIMPLE_OMP_TARGET
| |
| + gimple_statement_omp_sections
| | layout: GSS_OMP_SECTIONS, code: GIMPLE_OMP_SECTIONS
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 7bc673a..76abfb7 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -579,22 +579,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
@@ -612,12 +596,14 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
stmt->code == GIMPLE_OMP_PARALLEL. */
};
-/* GIMPLE_OMP_TARGET */
+/* GIMPLE_OMP_TARGET or GIMPLE_OACC_PARALLEL or GIMPLE_ACC_KERNELS */
struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
- gimple_statement_omp_target : public gimple_statement_omp_parallel_layout
+ gimple_statement_omp_targetreg : public gimple_statement_omp_parallel_layout
{
/* No extra fields; adds invariant:
- stmt->code == GIMPLE_OMP_TARGET. */
+ stmt->code == GIMPLE_OMP_TARGET
+ || stmt->code == GIMPLE_OACC_PARALLEL
+ || stmt->code == GIMPLE_OACC_KERNELS. */
};
/* GIMPLE_OMP_TASK */
@@ -637,6 +623,29 @@ struct GTY((tag("GSS_OMP_TASK")))
tree arg_align;
};
+/* GIMPLE_OACC_KERNELS */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+ gimple_statement_oacc_kernels : public gimple_statement_omp_targetreg
+{
+ /* 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_targetreg
+{
+ /* No extra fields; adds invariant:
+ stmt->code == GIMPLE_OACC_PARALLEL. */
+};
+
+/* GIMPLE_OMP_TARGET */
+struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
+ gimple_statement_omp_target : public gimple_statement_omp_targetreg
+{
+ /* No extra fields; adds invariant:
+ stmt->code == GIMPLE_OMP_TARGET. */
+};
/* GIMPLE_OMP_SECTION */
/* Uses struct gimple_statement_omp. */
@@ -944,6 +953,15 @@ is_a_helper <gimple_statement_omp_parallel *>::test (gimple gs)
template <>
template <>
inline bool
+is_a_helper <gimple_statement_omp_targetreg *>::test (gimple gs)
+{
+ return gs->code == GIMPLE_OMP_TARGET || gs->code == GIMPLE_OACC_PARALLEL
+ || gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
is_a_helper <gimple_statement_omp_target *>::test (gimple gs)
{
return gs->code == GIMPLE_OMP_TARGET;
@@ -1152,6 +1170,15 @@ is_a_helper <const gimple_statement_omp_parallel *>::test (const_gimple gs)
template <>
template <>
inline bool
+is_a_helper <const gimple_statement_omp_targetreg *>::test (const_gimple gs)
+{
+ return gs->code == GIMPLE_OMP_TARGET || gs->code == GIMPLE_OACC_PARALLEL
+ || gs->code == GIMPLE_OACC_KERNELS;
+}
+
+template <>
+template <>
+inline bool
is_a_helper <const gimple_statement_omp_target *>::test (const_gimple gs)
{
return gs->code == GIMPLE_OMP_TARGET;
@@ -1933,7 +1960,7 @@ gimple_references_memory_p (gimple stmt)
static inline unsigned
gimple_omp_subcode (const_gimple s)
{
- gcc_gimple_checking_assert (gimple_code (s) >= GIMPLE_OMP_ATOMIC_LOAD
+ gcc_gimple_checking_assert (gimple_code (s) >= GIMPLE_OACC_KERNELS
&& gimple_code (s) <= GIMPLE_OMP_TEAMS);
return s->subcode;
}
@@ -5318,6 +5345,127 @@ gimple_omp_single_set_clauses (gimple gs, tree clauses)
/* Return the clauses associated with OMP_TARGET GS. */
static inline tree
+gimple_omp_targetreg_clauses (const_gimple gs)
+{
+ const gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <const gimple_statement_omp_targetreg *> (gs);
+ return omp_targetreg_stmt->clauses;
+}
+
+
+/* Return a pointer to the clauses associated with OMP_TARGET GS. */
+
+static inline tree *
+gimple_omp_targetreg_clauses_ptr (gimple gs)
+{
+ gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <gimple_statement_omp_targetreg *> (gs);
+ return &omp_targetreg_stmt->clauses;
+}
+
+
+/* Set CLAUSES to be the clauses associated with OMP_TARGET GS. */
+
+static inline void
+gimple_omp_targetreg_set_clauses (gimple gs, tree clauses)
+{
+ gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <gimple_statement_omp_targetreg *> (gs);
+ omp_targetreg_stmt->clauses = clauses;
+}
+
+
+/* Return the child function used to hold the body of OMP_TARGET GS. */
+
+static inline tree
+gimple_omp_targetreg_child_fn (const_gimple gs)
+{
+ const gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <const gimple_statement_omp_targetreg *> (gs);
+ return omp_targetreg_stmt->child_fn;
+}
+
+/* Return a pointer to the child function used to hold the body of
+ OMP_TARGET GS. */
+
+static inline tree *
+gimple_omp_targetreg_child_fn_ptr (gimple gs)
+{
+ gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <gimple_statement_omp_targetreg *> (gs);
+ return &omp_targetreg_stmt->child_fn;
+}
+
+
+/* Set CHILD_FN to be the child function for OMP_TARGET GS. */
+
+static inline void
+gimple_omp_targetreg_set_child_fn (gimple gs, tree child_fn)
+{
+ gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <gimple_statement_omp_targetreg *> (gs);
+ omp_targetreg_stmt->child_fn = child_fn;
+}
+
+
+/* Return the artificial argument used to send variables and values
+ from the parent to the children threads in OMP_TARGET GS. */
+
+static inline tree
+gimple_omp_targetreg_data_arg (const_gimple gs)
+{
+ const gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <const gimple_statement_omp_targetreg *> (gs);
+ return omp_targetreg_stmt->data_arg;
+}
+
+
+/* Return a pointer to the data argument for OMP_TARGET GS. */
+
+static inline tree *
+gimple_omp_targetreg_data_arg_ptr (gimple gs)
+{
+ gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <gimple_statement_omp_targetreg *> (gs);
+ return &omp_targetreg_stmt->data_arg;
+}
+
+
+/* Set DATA_ARG to be the data argument for OMP_TARGET GS. */
+
+static inline void
+gimple_omp_targetreg_set_data_arg (gimple gs, tree data_arg)
+{
+ gimple_statement_omp_targetreg *omp_targetreg_stmt =
+ as_a <gimple_statement_omp_targetreg *> (gs);
+ omp_targetreg_stmt->data_arg = data_arg;
+}
+
+
+/* Return the kind of OMP targetreg statemement. */
+
+static inline int
+gimple_omp_targetreg_kind (const_gimple g)
+{
+ //GIMPLE_CHECK (g, GIMPLE_OMP_TARGET);
+ return (gimple_omp_subcode (g) & GF_OMP_TARGET_KIND_MASK);
+}
+
+
+/* Set the OMP targetreg kind. */
+
+static inline void
+gimple_omp_targetreg_set_kind (gimple g, int kind)
+{
+ //GIMPLE_CHECK (g, GIMPLE_OMP_TARGET);
+ g->subcode = (g->subcode & ~GF_OMP_TARGET_KIND_MASK)
+ | (kind & GF_OMP_TARGET_KIND_MASK);
+}
+
+
+/* Return the clauses associated with OMP_TARGET GS. */
+
+static inline tree
gimple_omp_target_clauses (const_gimple gs)
{
const gimple_statement_omp_target *omp_target_stmt =
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index d735e86..5e304fe 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -5407,24 +5407,9 @@ expand_oacc_offload (struct omp_region *region)
gimple_stmt_iterator gsi;
gimple entry_stmt, stmt;
edge e;
- tree (*gimple_omp_child_fn) (const_gimple);
- tree (*gimple_omp_data_arg) (const_gimple);
- switch (region->type)
- {
- case GIMPLE_OACC_KERNELS:
- gimple_omp_child_fn = gimple_oacc_kernels_child_fn;
- gimple_omp_data_arg = gimple_oacc_kernels_data_arg;
- break;
- case GIMPLE_OACC_PARALLEL:
- gimple_omp_child_fn = gimple_oacc_parallel_child_fn;
- gimple_omp_data_arg = gimple_oacc_parallel_data_arg;
- break;
- default:
- gcc_unreachable ();
- }
entry_stmt = last_stmt (region->entry);
- child_fn = gimple_omp_child_fn (entry_stmt);
+ child_fn = gimple_omp_targetreg_child_fn (entry_stmt);
child_cfun = DECL_STRUCT_FUNCTION (child_fn);
/* Supported by expand_omp_taskreg, but not here. */
@@ -5452,13 +5437,13 @@ expand_oacc_offload (struct omp_region *region)
a function call that has been inlined, the original PARM_DECL
.OMP_DATA_I may have been converted into a different local
variable. In which case, we need to keep the assignment. */
- if (gimple_omp_data_arg (entry_stmt))
+ if (gimple_omp_targetreg_data_arg (entry_stmt))
{
basic_block entry_succ_bb = single_succ (entry_bb);
gimple_stmt_iterator gsi;
tree arg;
gimple parcopy_stmt = NULL;
- tree sender = TREE_VEC_ELT (gimple_omp_data_arg (entry_stmt), 0);
+ tree sender = TREE_VEC_ELT (gimple_omp_targetreg_data_arg (entry_stmt), 0);
for (gsi = gsi_start_bb (entry_succ_bb); ; gsi_next (&gsi))
{
@@ -5725,7 +5710,7 @@ expand_oacc_offload (struct omp_region *region)
}
gsi = gsi_last_bb (new_bb);
- t = gimple_omp_data_arg (entry_stmt);
+ t = gimple_omp_targetreg_data_arg (entry_stmt);
if (t == NULL)
{
t1 = size_zero_node;
@@ -10319,23 +10304,8 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq par_body, olist, ilist, orlist, irlist, new_body;
location_t loc = gimple_location (stmt);
unsigned int map_cnt = 0;
- tree (*gimple_omp_clauses) (const_gimple);
- void (*gimple_omp_set_data_arg) (gimple, tree);
- switch (gimple_code (stmt))
- {
- case GIMPLE_OACC_KERNELS:
- gimple_omp_clauses = gimple_oacc_kernels_clauses;
- gimple_omp_set_data_arg = gimple_oacc_kernels_set_data_arg;
- break;
- case GIMPLE_OACC_PARALLEL:
- gimple_omp_clauses = gimple_oacc_parallel_clauses;
- gimple_omp_set_data_arg = gimple_oacc_parallel_set_data_arg;
- break;
- default:
- gcc_unreachable ();
- }
- clauses = gimple_omp_clauses (stmt);
+ clauses = gimple_omp_targetreg_clauses (stmt);
par_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
par_body = gimple_bind_body (par_bind);
child_fn = ctx->cb.dst_fn;
@@ -10428,7 +10398,7 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 2)) = 1;
- gimple_omp_set_data_arg (stmt, t);
+ gimple_omp_targetreg_set_data_arg (stmt, t);
vec<constructor_elt, va_gc> *vsize;
vec<constructor_elt, va_gc> *vkind;
diff --git a/gcc/testsuite/gcc.dg/goacc/nested-function-1.c b/gcc/testsuite/gcc.dg/goacc/nested-function-1.c
new file mode 100644
index 0000000..51a0e9f
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/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 b5d6543..e8ece9c 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -627,8 +627,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;
@@ -1325,10 +1323,6 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
break;
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- gcc_unreachable ();
-
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_suppress = info->suppress_expansion;
@@ -1359,7 +1353,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,
@@ -1385,12 +1378,14 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
info->suppress_expansion = save_suppress;
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
- gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
- if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
+ if (gimple_omp_targetreg_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
save_suppress = info->suppress_expansion;
- convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt),
+ convert_nonlocal_omp_clauses (gimple_omp_targetreg_clauses_ptr
+ (stmt),
wi);
info->suppress_expansion = save_suppress;
walk_body (convert_nonlocal_reference_stmt,
@@ -1399,7 +1394,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
}
save_suppress = info->suppress_expansion;
- if (convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt),
+ if (convert_nonlocal_omp_clauses (gimple_omp_targetreg_clauses_ptr (stmt),
wi))
{
tree c, decl;
@@ -1408,8 +1403,8 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
OMP_CLAUSE_DECL (c) = decl;
OMP_CLAUSE_MAP_KIND (c) = OMP_CLAUSE_MAP_TO;
OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
- OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt);
- gimple_omp_target_set_clauses (stmt, c);
+ OMP_CLAUSE_CHAIN (c) = gimple_omp_targetreg_clauses (stmt);
+ gimple_omp_targetreg_set_clauses (stmt, c);
}
save_local_var_chain = info->new_local_var_chain;
@@ -1898,10 +1893,6 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
switch (gimple_code (stmt))
{
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- gcc_unreachable ();
-
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
save_suppress = info->suppress_expansion;
@@ -1931,7 +1922,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,
@@ -1957,19 +1947,20 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
info->suppress_expansion = save_suppress;
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
- gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
- if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
+ if (gimple_omp_targetreg_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
save_suppress = info->suppress_expansion;
- convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
+ convert_local_omp_clauses (gimple_omp_targetreg_clauses_ptr (stmt), wi);
info->suppress_expansion = save_suppress;
walk_body (convert_local_reference_stmt, convert_local_reference_op,
info, gimple_omp_body_ptr (stmt));
break;
}
save_suppress = info->suppress_expansion;
- if (convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi))
+ if (convert_local_omp_clauses (gimple_omp_targetreg_clauses_ptr (stmt), wi))
{
tree c;
(void) get_frame_type (info);
@@ -1977,8 +1968,8 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
OMP_CLAUSE_DECL (c) = info->frame_decl;
OMP_CLAUSE_MAP_KIND (c) = OMP_CLAUSE_MAP_TOFROM;
OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (info->frame_decl);
- OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt);
- gimple_omp_target_set_clauses (stmt, c);
+ OMP_CLAUSE_CHAIN (c) = gimple_omp_targetreg_clauses (stmt);
+ gimple_omp_targetreg_set_clauses (stmt, c);
}
save_local_var_chain = info->new_local_var_chain;
@@ -2291,11 +2282,8 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
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)
+ if (gimple_omp_targetreg_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
*handled_ops_p = false;
return NULL_TREE;
@@ -2360,10 +2348,6 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
}
break;
- 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;
@@ -2396,9 +2380,10 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
info->static_chain_added |= save_static_chain_added;
break;
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
case GIMPLE_OMP_TARGET:
- gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
- if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
+ if (gimple_omp_targetreg_kind (stmt) != GF_OMP_TARGET_KIND_REGION)
{
walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
break;
@@ -2413,7 +2398,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
continue;
decl = i ? get_chain_decl (info) : info->frame_decl;
/* Don't add CHAIN.* or FRAME.* twice. */
- for (c = gimple_omp_target_clauses (stmt);
+ for (c = gimple_omp_targetreg_clauses (stmt);
c;
c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -2426,15 +2411,14 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
OMP_CLAUSE_MAP_KIND (c)
= i ? OMP_CLAUSE_MAP_TO : OMP_CLAUSE_MAP_TOFROM;
OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
- OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt);
- gimple_omp_target_set_clauses (stmt, c);
+ OMP_CLAUSE_CHAIN (c) = gimple_omp_targetreg_clauses (stmt);
+ gimple_omp_targetreg_set_clauses (stmt, c);
}
}
info->static_chain_added |= save_static_chain_added;
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 */
@@ -2446,7 +2430,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;
diff --git a/libgomp/testsuite/libgomp.oacc-c/sub-collapse-1.c b/libgomp/testsuite/libgomp.oacc-c/sub-collapse-1.c
new file mode 100644
index 0000000..f28348a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c/sub-collapse-1.c
@@ -0,0 +1,59 @@
+/* { dg-do run } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+int
+main (void)
+{
+ void test1 ()
+ {
+ int i, j, k;
+ int a[4][7][8];
+ int l = 0;
+
+ memset (a, 0, sizeof (a));
+
+#pragma acc parallel
+#pragma acc loop collapse(4 - 1)
+ for (i = 1; i <= 3; i++)
+ for (j = 4; j <= 6; j++)
+ for (k = 5; k <= 7; k++)
+ a[i][j][k] = i + j + k;
+#pragma acc end parallel
+
+ for (i = 1; i <= 3; i++)
+ for (j = 4; j <= 6; j++)
+ for (k = 5; k <= 7; k++)
+ if (a[i][j][k] != i + j + k)
+ abort();
+ }
+
+ void test2 ()
+ {
+ int i, j, k;
+ int a[4][4][4];
+
+ memset (a, 0, sizeof (a));
+
+#pragma acc parallel
+#pragma acc loop collapse(3)
+ for (i = 1; i <= 3; i++)
+ for (j = 1; j <= 3; j++)
+ for (k = 1; k <= 3; k++)
+ a[i][j][k] = 1;
+#pragma acc end parallel
+
+ for (i = 1; i <= 3; i++)
+ for (j = 1; j <= 3; j++)
+ for (k = 1; k <= 3; k++)
+ if (a[i][j][k] != 1)
+ abort ();
+ }
+
+ test1 ();
+ test2 ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c/sub-collapse-2.c b/libgomp/testsuite/libgomp.oacc-c/sub-collapse-2.c
new file mode 100644
index 0000000..00f8d4e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c/sub-collapse-2.c
@@ -0,0 +1,163 @@
+/* { dg-do run } */
+
+#include <string.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+int
+main (void)
+{
+ int p1 = 2, p2 = 6, p3 = 0, p4 = 4, p5 = 13, p6 = 18, p7 = 1, p8 = 1, p9 = 1;
+
+ void test1 ()
+ {
+ int i, j, k;
+ int a[4][4][4];
+
+ memset (a, '\0', sizeof (a));
+
+#pragma acc parallel
+#pragma acc loop collapse(3)
+ for (i = 1; i <= 3; i++)
+ for (j = 1; j <= 3; j++)
+ for (k = 2; k <= 3; k++)
+ a[i][j][k] = 1;
+#pragma acc end parallel
+
+ for (i = 1; i <= 3; i++)
+ for (j = 1; j <= 3; j++)
+ for (k = 2; k <= 3; k++)
+ if (a[i][j][k] != 1)
+ abort();
+ }
+
+ void test2 (int v1, int v2, int v3, int v4, int v5, int v6)
+ {
+ int i, j, k, l = 0, r = 0;
+ int a[7][5][19];
+ int b[7][5][19];
+
+ memset (a, '\0', sizeof (a));
+ memset (b, '\0', sizeof (b));
+
+#pragma acc parallel
+#pragma acc loop collapse(3) reduction (||:l)
+ for (i = v1; i <= v2; i++)
+ for (j = v3; j <= v4; j++)
+ for (k = v5; k <= v6; k++)
+ {
+ l = l || i < 2 || i > 6 || j < 0 || j > 4 || k < 13 || k > 18;
+ if (!l)
+ a[i][j][k] += 1;
+ }
+#pragma acc end parallel
+
+ for (i = v1; i <= v2; i++)
+ for (j = v3; j <= v4; j++)
+ for (k = v5; k <= v6; k++)
+ {
+ r = r || i < 2 || i > 6 || j < 0 || j > 4 || k < 13 || k > 18;
+ if (!r)
+ b[i][j][k] += 1;
+ }
+
+ if (l != r)
+ abort ();
+
+ for (i = v1; i <= v2; i++)
+ for (j = v3; j <= v4; j++)
+ for (k = v5; k <= v6; k++)
+ if (b[i][j][k] != a[i][j][k])
+ abort ();
+ }
+
+ void test3 (int v1, int v2, int v3, int v4, int v5, int v6, int v7, int v8,
+ int v9)
+ {
+ int i, j, k, l = 0, r = 0;
+ int a[7][5][19];
+ int b[7][5][19];
+
+ memset (a, '\0', sizeof (a));
+ memset (b, '\0', sizeof (b));
+
+#pragma acc parallel
+#pragma acc loop collapse(3) reduction (||:l)
+ for (i = v1; i <= v2; i += v7)
+ for (j = v3; j <= v4; j += v8)
+ for (k = v5; k <= v6; k += v9)
+ {
+ l = l || i < 2 || i > 6 || j < 0 || j > 4 || k < 13 || k > 18;
+ if (!l)
+ a[i][j][k] += 1;
+ }
+#pragma acc end parallel
+
+ for (i = v1; i <= v2; i += v7)
+ for (j = v3; j <= v4; j += v8)
+ for (k = v5; k <= v6; k += v9)
+ {
+ r = r || i < 2 || i > 6 || j < 0 || j > 4 || k < 13 || k > 18;
+ if (!r)
+ b[i][j][k] += 1;
+ }
+
+ if (l != r)
+ abort ();
+
+ for (i = v1; i <= v2; i++)
+ for (j = v3; j <= v4; j++)
+ for (k = v5; k <= v6; k++)
+ if (b[i][j][k] != a[i][j][k])
+ abort ();
+ }
+
+ void test4 ()
+ {
+ int i, j, k, l = 0, r = 0;
+ int a[7][5][19];
+ int b[7][5][19];
+ int v1 = p1, v2 = p2, v3 = p3, v4 = p4, v5 = p5, v6 = p6, v7 = p7, v8 = p8,
+ v9 = p9;
+
+ memset (a, '\0', sizeof (a));
+ memset (b, '\0', sizeof (b));
+
+#pragma acc parallel
+#pragma acc loop collapse(3) reduction (||:l)
+ for (i = v1; i <= v2; i += v7)
+ for (j = v3; j <= v4; j += v8)
+ for (k = v5; k <= v6; k += v9)
+ {
+ l = l || i < 2 || i > 6 || j < 0 || j > 4 || k < 13 || k > 18;
+ if (!l)
+ a[i][j][k] += 1;
+ }
+#pragma acc end parallel
+
+ for (i = v1; i <= v2; i += v7)
+ for (j = v3; j <= v4; j += v8)
+ for (k = v5; k <= v6; k += v9)
+ {
+ r = r || i < 2 || i > 6 || j < 0 || j > 4 || k < 13 || k > 18;
+ if (!r)
+ b[i][j][k] += 1;
+ }
+
+ if (l != r)
+ abort ();
+
+ for (i = v1; i <= v2; i++)
+ for (j = v3; j <= v4; j++)
+ for (k = v5; k <= v6; k++)
+ if (b[i][j][k] != a[i][j][k])
+ abort ();
+ }
+
+ test1 ();
+ test2 (p1, p2, p3, p4, p5, p6);
+ test3 (p1, p2, p3, p4, p5, p6, p7, p8, p9);
+ test4 ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-1.f90
new file mode 100644
index 0000000..169cd12
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-1.f90
@@ -0,0 +1,56 @@
+! { dg-do run }
+
+program collapse2
+ call test1
+ call test2
+contains
+ subroutine test1
+ integer :: i, j, k, a(1:3, 4:6, 5:7)
+ logical :: l
+ l = .false.
+ a(:, :, :) = 0
+ !$acc parallel
+ !$acc loop collapse(4 - 1)
+ do 164 i = 1, 3
+ do 164 j = 4, 6
+ do 164 k = 5, 7
+ a(i, j, k) = i + j + k
+164 end do
+ !$acc loop collapse(2) reduction(.or.:l)
+firstdo: do i = 1, 3
+ do j = 4, 6
+ do k = 5, 7
+ if (a(i, j, k) .ne. (i + j + k)) l = .true.
+ end do
+ end do
+ end do firstdo
+ !$acc end parallel
+ if (l) call abort
+ end subroutine test1
+
+ subroutine test2
+ integer :: a(3,3,3), k, kk, kkk, l, ll, lll
+ !$acc parallel
+ !$acc loop collapse(3)
+ do 115 k=1,3
+ dokk: do kk=1,3
+ do kkk=1,3
+ a(k,kk,kkk) = 1
+ enddo
+ enddo dokk
+115 continue
+ if (any(a(1:3,1:3,1:3).ne.1)) call abort
+
+ !$acc loop collapse(3)
+ dol: do 120 l=1,3
+ doll: do ll=1,3
+ do lll=1,3
+ a(l,ll,lll) = 2
+ enddo
+ enddo doll
+120 end do dol
+ if (any(a(1:3,1:3,1:3).ne.2)) call abort
+ !$acc end parallel
+ end subroutine test2
+
+end program collapse2
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-2.f90
new file mode 100644
index 0000000..a86e522
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-2.f90
@@ -0,0 +1,171 @@
+! { dg-do run }
+
+program collapse3
+ integer :: p1, p2, p3, p4, p5, p6, p7, p8, p9
+ p1 = 2
+ p2 = 6
+ p3 = -2
+ p4 = 4
+ p5 = 13
+ p6 = 18
+ p7 = 1
+ p8 = 1
+ p9 = 1
+ call test1
+ call test2 (p1, p2, p3, p4, p5, p6)
+ call test3 (p1, p2, p3, p4, p5, p6, p7, p8, p9)
+ call test4
+contains
+ subroutine test1
+ integer :: a(3,3,3), k, kk, kkk, l, ll, lll
+ !$acc parallel
+ !$acc loop collapse(3)
+ do 115 k=1,3
+dokk: do kk=1,3
+ do kkk=1,3
+ a(k,kk,kkk) = 1
+ enddo
+ enddo dokk
+115 continue
+ !$acc end parallel
+ if (any(a(1:3,1:3,1:3).ne.1)) call abort
+ !$acc parallel
+ !$acc loop collapse(3)
+dol: do 120 l=1,3
+doll: do ll=1,3
+ do lll=1,3
+ a(l,ll,lll) = 2
+ enddo
+ enddo doll
+120 end do dol
+ !$acc end parallel
+ if (any(a(1:3,1:3,1:3).ne.2)) call abort
+ end subroutine test1
+
+ subroutine test2(v1, v2, v3, v4, v5, v6)
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ !$acc parallel
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.l) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test2
+
+ subroutine test3(v1, v2, v3, v4, v5, v6, v7, v8, v9)
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6, v7, v8, v9
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ !$acc parallel
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.l) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test3
+
+ subroutine test4
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6, v7, v8, v9
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ v1 = p1
+ v2 = p2
+ v3 = p3
+ v4 = p4
+ v5 = p5
+ v6 = p6
+ v7 = p7
+ v8 = p8
+ v9 = p9
+ !$acc parallel
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.r) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test4
+
+end program collapse3
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-3.f90
new file mode 100644
index 0000000..f91f0be
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/sub-collapse-3.f90
@@ -0,0 +1,242 @@
+! { dg-do run }
+
+program sub_collapse_3
+ call test1
+ call test2 (2, 6, -2, 4, 13, 18)
+ call test3 (2, 6, -2, 4, 13, 18, 1, 1, 1)
+ call test4
+ call test5 (2, 6, -2, 4, 13, 18)
+ call test6 (2, 6, -2, 4, 13, 18, 1, 1, 1)
+contains
+ subroutine test1
+ integer :: a(3,3,3), k, kk, kkk, l, ll, lll
+ !$acc parallel
+ !$acc loop collapse(3)
+ do 115 k=1,3
+dokk: do kk=1,3
+ do kkk=1,3
+ a(k,kk,kkk) = 1
+ enddo
+ enddo dokk
+115 continue
+ !$acc end parallel
+ if (any(a(1:3,1:3,1:3).ne.1)) call abort
+ !$acc parallel
+ !$acc loop collapse(3)
+dol: do 120 l=1,3
+doll: do ll=1,3
+ do lll=1,3
+ a(l,ll,lll) = 2
+ enddo
+ enddo doll
+120 end do dol
+ !$acc end parallel
+ if (any(a(1:3,1:3,1:3).ne.2)) call abort
+ end subroutine test1
+
+ subroutine test2(v1, v2, v3, v4, v5, v6)
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ !$acc parallel copyin (v1, v2, v3, v4, v5, v6)
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.l) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test2
+
+ subroutine test3(v1, v2, v3, v4, v5, v6, v7, v8, v9)
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6, v7, v8, v9
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ !$acc parallel copyin (v1, v2, v3, v4, v5, v6, v7, v8, v9)
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.l) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test3
+
+ subroutine test4
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6, v7, v8, v9
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ v1 = 2
+ v2 = 6
+ v3 = -2
+ v4 = 4
+ v5 = 13
+ v6 = 18
+ v7 = 1
+ v8 = 1
+ v9 = 1
+ !$acc parallel copyin (v1, v2, v3, v4, v5, v6, v7, v8, v9)
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.r) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test4
+
+ subroutine test5(v1, v2, v3, v4, v5, v6)
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ !$acc parallel copyin (v1, v2, v3, v4, v5, v6)
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.r) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2
+ do j = v3, v4
+ do k = v5, v6
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test5
+
+ subroutine test6(v1, v2, v3, v4, v5, v6, v7, v8, v9)
+ integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19)
+ integer :: v1, v2, v3, v4, v5, v6, v7, v8, v9
+ logical :: l, r
+ l = .false.
+ r = .false.
+ a(:, :, :) = 0
+ b(:, :, :) = 0
+ !$acc parallel copyin (v1, v2, v3, v4, v5, v6, v7, v8, v9)
+ !$acc loop collapse (3) reduction (.or.:l)
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ l = l.or.k.lt.13.or.k.gt.18
+ if (.not.l) a(i, j, k) = a(i, j, k) + 1
+ m = i * 100 + j * 10 + k
+ end do
+ end do
+ end do
+ !$acc end parallel
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4
+ r = r.or.k.lt.13.or.k.gt.18
+ if (.not.r) b(i, j, k) = b(i, j, k) + 1
+ end do
+ end do
+ end do
+ if (l .neqv. r) call abort
+ do i = v1, v2, v7
+ do j = v3, v4, v8
+ do k = v5, v6, v9
+ if (a(i, j, k) .ne. b(i, j, k)) call abort
+ end do
+ end do
+ end do
+ end subroutine test6
+
+end program sub_collapse_3