2017-02-20 Cesar Philippidis <cesar@codesourcery.com>
gcc/
* common.opt (finform-parallelism): New option.
* omp-low.c (inform_oacc_loop): New function.
(execute_oacc_device_lower): Use it to report how ACC LOOPs are
assigned parallelism.
gcc/doc/
* invoke.texi: Document -finform-parallelism.
gcc/fortran/
* trans-openmp.c (gfc_trans_omp_clauses_1): Delete unused orig_decl.
(gfc_trans_oacc_combined_directive): Set the location of
combined acc loops.
gcc/testsuite/
* c-c++-common/goacc/inform-parallelism.c: New test.
* gfortran.dg/goacc/inform-parallelism.f90: New test.
@@ -1451,6 +1451,10 @@ fif-conversion2
Common Report Var(flag_if_conversion2) Optimization
Perform conversion of conditional jumps to conditional execution.
+finform-parallelism
+Common Var(flag_inform_parallelism) Init(0)
+Report all paralllelism detected inside offloaded regions.
+
fstack-reuse=
Common Joined RejectNegative Enum(stack_reuse_level) Var(flag_stack_reuse) Init(SR_ALL) Optimization
-fstack-reuse=[all|named_vars|none] Set stack reuse level for local variables.
@@ -351,7 +351,7 @@ Objective-C and Objective-C++ Dialects}.
-fforward-propagate -ffp-contract=@var{style} -ffunction-sections @gol
-fgcse -fgcse-after-reload -fgcse-las -fgcse-lm -fgraphite-identity @gol
-fgcse-sm -fhoist-adjacent-loads -fif-conversion @gol
--fif-conversion2 -findirect-inlining @gol
+-fif-conversion2 -findirect-inlining -finform-parallelism @gol
-finline-functions -finline-functions-called-once -finline-limit=@var{n} @gol
-finline-small-functions -fipa-cp -fipa-cp-clone -fipa-cp-alignment @gol
-fipa-pta -fipa-profile -fipa-pure-const -fipa-reference -fipa-icf @gol
@@ -6428,6 +6428,13 @@ or @option{-finline-small-functions} options.
Enabled at level @option{-O2}.
+@item -finform-parallelism
+@opindex finform-parallelism
+Report any parallelism detected by the compiler. Inside OpenACC
+offloaded regions, this includes the gang, worker and vector level
+parallelism associated with any @code{ACC LOOP}s. This option is disabled
+by default.
+
@item -finline-functions
@opindex finline-functions
Consider all functions for inlining, even if they are not declared inline.
@@ -1947,7 +1947,6 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
&& n->expr->ref->next->u.ar.type == AR_FULL)))
{
gfc_ref *ref = n->expr->ref;
- tree orig_decl = decl;
gfc_component *c = ref->u.c.component;
tree field;
tree context;
@@ -3819,6 +3818,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
enum tree_code construct_code;
bool scan_nodesc_arrays = false;
hash_set<gfc_symbol *> *array_set = NULL;
+ location_t loc = input_location;
switch (code->op)
{
@@ -3850,6 +3850,9 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
pushlevel ();
stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, loop_clauses, NULL);
+ if (CAN_HAVE_LOCATION_P (stmt))
+ SET_EXPR_LOCATION (stmt, loc);
+
if (array_set && array_set->elements ())
gfc_add_expr_to_block (&inner, stmt);
@@ -3865,8 +3868,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
delete array_set;
}
- stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
- oacc_clauses);
+ stmt = build2_loc (loc, construct_code, void_type_node, stmt, oacc_clauses);
gfc_add_expr_to_block (&block, stmt);
gfc_free_omp_clauses (loop_clauses);
@@ -20399,6 +20399,28 @@ debug_oacc_loop (oacc_loop *loop)
dump_oacc_loop (stderr, loop, 0);
}
+/* Provide diagnostics on OpenACC loops LOOP, its siblings and its
+ children. */
+
+static void
+inform_oacc_loop (oacc_loop *loop)
+{
+ const char *seq = loop->mask == 0 ? " SEQ" : "";
+ const char *gang = loop->mask & GOMP_DIM_MASK (GOMP_DIM_GANG)
+ ? " GANG" : "";
+ const char *worker = loop->mask & GOMP_DIM_MASK (GOMP_DIM_WORKER)
+ ? " WORKER" : "";
+ const char *vector = loop->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR)
+ ? " VECTOR" : "";
+
+ inform (loop->loc, "ACC LOOP%s%s%s%s", seq, gang, worker, vector);
+
+ if (loop->child)
+ inform_oacc_loop (loop->child);
+ if (loop->sibling)
+ inform_oacc_loop (loop->sibling);
+}
+
/* DFS walk of basic blocks BB onwards, creating OpenACC loop
structures as we go. By construction these loops are properly
nested. */
@@ -21069,6 +21091,8 @@ execute_oacc_device_lower ()
dump_oacc_loop (dump_file, loops, 0);
fprintf (dump_file, "\n");
}
+ if (flag_inform_parallelism && loops->child)
+ inform_oacc_loop (loops->child);
/* Offloaded targets may introduce new basic blocks, which require
dominance information to update SSA. */
new file mode 100644
@@ -0,0 +1,61 @@
+/* Test the output of -finform-parallelism. */
+
+/* { dg-additional-options "-finform-parallelism" } */
+
+int
+main ()
+{
+ int x, y, z;
+
+#pragma acc parallel loop seq /* { dg-message "ACC LOOP SEQ" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop gang /* { dg-message "ACC LOOP GANG" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop worker /* { dg-message "ACC LOOP WORKER" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop vector /* { dg-message "ACC LOOP VECTOR" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop gang vector /* { dg-message "ACC LOOP GANG VECTOR" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop gang worker /* { dg-message "ACC LOOP GANG WORKER" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop worker vector /* { dg-message "ACC LOOP WORKER VECTOR" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop gang worker vector /* { dg-message "ACC LOOP GANG WORKER VECTOR" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop /* { dg-message "ACC LOOP GANG VECTOR" } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc parallel loop /* { dg-message "ACC LOOP GANG WORKER" } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop /* { dg-message "ACC LOOP VECTOR" } */
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc parallel loop gang /* { dg-message "ACC LOOP GANG" } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop worker /* { dg-message "ACC LOOP WORKER" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop vector /* { dg-message "ACC LOOP VECTOR" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,62 @@
+! Test the output of -finform-parallellism.
+
+! { dg-additional-options "-finform-parallelism" }
+
+program test
+ implicit none
+
+ integer x, y, z
+
+ !$acc parallel loop seq ! { dg-message "ACC LOOP SEQ" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop gang ! { dg-message "ACC LOOP GANG" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop worker ! { dg-message "ACC LOOP WORKER" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop vector ! { dg-message "ACC LOOP VECTOR" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop gang vector ! { dg-message "ACC LOOP GANG VECTOR" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop gang worker ! { dg-message "ACC LOOP GANG WORKER" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop worker vector ! { dg-message "ACC LOOP WORKER VECTOR" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop gang worker vector ! { dg-message "ACC LOOP GANG WORKER VECTOR" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop ! { dg-message "ACC LOOP GANG VECTOR" }
+ do x = 1, 10
+ end do
+
+ !$acc parallel loop ! { dg-message "ACC LOOP GANG WORKER" }
+ do x = 1, 10
+ !$acc loop ! { dg-message "ACC LOOP VECTOR" }
+ do y = 1, 10
+ end do
+ end do
+
+ !$acc parallel loop gang ! { dg-message "ACC LOOP GANG" }
+ do x = 1, 10
+ !$acc loop worker ! { dg-message "ACC LOOP WORKER" }
+ do y = 1, 10
+ !$acc loop vector ! { dg-message "ACC LOOP VECTOR" }
+ do z = 1, 10
+ end do
+ end do
+ end do
+end program test