+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
}
}
-/* 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)
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. */
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. */
+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
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" } */
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)];
}
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" } */
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];
}
parallel. */
/* { dg-additional-options "-O2" }
+ { dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
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];
}
routine. */
/* { dg-additional-options "-O2" }
+ { dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-oaccdevlow" } */
#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];
}
+/* { 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;
int
kernels_eternal (void)
{
-#pragma acc kernels
+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
{
while (1)
;
int
kernels_noreturn (void)
{
-#pragma acc kernels
+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loop parallelism" } */
__builtin_abort ();
return 0;
{
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++)
/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
/* { dg-additional-options "-fdump-tree-parloops1-all" } */
/* { dg-additional-options "-fdump-tree-optimized" } */
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)
/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
/* { dg-additional-options "-fdump-tree-parloops1-all" } */
/* { dg-additional-options "-fdump-tree-optimized" } */
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)
--- /dev/null
+/* 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;
+}
! 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" }
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
! 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" }
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
! parallel.
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
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
! routine.
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
! { dg-additional-options "-fdump-tree-ompexp" }
! { dg-additional-options "-fdump-tree-oaccdevlow" }
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
! { dg-additional-options "-O2" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
program main
implicit none
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
--- /dev/null
+! 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