+2015-11-04 Nathan Sidwell <nathan@codesourcery.com>
+
+ * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Add checking.
+
2015-11-04 Nathan Sidwell <nathan@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
{
bool changed = false;
- /* TODO: Leave dimensions unaltered. Reductions need
- porting before filtering dimensions makes sense. */
+ /* The vector size must be 32, unless this is a SEQ routine. */
+ if (fn_level <= GOMP_DIM_VECTOR
+ && dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH)
+ {
+ if (dims[GOMP_DIM_VECTOR] >= 0 && fn_level < 0)
+ warning_at (DECL_SOURCE_LOCATION (decl), 0,
+ dims[GOMP_DIM_VECTOR]
+ ? "using vector_length (%d), ignoring %d"
+ : "using vector_length (%d), ignoring runtime setting",
+ PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]);
+ dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH;
+ changed = true;
+ }
+
+ /* Check the num workers is not too large. */
+ if (dims[GOMP_DIM_WORKER] > PTX_WORKER_LENGTH)
+ {
+ warning_at (DECL_SOURCE_LOCATION (decl), 0,
+ "using num_workers (%d), ignoring %d",
+ PTX_WORKER_LENGTH, dims[GOMP_DIM_WORKER]);
+ dims[GOMP_DIM_WORKER] = PTX_WORKER_LENGTH;
+ changed = true;
+ }
return changed;
}
+2015-11-04 Nathan Sidwell <nathan@codesourcery.com>
+
+ * testsuite/libgomp.oacc-fortran/reduction-1.f90: Fix dimensions
+ and reduction copy.
+ * testsuite/libgomp.oacc-fortran/reduction-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-3.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-4.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New.
+
2015-11-04 Nathan Sidwell <nathan@codesourcery.com>
* libgomp.oacc-c-c++-common/loop-red-g-1.c: New.
int i, j, k, l = 0, f = 0, x = 0;
int m1 = 4, m2 = -5, m3 = 17;
- #pragma acc parallel
+#pragma acc parallel copy(l)
#pragma acc loop collapse(3) reduction(+:l)
for (i = -2; i < m1; i++)
for (j = m2; j < -2; j++)
#else
# define GANGS 256
#endif
- #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \
- copy(res2)
+ #pragma acc parallel num_gangs(GANGS) copy(res2)
{
#pragma acc atomic
res2 += 5;
#else
# define GANGS 8
#endif
- #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \
- copy(res2)
+ #pragma acc parallel num_gangs(GANGS) copy(res2)
{
#pragma acc atomic
res2 *= 5;
#else
# define GANGS 256
#endif
- #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \
- copy(res2) async(1)
+ #pragma acc parallel num_gangs(GANGS) copy(res2) async(1)
{
#pragma acc atomic
res2 += 5;
#else
# define GANGS 8
#endif
- #pragma acc parallel num_gangs(GANGS) num_workers(1) vector_length(1) \
- copy(res2) async(1)
+ #pragma acc parallel num_gangs(GANGS) copy(res2) async(1)
{
#pragma acc atomic
res2 *= 5;
--- /dev/null
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+/* Worker and vector size checks. Picked an outrageously large
+ value. */
+
+int main ()
+{
+#pragma acc parallel num_workers (2<<20) /* { dg-error "using num_workers" } */
+ {
+ }
+
+#pragma acc parallel vector_length (2<<20) /* { dg-error "using vector_length" } */
+ {
+ }
+
+ return 0;
+}
{ \
type res, vres; \
res = (init); \
-DO_PRAGMA (acc parallel vector_length (vl))\
+ DO_PRAGMA (acc parallel vector_length (vl) copy(res)) \
DO_PRAGMA (acc loop reduction (op:res))\
for (i = 0; i < n; i++) \
res = res op (b); \
{ \
type res, vres; \
res = (init); \
-DO_PRAGMA (acc parallel vector_length (vl))\
+DO_PRAGMA (acc parallel vector_length (vl) copy(res))\
DO_PRAGMA (acc loop reduction (op:res))\
for (i = 0; i < n; i++) \
res = op (res, (b)); \
vresult = 0;
/* '+' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(result)
#pragma acc loop reduction (+:result)
for (i = 0; i < n; i++)
result += array[i];
vresult = 0;
/* '*' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(result)
#pragma acc loop reduction (*:result)
for (i = 0; i < n; i++)
result *= array[i];
lvresult = false;
/* '&&' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(lresult)
#pragma acc loop reduction (&&:lresult)
for (i = 0; i < n; i++)
lresult = lresult && (result > array[i]);
lvresult = false;
/* '||' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(lresult)
#pragma acc loop reduction (||:lresult)
for (i = 0; i < n; i++)
lresult = lresult || (result > array[i]);
vresult = 0;
/* '+' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(result)
#pragma acc loop reduction (+:result)
for (i = 0; i < n; i++)
result += array[i];
vresult = 0;
/* '*' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(result)
#pragma acc loop reduction (*:result)
for (i = 0; i < n; i++)
result *= array[i];
lvresult = false;
/* '&&' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(lresult)
#pragma acc loop reduction (&&:lresult)
for (i = 0; i < n; i++)
lresult = lresult && (result > array[i]);
lvresult = false;
/* '||' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(lresult)
#pragma acc loop reduction (||:lresult)
for (i = 0; i < n; i++)
lresult = lresult || (result > array[i]);
vresult = 0;
/* '+' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(result)
#pragma acc loop reduction (+:result)
for (i = 0; i < n; i++)
result += array[i];
lvresult = false;
/* '&&' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(lresult)
#pragma acc loop reduction (&&:lresult)
for (i = 0; i < n; i++)
lresult = lresult && (creal(result) > creal(array[i]));
lvresult = false;
/* '||' reductions. */
-#pragma acc parallel vector_length (vl)
+#pragma acc parallel vector_length (vl) copy(lresult)
#pragma acc loop reduction (||:lresult)
for (i = 0; i < n; i++)
lresult = lresult || (creal(result) > creal(array[i]));
int n = 100;
int i;
-#pragma acc parallel vector_length (32)
+#pragma acc parallel vector_length (32) copy(s1,s2)
#pragma acc loop reduction (+:s1, s2)
for (i = 0; i < n; i++)
{
main(void)
{
#define I 5
-#define N 11
+#define N 32
#define A 8
int a = A;
int s = I;
-#pragma acc parallel vector_length(N)
+#pragma acc parallel vector_length(N) copy(s)
{
int i;
#pragma acc loop reduction(+:s)
program reduction_1
implicit none
- integer, parameter :: n = 10, vl = 2
+ integer, parameter :: n = 10, vl = 32
integer :: i, vresult, result
logical :: lresult, lvresult
integer, dimension (n) :: array
! '+' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(+:result)
do i = 1, n
result = result + array(i)
! '*' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(*:result)
do i = 1, n
result = result * array(i)
! 'max' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(max:result)
do i = 1, n
result = max (result, array(i))
! 'min' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(min:result)
do i = 1, n
result = min (result, array(i))
! 'iand' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(iand:result)
do i = 1, n
result = iand (result, array(i))
! 'ior' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(ior:result)
do i = 1, n
result = ior (result, array(i))
! 'ieor' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(ieor:result)
do i = 1, n
result = ieor (result, array(i))
! '.and.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.and.:lresult)
do i = 1, n
lresult = lresult .and. (array(i) .ge. 5)
! '.or.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.or.:lresult)
do i = 1, n
lresult = lresult .or. (array(i) .ge. 5)
! '.eqv.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.eqv.:lresult)
do i = 1, n
lresult = lresult .eqv. (array(i) .ge. 5)
! '.neqv.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.neqv.:lresult)
do i = 1, n
lresult = lresult .neqv. (array(i) .ge. 5)
program reduction_2
implicit none
- integer, parameter :: n = 10, vl = 2
+ integer, parameter :: n = 10, vl = 32
integer :: i
real, parameter :: e = .001
real :: vresult, result
! '+' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(+:result)
do i = 1, n
result = result + array(i)
! '*' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(*:result)
do i = 1, n
result = result * array(i)
! 'max' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(max:result)
do i = 1, n
result = max (result, array(i))
! 'min' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(min:result)
do i = 1, n
result = min (result, array(i))
! '.and.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.and.:lresult)
do i = 1, n
lresult = lresult .and. (array(i) .ge. 5)
! '.or.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.or.:lresult)
do i = 1, n
lresult = lresult .or. (array(i) .ge. 5)
! '.eqv.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.eqv.:lresult)
do i = 1, n
lresult = lresult .eqv. (array(i) .ge. 5)
! '.neqv.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.neqv.:lresult)
do i = 1, n
lresult = lresult .neqv. (array(i) .ge. 5)
program reduction_3
implicit none
- integer, parameter :: n = 10, vl = 2
+ integer, parameter :: n = 10, vl = 32
integer :: i
double precision, parameter :: e = .001
double precision :: vresult, result
! '+' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(+:result)
do i = 1, n
result = result + array(i)
! '*' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(*:result)
do i = 1, n
result = result * array(i)
! 'max' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(max:result)
do i = 1, n
result = max (result, array(i))
! 'min' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(min:result)
do i = 1, n
result = min (result, array(i))
! '.and.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.and.:lresult)
do i = 1, n
lresult = lresult .and. (array(i) .ge. 5)
! '.or.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.or.:lresult)
do i = 1, n
lresult = lresult .or. (array(i) .ge. 5)
! '.eqv.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.eqv.:lresult)
do i = 1, n
lresult = lresult .eqv. (array(i) .ge. 5)
! '.neqv.' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(lresult)
!$acc loop reduction(.neqv.:lresult)
do i = 1, n
lresult = lresult .neqv. (array(i) .ge. 5)
! '+' reductions
- !$acc parallel vector_length(vl) num_gangs(1)
+ !$acc parallel vector_length(vl) num_gangs(1) copy(result)
!$acc loop reduction(+:result)
do i = 1, n
result = result + array(i)
vs1 = 0
vs2 = 0
- !$acc parallel vector_length (32)
+ !$acc parallel vector_length (32) copy(s1, s2)
!$acc loop reduction(+:s1, s2)
do i = 1, n
s1 = s1 + 1