Typo fix relative to last posted version.
2018-10-05 Nathan Sidwell <nathan@acm.org>
Tom de Vries <tdevries@suse.de>
Thomas Schwinge <thomas@codesourcery.com>
Julian Brown <julian@codesourcery.com>
gcc/
* doc/invoke.texi (fopenacc-dim): Update.
* omp-offload.c (oacc_parse_default_dims): Update.
(oacc_validate_dims): Emit warnings about strange partitioning choices.
gcc/testsuite/
* c-c++-common/goacc/acc-icf.c: Update.
* c-c++-common/goacc/parallel-dims-1.c: Likewise.
* c-c++-common/goacc/parallel-reduction.c: Likewise.
* c-c++-common/goacc/pr70688.c: Likewise.
* c-c++-common/goacc/routine-1.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/routine-4.f90: Likewise.
* gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Add -w.
* testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New.
* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update.
* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise.
* testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
* testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New.
(cherry picked from openacc-gcc-9-branch commit
0dd9e6ac878a47ebf021b700e56d4a8d31938bd1)
+2018-10-05 Nathan Sidwell <nathan@acm.org>
+ Tom de Vries <tdevries@suse.de>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * doc/invoke.texi (fopenacc-dim): Update.
+ * omp-offload.c (oacc_parse_default_dims): Update.
+ (oacc_validate_dims): Emit warnings about strange partitioning choices.
+
2019-09-20 Chung-Lin Tang <cltang@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
@cindex OpenACC accelerator programming
Specify default compute dimensions for parallel offload regions that do
not explicitly specify. The @var{geom} value is a triple of
-':'-separated sizes, in order 'gang', 'worker' and, 'vector'. A size
-can be omitted, to use a target-specific default value.
+':'-separated sizes, in order 'gang', 'worker' and, 'vector'. If a size
+is to be deferred until execution '-' can be used, alternatively a size
+can be omitted to use a target-specific default value. When deferring
+to runtime, the environment variable @var{GOMP_OPENACC_DIM} can be set.
+It has the same format as the option value, except that '-' is not
+permitted.
@item -fopenmp
@opindex fopenmp
}
/* Parse the default dimension parameter. This is a set of
- :-separated optional compute dimensions. Each specified dimension
- is a positive integer. When device type support is added, it is
+ :-separated optional compute dimensions. Each dimension is either
+ a positive integer, or '-' for a dynamic value computed at
+ runtime. When device type support is added, it is
planned to be a comma separated list of such compute dimensions,
with all but the first prefixed by the colon-terminated device
type. */
if (*pos != ':')
{
- long val;
- const char *eptr;
+ long val = 0;
- errno = 0;
- val = strtol (pos, CONST_CAST (char **, &eptr), 10);
- if (errno || val <= 0 || (int) val != val)
- goto malformed;
- pos = eptr;
+ if (*pos == '-')
+ pos++;
+ else
+ {
+ const char *eptr;
+
+ errno = 0;
+ val = strtol (pos, CONST_CAST (char **, &eptr), 10);
+ if (errno || val <= 0 || (int) val != val)
+ goto malformed;
+ pos = eptr;
+ }
oacc_default_dims[ix] = (int) val;
}
}
pos = TREE_CHAIN (pos);
}
+ bool check = true;
+#ifdef ACCEL_COMPILER
+ check = false;
+#endif
+ if (check
+ && !lookup_attribute ("oacc kernels", DECL_ATTRIBUTES (fn)))
+ {
+ static char const *const axes[] =
+ /* Must be kept in sync with GOMP_DIM enumeration. */
+ { "gang", "worker", "vector" };
+ for (ix = level >= 0 ? level : 0; ix != GOMP_DIM_MAX; ix++)
+ if (dims[ix] < 0)
+ ; /* Defaulting axis. */
+ else if ((used & GOMP_DIM_MASK (ix)) && dims[ix] == 1)
+ /* There is partitioned execution, but the user requested a
+ dimension size of 1. They're probably confused. */
+ warning_at (DECL_SOURCE_LOCATION (fn), 0,
+ "region contains %s partitioned code but"
+ " is not %s partitioned", axes[ix], axes[ix]);
+ else if (!(used & GOMP_DIM_MASK (ix)) && dims[ix] != 1)
+ /* The dimension is explicitly partitioned to non-unity, but
+ no use is made within the region. */
+ warning_at (DECL_SOURCE_LOCATION (fn), 0,
+ "region is %s partitioned but"
+ " does not contain %s partitioned code",
+ axes[ix], axes[ix]);
+ }
+
bool changed = targetm.goacc.validate_dims (fn, dims, level, used);
/* Default anything left to 1 or a partitioned default. */
+2018-10-05 Nathan Sidwell <nathan@acm.org>
+ Tom de Vries <tdevries@suse.de>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * c-c++-common/goacc/acc-icf.c: Update.
+ * c-c++-common/goacc/parallel-dims-1.c: Likewise.
+ * c-c++-common/goacc/parallel-reduction.c: Likewise.
+ * c-c++-common/goacc/pr70688.c: Likewise.
+ * c-c++-common/goacc/routine-1.c: Likewise.
+ * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
+ * gfortran.dg/goacc/parallel-tree.f95: Likewise.
+ * gfortran.dg/goacc/routine-4.f90: Likewise.
+ * gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise.
+ * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
+
2019-09-20 Chung-Lin Tang <cltang@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
#pragma acc routine gang
int
-routine1 (int n)
+routine1 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */
{
int i;
#pragma acc routine gang
int
-routine2 (int n)
+routine2 (int n) /* { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "" { xfail *-*-* } } */
{
int i;
void f(int i)
{
-#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
+#pragma acc kernels \
+ num_gangs(i) num_workers(i) vector_length(i)
;
-#pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
+#pragma acc parallel /* { dg-bogus "region is (gang|worker|vector) partitioned" "" { xfail *-*-* } } */ \
+ num_gangs(i) num_workers(i) vector_length(i)
;
}
#pragma acc data copy (dummy)
{
-#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) /* { dg-warning "gang partitioned" } */
{
int v = 5;
sum += 10 + v;
#pragma acc data copy (dummy)
{
-#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum)
+#pragma acc parallel num_gangs (10) copy (sum) reduction (+:sum) /* { dg-warning "region is gang partitioned" } */
{
int v = 5;
sum += 10 + v;
{
int i, s = 0;
-#pragma acc parallel num_gangs (10) copy (s) reduction (+:s)
+#pragma acc parallel num_gangs (10) copy (s) reduction (+:s) /* { dg-warning "region is gang partitioned" } */
for (i = 0; i < n; i++)
s += i+1;
-#pragma acc parallel num_gangs (10) reduction (+:s) copy (s)
+#pragma acc parallel num_gangs (10) reduction (+:s) copy (s) /* { dg-warning "region is gang partitioned" } */
for (i = 0; i < n; i++)
s += i+1;
/* Test valid use of clauses with routine. */
#pragma acc routine gang
-void gang (void)
+void gang (void) /* { dg-warning "partitioned" 3 } */
{
}
#pragma acc routine worker
-void worker (void)
+void worker (void) /* { dg-warning "partitioned" 2 } */
{
}
#pragma acc routine vector
-void vector (void)
+void vector (void) /* { dg-warning "partitioned" 1 } */
{
}
{
int i, j, k;
- #pragma acc parallel num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */
- ;
+ #pragma acc parallel loop gang num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */
+ for (i = 0; i < 1; i++)
+ ;
- #pragma acc parallel num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
- ;
+ #pragma acc parallel loop worker num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
+ for (j = 0; j < 1; j++)
+ ;
- #pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
- ;
+ #pragma acc parallel loop vector vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
+ for (k = 0; k < 1; k++)
+ ;
}
void acc_kernels()
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
+ ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" "" { target *-*-* } .-1 }
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } .-2 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } .-3 }
!$acc end parallel
end program test
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop gang worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop vector
do i = 1, N
a(i) = a(i) - a(i)
end do
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop seq
do i = 1, N
a(i) = a(i) - a(i)
end do
! parallelism with the OpenACC routine directive. The Fortran counterpart is
! c-c++-common/goacc/routine-level-of-parallelism-2.c
-subroutine g_1
+subroutine g_1 ! { dg-warning "region is gang partitioned but does not contain gang partitioned code" }
!$acc routine gang
+! { dg-bogus "region is worker partitioned but does not contain worker partitioned code" "worker partitioned" { xfail *-*-* } .-2 }
+! { dg-bogus "region is vector partitioned but does not contain vector partitioned code" "worker partitioned" { xfail *-*-* } .-3 }
end subroutine g_1
subroutine s_1_2a
!$ACC ROUTINE(s_2)
END SUBROUTINE s_2
- SUBROUTINE v_1
+ SUBROUTINE v_1 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" }
!$ACC ROUTINE VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_1) VECTOR
!$ACC ROUTINE VECTOR
END SUBROUTINE v_1
- SUBROUTINE v_2
+ SUBROUTINE v_2 ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" }
!$ACC ROUTINE(v_2) VECTOR
!$ACC ROUTINE VECTOR
!$ACC ROUTINE(v_2) VECTOR
implicit none
integer :: i, j, k
- !$acc parallel num_gangs(i) ! { dg-warning "is used uninitialized in this function" }
- !$acc end parallel
-
- !$acc parallel num_workers(j) ! { dg-warning "is used uninitialized in this function" }
- !$acc end parallel
-
- !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized in this function" }
- !$acc end parallel
+ !$acc parallel loop gang num_gangs(i) ! { dg-warning "is used uninitialized in this function" }
+ do i = 0, 1
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop worker num_workers(j) ! { dg-warning "is used uninitialized in this function" }
+ do j = 0, 1
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop vector vector_length(k) ! { dg-warning "is used uninitialized in this function" }
+ do k = 0, 1
+ end do
+ !$acc end parallel loop
end subroutine acc_parallel
subroutine acc_kernels
+2018-10-05 Nathan Sidwell <nathan@acm.org>
+ Tom de Vries <tdevries@suse.de>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Add -w.
+ * testsuite/libgomp.oacc-c-c++-common/loop-g-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/loop-warn-1.c: New.
+ * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update.
+ * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise.
+ * testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/loop-default-compile.c: New.
+
2018-10-22 James Norris <jnorris@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
Tom de Vries <tom@codesourcery.com>
arr[i] = 3;
#pragma acc parallel firstprivate(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 119 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 119 } */
{
#pragma acc loop gang
for (i = 0; i < 32; i++)
{
clear (ary, size);
-#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+#pragma acc parallel num_gangs (32) num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } } */
{
#pragma acc loop auto
for (int jx = 0; jx < size / 64; jx++)
--- /dev/null
+/* { dg-additional-options "-fopenacc-dim=16:16" } */
+/* This code uses nvptx inline assembly guarded with acc_on_device, which is
+ not optimized away at -O0, and then confuses the target assembler.
+ { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
+/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "8:8" } */
+
+#include "loop-default.h"
+
+int main ()
+{
+ /* Environment should be ignored. */
+ return test_1 (16, 16, 32);
+}
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
ary[ix] = -1;
#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 16 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
--- /dev/null
+
+/* Check warnings about suboptimal partitioning choices. */
+
+int main ()
+{
+ int ary[10];
+
+#pragma acc parallel copy(ary) num_gangs (1) /* { dg-warning "is not gang partitioned" } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_workers (1) /* { dg-warning "is not worker partitioned" } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_gangs (8) /* { dg-warning "is gang partitioned" } */
+ {
+ #pragma acc loop worker
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+#pragma acc parallel copy(ary) num_workers (8) /* { dg-warning "is worker partitioned" } */
+ {
+ #pragma acc loop gang
+ for (int i = 0; i < 10; i++)
+ ary[i] = i;
+ }
+
+ return 0;
+}
int n = 0;
#pragma acc parallel copy(n) \
num_gangs(1) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 288 } */
{
n++;
}
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 312 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 342 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(gangs) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 375 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(1024) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 409 } */
{
int j;
#pragma acc parallel copyout(fizz, buzz, fizzbuzz) \
num_gangs(NUM_GANGS) num_workers(1) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 448 } */
{
int j;
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 495 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(n, arr) \
num_gangs(8) num_workers(16) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 621 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copyin(arr_a) copyout(arr_b) \
num_gangs(num_gangs) num_workers(num_workers) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 674 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 892 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 892 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 917 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 917 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 945 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 945 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 973 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 973 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1006 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(arr) \
num_gangs(8) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1039 } */
{
int j;
#pragma acc loop gang
#pragma acc parallel copy(n, arr) copyout(ondev) \
num_gangs(ACTUAL_GANGS) num_workers(8) vector_length(32)
+ /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" "gang" { target *-*-* } 1090 } */
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 1090 } */
{
int j;
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (gangs_actual) \
+#pragma acc parallel copy (gangs_actual) /* { dg-warning "region contains gang partitioned code but is not gang partitioned" } */ \
num_gangs (GANGS) /* { dg-warning "'num_gangs' value must be positive" "" { target c++ } } */
{
/* We're actually executing with num_gangs (1). */
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (workers_actual) \
+#pragma acc parallel copy (workers_actual) /* { dg-warning "region contains worker partitioned code but is not worker partitioned" } */ \
num_workers (WORKERS) /* { dg-warning "'num_workers' value must be positive" "" { target c++ } } */
{
/* We're actually executing with num_workers (1). */
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
+#pragma acc parallel copy (vectors_actual) /* { dg-warning "region contains vector partitioned code but is not vector partitioned" } */ \
+ /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 157 } */ \
vector_length (VECTORS) /* { dg-warning "'vector_length' value must be positive" "" { target c++ } } */
{
/* We're actually executing with vector_length (1), just the GCC nvptx
int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
gangs_min = workers_min = vectors_min = INT_MAX;
gangs_max = workers_max = vectors_max = INT_MIN;
-#pragma acc parallel copy (gangs_actual) \
+#pragma acc parallel copy (gangs_actual) /* { dg-warning "region is gang partitioned but does not contain gang partitioned code" } */ \
reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) \
num_gangs (gangs)
{
;
#pragma acc parallel async (1) num_gangs (320)
- #pragma loop gang
+ #pragma acc loop gang
for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[N - ii - 1]);
arr[i] = 3;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 24 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 24 } */
{
int x;
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 299 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 299 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 326 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 355 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 384 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 417 } */
{
#pragma acc loop gang private(x)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 448 } */
{
#pragma acc loop gang private(pt)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 570 } */
{
int j;
arr[i] = 3;
#pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 887 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 887 } */
{
#pragma acc loop gang(static:1)
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 918 } */
{
#pragma acc loop gang
for (i = 0; i < 32; i++)
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 16 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 16 } */
{
#pragma acc loop gang reduction(+:res)
for (i = 0; i < 1024; i++)
res = hres = 1;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 32 } */
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 32 } */
{
#pragma acc loop gang reduction(*:res)
for (i = 0; i < 12; i++)
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 58 } */
{
#pragma acc loop gang vector reduction(+:res)
for (i = 0; i < 1024; i++)
arr[i] = i;
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 83 } */
{
#pragma acc loop gang worker reduction(+:res)
for (i = 0; i < 1024; i++)
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
private(res) copyout(out)
+ /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 246 } */
{
#pragma acc loop gang
for (j = 0; j < 32; j++)
#pragma acc parallel num_gangs(32) num_workers(32) vector_length(32) \
private(res) copyout(out)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 323 } */
{
#pragma acc loop gang
for (j = 0; j < 32; j++)
#pragma acc routine gang
void __attribute__ ((noinline)) gang (int ary[N])
+/* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 8 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 8 } */
{
#pragma acc loop gang
for (unsigned ix = 0; ix < N; ix++)
#pragma acc routine worker
void __attribute__ ((noinline)) worker (int ary[N])
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 8 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
RES2 = 0
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 + 5
!$ACC ATOMIC
RES2 = 1
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 * 5
!$ACC ATOMIC
RES2 = 0
!$ACC PARALLEL NUM_GANGS(256) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(+:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 + 5
!$ACC ATOMIC
RES2 = 1
!$ACC PARALLEL NUM_GANGS(8) NUM_WORKERS(32) VECTOR_LENGTH(32)
-!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1)
+!$ACC& REDUCTION(*:RES1) COPY(RES1, RES2) ASYNC(1) ! { dg-warning "region is (gang|worker|vector) partitioned" }
res1 = res1 * 5
!$ACC ATOMIC
a = 1
- !$acc parallel num_gangs(1) num_workers(2)
+ !$acc parallel num_gangs(1) num_workers(2) ! { dg-warning "region is worker partitioned" }
if (any(a(1:3,1:3,1:3).ne.1)) STOP 1
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 15 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 15 }
!$acc loop gang private(x)
do i = 1, 32
x = i * 2;
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 41 }
!$acc loop gang private(x)
do i = 0, 31
x = i * 2;
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 70 }
!$acc loop gang private(x)
do i = 0, 31
x = i * 2;
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 104 }
!$acc loop gang private(pt)
do i = 0, 31
pt%x = i
end do
!$acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 215 }
!$acc loop gang private(x)
do i = 0, 31
!$acc loop worker private(x)
end do
!$acc parallel private(x) copy(arr) num_gangs(n) num_workers(8) vector_length(32)
+ ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 515 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 515 }
!$acc loop gang(static:1)
do i = 1, n
x = i * 2;
-
! { dg-do run }
! { dg-additional-options "-cpp" }
integer, intent (inout) :: a(N)
integer :: i
- !$acc loop gang
+ !$acc loop gang worker vector
do i = 1, N
a(i) = a(i) - i
end do