@@ -3545,6 +3545,7 @@ void gfc_resolve_oacc_declare (gfc_namespace *);
void gfc_resolve_oacc_parallel_loop_blocks (gfc_code *, gfc_namespace *);
void gfc_resolve_oacc_blocks (gfc_code *, gfc_namespace *);
void gfc_resolve_oacc_routines (gfc_namespace *);
+void gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *);
/* expr.c */
void gfc_free_actual_arglist (gfc_actual_arglist *);
@@ -289,6 +289,10 @@ Wopenacc-parallelism
Fortran
; Documented in C
+Wopenacc-kernels-annotate-loops
+Fortran
+; Documented in C
+
Wopenmp-simd
Fortran
; Documented in C
@@ -695,6 +699,10 @@ fopenacc-dim=
Fortran LTO Joined Var(flag_openacc_dims)
; Documented in C
+fopenacc-kernels-annotate-loops
+Fortran LTO Optimization
+; Documented in C
+
fopenmp
Fortran LTO
; Documented in C
@@ -29,6 +29,7 @@ along with GCC; see the file COPYING3. If not see
#include "diagnostic.h"
#include "gomp-constants.h"
#include "target-memory.h" /* For gfc_encode_character. */
+#include "options.h"
/* Match an end of OpenMP directive. End of OpenMP directive is optional
whitespace, followed by '\n' or comment '!'. */
@@ -9090,3 +9091,366 @@ gfc_resolve_omp_udrs (gfc_symtree *st)
for (omp_udr = st->n.omp_udr; omp_udr; omp_udr = omp_udr->next)
gfc_resolve_omp_udr (omp_udr);
}
+
+
+/* The following functions implement automatic recognition and annotation of
+ DO loops in OpenACC kernels regions. Inside a kernels region, a nest of
+ DO loops that does not contain any annotated OpenACC loops, nor EXIT
+ or GOTO statements, gets an automatic "acc loop auto" annotation
+ on each loop.
+ This feature is controlled by flag_openacc_kernels_annotate_loops. */
+
+
+/* State of annotation state traversal for DO loops in kernels regions. */
+enum annotation_state {
+ as_outer,
+ as_in_kernels_region,
+ as_in_kernels_loop,
+ as_in_kernels_inner_loop
+};
+
+/* Return status of annotation traversal. */
+enum annotation_result {
+ ar_ok,
+ ar_invalid_loop,
+ ar_invalid_nest
+};
+
+/* Code walk function for check_for_invalid_calls. */
+
+static int
+check_code_for_invalid_calls (gfc_code **codep, int *walk_subtrees,
+ void *data ATTRIBUTE_UNUSED)
+{
+ gfc_code *code = *codep;
+ switch (code->op)
+ {
+ case EXEC_CALL:
+ /* Calls to openacc routines are permitted. */
+ if (code->resolved_sym
+ && (code->resolved_sym->attr.oacc_routine_lop
+ != OACC_ROUTINE_LOP_NONE))
+ return 0;
+ /* Else fall through. */
+
+ case EXEC_CALL_PPC:
+ case EXEC_ASSIGN_CALL:
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "Subroutine call at %L prevents annotation of loop nest",
+ &code->loc);
+ *walk_subtrees = 0;
+ return 1;
+
+ default:
+ return 0;
+ }
+}
+
+/* Expr walk function for check_for_invalid_calls. */
+
+static int
+check_expr_for_invalid_calls (gfc_expr **exprp, int *walk_subtrees,
+ void *data ATTRIBUTE_UNUSED)
+{
+ gfc_expr *expr = *exprp;
+ switch (expr->expr_type)
+ {
+ case EXPR_FUNCTION:
+ if (expr->value.function.esym
+ && (expr->value.function.esym->attr.oacc_routine_lop
+ != OACC_ROUTINE_LOP_NONE))
+ return 0;
+ /* Else fall through. */
+
+ case EXPR_COMPCALL:
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "Function call at %L prevents annotation of loop nest",
+ &expr->where);
+ *walk_subtrees = 0;
+ return 1;
+
+ default:
+ return 0;
+ }
+}
+
+/* Return TRUE if the DO loop CODE contains function or procedure
+ calls that ought to prohibit annotation. This traversal is
+ separate from the main annotation tree walk because we need to walk
+ expressions as well as executable statements. */
+
+static bool
+check_for_invalid_calls (gfc_code *code)
+{
+ gcc_assert (code->op == EXEC_DO);
+ return gfc_code_walker (&code, check_code_for_invalid_calls,
+ check_expr_for_invalid_calls, NULL);
+}
+
+/* Annotate DO loop CODE with OpenACC "loop auto". */
+
+static void
+annotate_do_loop (gfc_code *code, gfc_code *parent)
+{
+
+ /* A DO loop's body is another phony DO node whose next pointer starts
+ the actual body. */
+ gcc_assert (code->op == EXEC_DO);
+ gcc_assert (code->block->op == EXEC_DO);
+
+ /* Build the "acc loop auto" annotation and add the loop as its
+ body. */
+ gfc_omp_clauses *clauses = gfc_get_omp_clauses ();
+ clauses->par_auto = 1;
+ gfc_code *oacc_loop = gfc_get_code (EXEC_OACC_LOOP);
+ oacc_loop->block = gfc_get_code (EXEC_OACC_LOOP);
+ oacc_loop->block->next = code;
+ oacc_loop->ext.omp_clauses = clauses;
+ oacc_loop->loc = code->loc;
+ oacc_loop->block->loc = code->loc;
+
+ /* Splice the annotation into the place of the original loop. */
+ if (parent->block == code)
+ parent->block = oacc_loop;
+ else
+ {
+ gfc_code *prev = parent->block;
+ while (prev != code && prev->next != code)
+ {
+ prev = prev->next;
+ gcc_assert (prev != NULL);
+ }
+ prev->next = oacc_loop;
+ }
+ oacc_loop->next = code->next;
+ code->next = NULL;
+}
+
+/* Recursively traverse CODE in block PARENT, finding OpenACC kernels
+ regions. GOTO_TARGETS keeps track of statement labels that are
+ targets of gotos in the current function, while STATE keeps track
+ of the current context of the traversal. If the traversal
+ encounters a DO loop inside a kernels region, annotate it with
+ OpenACC loop directives if appropriate. Return the status of the
+ traversal. */
+
+static enum annotation_result
+annotate_do_loops_in_kernels (gfc_code *code, gfc_code *parent,
+ hash_set <gfc_st_label *> *goto_targets,
+ annotation_state state)
+{
+ gfc_code *next_code = NULL;
+ enum annotation_result retval = ar_ok;
+
+ for ( ; code; code = next_code)
+ {
+ bool walk_block = true;
+ next_code = code->next;
+
+ if (state >= as_in_kernels_loop
+ && code->here && goto_targets->contains (code->here))
+ /* This statement has a label that is the target of a GOTO or some
+ other jump. Do not try to sort out the details, just reject
+ this loop nest. */
+ {
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "Possible control transfer to label at %L "
+ "prevents annotation of loop nest",
+ &code->loc);
+ return ar_invalid_nest;
+ }
+
+ switch (code->op)
+ {
+ case EXEC_OACC_KERNELS:
+ /* Enter kernels region. */
+ annotate_do_loops_in_kernels (code->block->next, code,
+ goto_targets,
+ as_in_kernels_region);
+ walk_block = false;
+ break;
+
+ case EXEC_OACC_PARALLEL_LOOP:
+ case EXEC_OACC_PARALLEL:
+ case EXEC_OACC_KERNELS_LOOP:
+ case EXEC_OACC_LOOP:
+ /* Do not try to add automatic OpenACC annotations inside manually
+ annotated loops. Presumably, the user avoided doing it on
+ purpose; for example, all available levels of parallelism may
+ have been used up. */
+ if (state >= as_in_kernels_region)
+ {
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "Explicit loop annotation at %L "
+ "prevents annotation of loop nest",
+ &code->loc);
+ return ar_invalid_nest;
+ }
+ walk_block = false;
+ break;
+
+ case EXEC_DO:
+ if (state >= as_in_kernels_region)
+ {
+ /* A DO loop's body is another phony DO node whose next
+ pointer starts the actual body. Skip the phony node. */
+ gcc_assert (code->block->op == EXEC_DO);
+ enum annotation_result result
+ = annotate_do_loops_in_kernels (code->block->next, code,
+ goto_targets,
+ as_in_kernels_loop);
+ /* Check for function/procedure calls in the body of the
+ loop that would prevent parallelization. Unlike in C/C++,
+ we do not have to check that there is no modification of
+ the loop variable or loop count since they are already
+ handled by the semantics of DO loops in the FORTRAN
+ language. */
+ if (result != ar_invalid_nest && check_for_invalid_calls (code))
+ result = ar_invalid_nest;
+ if (result == ar_ok)
+ annotate_do_loop (code, parent);
+ else if (result == ar_invalid_nest
+ && state >= as_in_kernels_loop)
+ /* The outer loop is invalid, too, so stop traversal. */
+ return result;
+ walk_block = false;
+ }
+ break;
+
+ case EXEC_DO_WHILE:
+ case EXEC_DO_CONCURRENT:
+ /* Traverse the body in a special state to allow EXIT statements
+ from these loops. */
+ if (state >= as_in_kernels_loop)
+ {
+ enum annotation_result result
+ = annotate_do_loops_in_kernels (code->block, code,
+ goto_targets,
+ as_in_kernels_inner_loop);
+ if (result == ar_invalid_nest)
+ return result;
+ else if (result != ar_ok)
+ retval = result;
+ walk_block = false;
+ }
+ break;
+
+ case EXEC_GOTO:
+ case EXEC_ARITHMETIC_IF:
+ case EXEC_STOP:
+ case EXEC_ERROR_STOP:
+ /* A jump that may leave this loop. */
+ if (state >= as_in_kernels_loop)
+ {
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "Possible unstructured control flow at %L "
+ "prevents annotation of loop nest",
+ &code->loc);
+ return ar_invalid_nest;
+ }
+ break;
+
+ case EXEC_RETURN:
+ /* A return from a kernels region is diagnosed elsewhere as a
+ hard error, so no warning is needed here. */
+ if (state >= as_in_kernels_loop)
+ return ar_invalid_nest;
+ break;
+
+ case EXEC_EXIT:
+ if (state == as_in_kernels_loop)
+ {
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "Exit at %L prevents annotation of loop",
+ &code->loc);
+ retval = ar_invalid_loop;
+ }
+ break;
+
+ case EXEC_BACKSPACE:
+ case EXEC_CLOSE:
+ case EXEC_ENDFILE:
+ case EXEC_FLUSH:
+ case EXEC_INQUIRE:
+ case EXEC_OPEN:
+ case EXEC_READ:
+ case EXEC_REWIND:
+ case EXEC_WRITE:
+ /* Executing side-effecting I/O statements in parallel doesn't
+ make much sense. If this is what users want, they can always
+ add explicit annotations on the loop nest. */
+ if (state >= as_in_kernels_loop)
+ {
+ gfc_warning (OPT_Wopenacc_kernels_annotate_loops,
+ "I/O statement at %L prevents annotation of loop",
+ &code->loc);
+ return ar_invalid_nest;
+ }
+ break;
+
+ default:
+ break;
+ }
+
+ /* Visit nested statements, if any, returning early if we hit
+ any problems. */
+ if (walk_block)
+ {
+ enum annotation_result result
+ = annotate_do_loops_in_kernels (code->block, code,
+ goto_targets, state);
+ if (result == ar_invalid_nest)
+ return result;
+ else if (result != ar_ok)
+ retval = result;
+ }
+ }
+ return retval;
+}
+
+/* Traverse CODE to find all the labels referenced by GOTO and similar
+ statements and store them in GOTO_TARGETS. */
+
+static void
+compute_goto_targets (gfc_code *code, hash_set <gfc_st_label *> *goto_targets)
+{
+ for ( ; code; code = code->next)
+ {
+ switch (code->op)
+ {
+ case EXEC_GOTO:
+ case EXEC_LABEL_ASSIGN:
+ goto_targets->add (code->label1);
+ gcc_fallthrough ();
+
+ case EXEC_ARITHMETIC_IF:
+ goto_targets->add (code->label2);
+ goto_targets->add (code->label3);
+ gcc_fallthrough ();
+
+ default:
+ /* Visit nested statements, if any. */
+ if (code->block != NULL)
+ compute_goto_targets (code->block, goto_targets);
+ }
+ }
+}
+
+/* Find DO loops in OpenACC kernels regions that do not have OpenACC
+ annotations but look like they might benefit from automatic
+ parallelization. Add "acc loop auto" annotations for them. Assumes
+ flag_openacc_kernels_annotate_loops is set. */
+
+void
+gfc_oacc_annotate_loops_in_kernels_regions (gfc_namespace *ns)
+{
+ if (ns->proc_name)
+ {
+ hash_set <gfc_st_label *> goto_targets;
+ compute_goto_targets (ns->code, &goto_targets);
+ annotate_do_loops_in_kernels (ns->code, NULL, &goto_targets, as_outer);
+ }
+
+ for (ns = ns->contained; ns; ns = ns->sibling)
+ gfc_oacc_annotate_loops_in_kernels_regions (ns);
+}
@@ -6912,6 +6912,15 @@ done:
if (flag_c_prototypes || flag_c_prototypes_external)
fprintf (stdout, "\n#ifdef __cplusplus\n}\n#endif\n");
+ /* Add annotations on loops in OpenACC kernels regions if requested. This
+ is most easily done on this representation close to the source code. */
+ if (flag_openacc && flag_openacc_kernels_annotate_loops)
+ {
+ gfc_current_ns = gfc_global_ns_list;
+ for (; gfc_current_ns; gfc_current_ns = gfc_current_ns->sibling)
+ gfc_oacc_annotate_loops_in_kernels_regions (gfc_current_ns);
+ }
+
/* Do the translation. */
translate_all_program_units (gfc_global_ns_list);
@@ -2,6 +2,7 @@
! OpenACC kernels.
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
@@ -2,6 +2,7 @@
! kernels.
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
@@ -1,4 +1,5 @@
! { dg-options "-fopenacc -fdump-tree-omplower" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
module consts
integer, parameter :: n = 100
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
new file mode 100644
@@ -0,0 +1,33 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that all loops in the nest are annotated.
+
+subroutine f (a, b, c)
+ implicit none
+
+ real, intent (in), dimension(16,16) :: a
+ real, intent (in), dimension(16,16) :: b
+ real, intent (out), dimension(16,16) :: c
+
+ integer :: i, j, k
+ real :: t
+
+!$acc kernels copyin(a(1:16,1:16), b(1:16,1:16)) copyout(c(1:16,1:16))
+
+ do i = 1, 16
+ do j = 1, 16
+ t = 0
+ do k = 1, 16
+ t = t + a(i,k) * b(k,j)
+ end do
+ c(i,j) = t;
+ end do
+ end do
+
+!$acc end kernels
+end subroutine f
+
+! { dg-final { scan-tree-dump-times "acc loop private\\(.\\) auto" 3 "original" } }
new file mode 100644
@@ -0,0 +1,32 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a random goto in the body can't be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ if (a(i) < 0 .or. b(i) < 0) then
+ go to 10 ! { dg-warning "Possible unstructured control flow" }
+ end if
+ t = t + a(i) * b(i)
+ end do
+
+10 f = t
+
+!$acc end kernels
+
+end function f
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-std=legacy" }
+! { dg-do compile }
+
+! Test that a loop with a random label in the body cannot be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ goto 10
+
+ do i = 1, 16
+10 t = t + a(i) * b(i) ! { dg-warning "Possible control transfer to label" }
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
new file mode 100644
@@ -0,0 +1,39 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that in a situation with nested loops, a problem that prevents
+! annotation of the inner loop only still allows the outer loop to be
+! annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i, j
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ do j = 1, 16
+ if (a(i) < 0 .or. b(j) < 0) then
+ exit ! { dg-warning "Exit" }
+ else
+ t = t + a(i) * b(j)
+ end if
+ end do
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
new file mode 100644
@@ -0,0 +1,38 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that in a situation with nested loops, a problem that prevents
+! annotation of the outer loop only still allows the inner loop to be
+! annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i, j
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ if (a(i) < 0) then
+ exit ! { dg-warning "Exit" }
+ end if
+ do j = 1, 16
+ t = t + a(i) * b(j)
+ end do
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
new file mode 100644
@@ -0,0 +1,35 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that an explicit annotation on an outer loop suppresses annotation
+! of inner loops, and produces a diagnostic.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i, j
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+!$acc loop seq ! { dg-warning "Explicit loop annotation" }
+ do i = 1, 16
+ do j = 1, 16
+ t = t + a(i) * b(j)
+ end do
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
new file mode 100644
@@ -0,0 +1,35 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that an explicit annotation on an inner loop suppresses annotation
+! of the outer loop, and produces a diagnostic.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i, j
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ !$acc loop seq ! { dg-warning "Explicit loop annotation" }
+ do j = 1, 16
+ t = t + a(i) * b(j)
+ end do
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that loops containing I/O statements can't be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i, j
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ do j = 1, 16
+ print *, " i =", i, " j =", j ! { dg-warning "I/O statement" }
+ t = t + a(i) * b(j)
+ end do
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
new file mode 100644
@@ -0,0 +1,32 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a variable bound can be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (:) :: a, b
+
+ integer :: i, n
+ real :: t
+
+ t = 0.0
+ n = size (a)
+
+!$acc kernels
+
+ do i = 1, n
+ t = t + a(i) * b(i)
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
new file mode 100644
@@ -0,0 +1,33 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a conditional in the body can be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ if (a(i) > 0 .and. b(i) > 0) then
+ t = t + a(i) * b(i)
+ end if
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a case construct in the body can be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+!$acc kernels
+
+ do i = 1, 16
+ select case (i)
+ case (1)
+ t = a(i) * b(i)
+ case default
+ t = t + a(i) * b(i)
+ end select
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
new file mode 100644
@@ -0,0 +1,35 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a cycle statement in the body can be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ if (a(i) < 0 .or. b(i) < 0) then
+ cycle
+ end if
+ t = t + a(i) * b(i)
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 1 "original" } }
+
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a exit statement in the body cannot be annotated.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ if (a(i) < 0 .or. b(i) < 0) then
+ exit ! { dg-warning "Exit" }
+ end if
+ t = t + a(i) * b(i)
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
new file mode 100644
@@ -0,0 +1,48 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a random function call in the body cannot
+! be annotated.
+
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ interface
+ function g (x)
+ real :: g
+ real, intent (in) :: x
+ end function g
+
+ subroutine h (x)
+ real, intent (in) :: x
+ end subroutine h
+ end interface
+
+ t = 0.0
+
+!$acc kernels
+ do i = 1, 16
+ t = t + g (a(i) * b(i)) ! { dg-warning "Function call" }
+ end do
+
+ do i = 1, 16
+ call h (t) ! { dg-warning "Subroutine call" }
+ t = t + a(i) * b(i)
+ end do
+
+ f = t
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private.* auto" 0 "original" } }
+
new file mode 100644
@@ -0,0 +1,50 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a call to a declared openacc function/subroutine
+! can be annotated.
+
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ interface
+ function g (x)
+ !$acc routine worker
+ real :: g
+ real, intent (in) :: x
+ end function g
+
+ subroutine h (x)
+ !$acc routine worker
+ real, intent (in) :: x
+ end subroutine h
+ end interface
+
+ t = 0.0
+
+!$acc kernels
+ do i = 1, 16
+ t = t + g (a(i) * b(i))
+ end do
+
+ do i = 1, 16
+ call h (t)
+ t = t + a(i) * b(i)
+ end do
+
+ f = t
+!$acc end kernels
+
+end function f
+
+! { dg-final { scan-tree-dump-times "acc loop private\\(i\\) auto" 2 "original" } }
+
new file mode 100644
@@ -0,0 +1,34 @@
+! { dg-additional-options "-fopenacc -fopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-Wopenacc-kernels-annotate-loops" }
+! { dg-additional-options "-fdump-tree-original" }
+! { dg-do compile }
+
+! Test that a loop with a return statement in the body gives a hard
+! error.
+
+function f (a, b)
+ implicit none
+
+ real :: f
+ real, intent (in), dimension (16) :: a, b
+
+ integer :: i
+ real :: t
+
+ t = 0.0
+
+!$acc kernels
+
+ do i = 1, 16
+ if (a(i) < 0 .or. b(i) < 0) then
+ f = 0.0
+ return ! { dg-error "invalid branch" }
+ end if
+ t = t + a(i) * b(i)
+ end do
+
+ f = t
+
+!$acc end kernels
+
+end function f
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
@@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fno-openacc-kernels-annotate-loops" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-optimized" }
From: Sandra Loosemore <sandra@codesourcery.com> This patch implements the Fortran support for adding "#pragma acc loop auto" annotations to loops in OpenACC kernels regions. It implements the same -fopenacc-kernels-annotate-loops and -Wopenacc-kernels-annotate-loops options that were previously added (and documented) for the C/C++ front ends. Co-Authored-By: Gergö Barany <gergo@codesourcery.com> gcc/fortran/ * gfortran.h (gfc_oacc_annotate_loops_in_kernels_regions): Declare. * lang.opt (Wopenacc-kernels-annotate-loops): New. (fopenacc-kernels-annotate-loops): New. * openmp.c: Include options.h. (enum annotation_state, enum annotation_result): New. (check_code_for_invalid_calls): New. (check_expr_for_invalid_calls): New. (check_for_invalid_calls): New. (annotate_do_loop): New. (annotate_do_loops_in_kernels): New. (compute_goto_targets): New. (gfc_oacc_annotate_loops_in_kernels_regions): New. * parse.c (gfc_parse_file): Handle -fopenacc-kernels-annotate-loops. gcc/testsuite/ * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Add -fno-openacc-kernels-annotate-loops option. * gfortran.dg/goacc/classify-kernels.f95: Likewise. * gfortran.dg/goacc/common-block-3.f90: Likewise. * gfortran.dg/goacc/kernels-loop-2.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-2.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-enter-exit-2.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-enter-exit.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data-update.f95: Likewise. * gfortran.dg/goacc/kernels-loop-data.f95: Likewise. * gfortran.dg/goacc/kernels-loop-n.f95: Likewise. * gfortran.dg/goacc/kernels-loop.f95: Likewise. * gfortran.dg/goacc/kernels-parallel-loop-data-enter-exit.f95: Likewise. * gfortran.dg/goacc/kernels-loop-annotation-1.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-2.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-3.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-4.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-5.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-6.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-7.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-8.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-9.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-10.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-11.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-12.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-13.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-14.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-15.f95: New. * gfortran.dg/goacc/kernels-loop-annotation-16.f95: New. --- gcc/fortran/gfortran.h | 1 + gcc/fortran/lang.opt | 8 + gcc/fortran/openmp.c | 364 ++++++++++++++++++ gcc/fortran/parse.c | 9 + .../goacc/classify-kernels-unparallelized.f95 | 1 + .../gfortran.dg/goacc/classify-kernels.f95 | 1 + .../gfortran.dg/goacc/common-block-3.f90 | 1 + .../gfortran.dg/goacc/kernels-loop-2.f95 | 1 + .../goacc/kernels-loop-annotation-1.f95 | 33 ++ .../goacc/kernels-loop-annotation-10.f95 | 32 ++ .../goacc/kernels-loop-annotation-11.f95 | 34 ++ .../goacc/kernels-loop-annotation-12.f95 | 39 ++ .../goacc/kernels-loop-annotation-13.f95 | 38 ++ .../goacc/kernels-loop-annotation-14.f95 | 35 ++ .../goacc/kernels-loop-annotation-15.f95 | 35 ++ .../goacc/kernels-loop-annotation-16.f95 | 34 ++ .../goacc/kernels-loop-annotation-2.f95 | 32 ++ .../goacc/kernels-loop-annotation-3.f95 | 33 ++ .../goacc/kernels-loop-annotation-4.f95 | 34 ++ .../goacc/kernels-loop-annotation-5.f95 | 35 ++ .../goacc/kernels-loop-annotation-6.f95 | 34 ++ .../goacc/kernels-loop-annotation-7.f95 | 48 +++ .../goacc/kernels-loop-annotation-8.f95 | 50 +++ .../goacc/kernels-loop-annotation-9.f95 | 34 ++ .../gfortran.dg/goacc/kernels-loop-data-2.f95 | 1 + .../goacc/kernels-loop-data-enter-exit-2.f95 | 1 + .../goacc/kernels-loop-data-enter-exit.f95 | 1 + .../goacc/kernels-loop-data-update.f95 | 1 + .../gfortran.dg/goacc/kernels-loop-data.f95 | 1 + .../gfortran.dg/goacc/kernels-loop-n.f95 | 1 + .../gfortran.dg/goacc/kernels-loop.f95 | 1 + .../kernels-parallel-loop-data-enter-exit.f95 | 1 + 32 files changed, 974 insertions(+) create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-1.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-10.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-11.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-12.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-13.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-14.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-15.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-16.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-2.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-3.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-4.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-5.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-6.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-7.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-8.f95 create mode 100644 gcc/testsuite/gfortran.dg/goacc/kernels-loop-annotation-9.f95 -- 2.33.0 ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955