2017-02-20 Cesar Philippidis 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. diff --git a/gcc/common.opt b/gcc/common.opt index 42c0b2f..a7e5494 100644 --- a/gcc/common.opt +++ b/gcc/common.opt @@ -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. diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi index fd8ba42..9cc8a8d 100644 --- a/gcc/doc/invoke.texi +++ b/gcc/doc/invoke.texi @@ -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. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 295f172..8688425 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -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 *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); diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 0f79533..6ea8738 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -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. */ diff --git a/gcc/testsuite/c-c++-common/goacc/inform-parallelism.c b/gcc/testsuite/c-c++-common/goacc/inform-parallelism.c new file mode 100644 index 0000000..b892bf0 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/inform-parallelism.c @@ -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; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/inform-parallelism.f90 b/gcc/testsuite/gfortran.dg/goacc/inform-parallelism.f90 new file mode 100644 index 0000000..6e11331 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/inform-parallelism.f90 @@ -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