Add user-friendly diagnostics for OpenACC loop parallelism assigned

gcc/
	* omp-offload.c (inform_oacc_loop): New function.
	(execute_oacc_device_lower): Use it to display loop parallelism.
	gcc/testsuite/
	* c-c++-common/goacc/note-parallelism.c: New test.
	* gfortran.dg/goacc/note-parallelism.f90: New test.
	* c-c++-common/goacc/classify-kernels-unparallelized.c: Update.
	* c-c++-common/goacc/classify-kernels.c: Likewise.
	* c-c++-common/goacc/classify-parallel.c: Likewise.
	* c-c++-common/goacc/classify-routine.c: Likewise.
	* c-c++-common/goacc/kernels-1.c: Likewise.
	* c-c++-common/goacc/kernels-double-reduction-n.c: Likewise.
	* c-c++-common/goacc/kernels-double-reduction.c: Likewise.
	* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise.
	* gfortran.dg/goacc/classify-kernels.f95: Likewise.
	* gfortran.dg/goacc/classify-parallel.f95: Likewise.
	* gfortran.dg/goacc/classify-routine.f95: Likewise.
	* gfortran.dg/goacc/kernels-loop-inner.f95: Likewise.

Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>

From-SVN: r267146
This commit is contained in:
Thomas Schwinge 2018-12-14 21:41:58 +01:00 committed by Thomas Schwinge
parent 890b87d174
commit 5d390fd3ae
17 changed files with 346 additions and 16 deletions

View File

@ -1,3 +1,9 @@
2018-12-14 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
* omp-offload.c (inform_oacc_loop): New function.
(execute_oacc_device_lower): Use it to display loop parallelism.
2018-12-14 Jakub Jelinek <jakub@redhat.com>
PR c++/82294

View File

@ -823,7 +823,7 @@ dump_oacc_loop_part (FILE *file, gcall *from, int depth,
}
}
/* Dump OpenACC loops LOOP, its siblings and its children. */
/* Dump OpenACC loop LOOP, its children, and its siblings. */
static void
dump_oacc_loop (FILE *file, oacc_loop *loop, int depth)
@ -866,6 +866,31 @@ debug_oacc_loop (oacc_loop *loop)
dump_oacc_loop (stderr, loop, 0);
}
/* Provide diagnostics on OpenACC loop LOOP, its children, and its
siblings. */
static void
inform_oacc_loop (const oacc_loop *loop)
{
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" : "";
const char *seq = loop->mask == 0 ? " seq" : "";
const dump_user_location_t loc
= dump_user_location_t::from_location_t (loop->loc);
dump_printf_loc (MSG_OPTIMIZED_LOCATIONS, loc,
"assigned OpenACC%s%s%s%s loop parallelism\n", gang, worker,
vector, seq);
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. */
@ -1533,6 +1558,28 @@ execute_oacc_device_lower ()
dump_oacc_loop (dump_file, loops, 0);
fprintf (dump_file, "\n");
}
if (dump_enabled_p ())
{
oacc_loop *l = loops;
/* OpenACC kernels constructs are special: they currently don't use the
generic oacc_loop infrastructure. */
if (is_oacc_kernels)
{
/* Create a fake oacc_loop for diagnostic purposes. */
l = new_oacc_loop_raw (NULL,
DECL_SOURCE_LOCATION (current_function_decl));
l->mask = used_mask;
}
else
{
/* Skip the outermost, dummy OpenACC loop */
l = l->child;
}
if (l)
inform_oacc_loop (l);
if (is_oacc_kernels)
free_oacc_loop (l);
}
/* Offloaded targets may introduce new basic blocks, which require
dominance information to update SSA. */

View File

@ -1,3 +1,21 @@
2018-12-14 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
* c-c++-common/goacc/note-parallelism.c: New test.
* gfortran.dg/goacc/note-parallelism.f90: New test.
* c-c++-common/goacc/classify-kernels-unparallelized.c: Update.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/classify-parallel.c: Likewise.
* c-c++-common/goacc/classify-routine.c: Likewise.
* c-c++-common/goacc/kernels-1.c: Likewise.
* c-c++-common/goacc/kernels-double-reduction-n.c: Likewise.
* c-c++-common/goacc/kernels-double-reduction.c: Likewise.
* gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise.
* gfortran.dg/goacc/classify-kernels.f95: Likewise.
* gfortran.dg/goacc/classify-parallel.f95: Likewise.
* gfortran.dg/goacc/classify-routine.f95: Likewise.
* gfortran.dg/goacc/kernels-loop-inner.f95: Likewise.
2018-12-14 Alexandre Oliva <aoliva@redhat.com>
PR c++/86823

View File

@ -2,6 +2,7 @@
OpenACC kernels. */
/* { dg-additional-options "-O2" }
{ dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-parloops1-all" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
@ -18,7 +19,7 @@ extern unsigned int f (unsigned int);
void KERNELS ()
{
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (unsigned int i = 0; i < N; i++)
c[i] = a[f (i)] + b[f (i)];
}

View File

@ -2,6 +2,7 @@
kernels. */
/* { dg-additional-options "-O2" }
{ dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-parloops1-all" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
@ -14,7 +15,7 @@ extern unsigned int *__restrict c;
void KERNELS ()
{
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
for (unsigned int i = 0; i < N; i++)
c[i] = a[i] + b[i];
}

View File

@ -2,6 +2,7 @@
parallel. */
/* { dg-additional-options "-O2" }
{ dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
@ -13,7 +14,7 @@ extern unsigned int *__restrict c;
void PARALLEL ()
{
#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N])
#pragma acc parallel loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
for (unsigned int i = 0; i < N; i++)
c[i] = a[i] + b[i];
}

View File

@ -2,6 +2,7 @@
routine. */
/* { dg-additional-options "-O2" }
{ dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
@ -15,7 +16,7 @@ extern unsigned int *__restrict c;
#pragma acc routine worker
void ROUTINE ()
{
#pragma acc loop
#pragma acc loop /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */
for (unsigned int i = 0; i < N; i++)
c[i] = a[i] + b[i];
}

View File

@ -1,7 +1,9 @@
/* { dg-additional-options "-fopt-info-optimized-omp" } */
int
kernels_empty (void)
{
#pragma acc kernels
#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
;
return 0;
@ -10,7 +12,7 @@ kernels_empty (void)
int
kernels_eternal (void)
{
#pragma acc kernels
#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
{
while (1)
;
@ -22,7 +24,7 @@ kernels_eternal (void)
int
kernels_noreturn (void)
{
#pragma acc kernels
#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
__builtin_abort ();
return 0;
@ -36,7 +38,7 @@ kernels_loop_ptr_it (void)
{
float *i;
#pragma acc kernels
#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
{
#pragma acc loop
for (i = &b[0][0][0]; i < &b[0][0][10]; i++)

View File

@ -1,4 +1,5 @@
/* { dg-additional-options "-O2" } */
/* { dg-additional-options "-fopt-info-optimized-omp" } */
/* { dg-additional-options "-fdump-tree-parloops1-all" } */
/* { dg-additional-options "-fdump-tree-optimized" } */
@ -14,7 +15,7 @@ foo (unsigned int n)
int i, j;
unsigned int sum = 1;
#pragma acc kernels copyin (a[0:n]) copy (sum)
#pragma acc kernels copyin (a[0:n]) copy (sum) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
{
for (i = 0; i < n; ++i)
for (j = 0; j < n; ++j)

View File

@ -1,4 +1,5 @@
/* { dg-additional-options "-O2" } */
/* { dg-additional-options "-fopt-info-optimized-omp" } */
/* { dg-additional-options "-fdump-tree-parloops1-all" } */
/* { dg-additional-options "-fdump-tree-optimized" } */
@ -14,7 +15,7 @@ foo (void)
int i, j;
unsigned int sum = 1;
#pragma acc kernels copyin (a[0:N]) copy (sum)
#pragma acc kernels copyin (a[0:N]) copy (sum) /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
{
for (i = 0; i < N; ++i)
for (j = 0; j < N; ++j)

View File

@ -0,0 +1,115 @@
/* Test the output of "-fopt-info-optimized-omp". */
/* { dg-additional-options "-fopt-info-optimized-omp" } */
/* See also "../../gfortran.dg/goacc/note-parallelism.f90". */
int
main ()
{
int x, y, z;
#pragma acc parallel
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop gang vector /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop gang worker /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop worker vector /* { dg-message "optimized: assigned OpenACC worker vector loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop gang worker vector /* { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop gang /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (z = 0; z < 10; z++)
;
#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
for (x = 0; x < 10; x++)
;
#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (y = 0; y < 10; y++)
;
#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC worker loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (z = 0; z < 10; z++)
;
#pragma acc parallel
for (x = 0; x < 10; x++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (z = 0; z < 10; z++)
;
#pragma acc parallel loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (z = 0; z < 10; z++)
;
#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (z = 0; z < 10; z++)
;
#pragma acc parallel loop /* { dg-message "optimized: assigned OpenACC gang worker loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC vector loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (z = 0; z < 10; z++)
;
#pragma acc parallel loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (x = 0; x < 10; x++)
#pragma acc loop /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
for (y = 0; y < 10; y++)
#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
for (z = 0; z < 10; z++)
;
return 0;
}

View File

@ -2,6 +2,7 @@
! OpenACC kernels.
! { dg-additional-options "-O2" }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
@ -19,7 +20,7 @@ program main
call setup(a, b)
!$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
do i = 0, n - 1
do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
c(i) = a(f (i)) + b(f (i))
end do
!$acc end kernels

View File

@ -2,6 +2,7 @@
! kernels.
! { dg-additional-options "-O2" }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-parloops1-all" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
@ -15,7 +16,7 @@ program main
call setup(a, b)
!$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
do i = 0, n - 1
do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
c(i) = a(i) + b(i)
end do
!$acc end kernels

View File

@ -2,6 +2,7 @@
! parallel.
! { dg-additional-options "-O2" }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
@ -13,7 +14,7 @@ program main
call setup(a, b)
!$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
!$acc parallel loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do

View File

@ -2,6 +2,7 @@
! routine.
! { dg-additional-options "-O2" }
! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
@ -13,7 +14,7 @@ subroutine ROUTINE
call setup(a, b)
!$acc loop
!$acc loop ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do

View File

@ -1,4 +1,5 @@
! { dg-additional-options "-O2" }
! { dg-additional-options "-fopt-info-optimized-omp" }
program main
implicit none
@ -6,7 +7,7 @@ program main
integer :: a(100,100), b(100,100)
integer :: i, j, d
!$acc kernels
!$acc kernels ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do i=1,100
do j=1,100
a(i,j) = 1

View File

@ -0,0 +1,131 @@
! Test the output of "-fopt-info-optimized-omp".
! { dg-additional-options "-fopt-info-optimized-omp" }
! See also "../../c-c++-common/goacc/note-parallelism.c".
program test
implicit none
integer x, y, z
!$acc parallel
do x = 1, 10
end do
!$acc end parallel
!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelis" }
do x = 1, 10
end do
!$acc parallel loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop gang vector ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop gang worker ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop worker vector ! { dg-message "optimized: assigned OpenACC worker vector loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop gang worker vector ! { dg-message "optimized: assigned OpenACC gang worker vector loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop gang ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
do x = 1, 10
!$acc loop worker ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
do y = 1, 10
!$acc loop vector ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do z = 1, 10
end do
end do
end do
!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
do x = 1, 10
end do
!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
do x = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do y = 1, 10
end do
end do
!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang loop parallelism" }
do x = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC worker loop parallelism" }
do y = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do z = 1, 10
end do
end do
end do
!$acc parallel
do x = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
do y = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do z = 1, 10
end do
end do
end do
!$acc end parallel
!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do x = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
do y = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do z = 1, 10
end do
end do
end do
!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
do x = 1, 10
!$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do y = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do z = 1, 10
end do
end do
end do
!$acc parallel loop ! { dg-message "optimized: assigned OpenACC gang worker loop parallelism" }
do x = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC vector loop parallelism" }
do y = 1, 10
!$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do z = 1, 10
end do
end do
end do
!$acc parallel loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do x = 1, 10
!$acc loop ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
do y = 1, 10
!$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }
do z = 1, 10
end do
end do
end do
end program test