@@ -928,6 +928,34 @@ oacc_validate_dims (tree fn, tree attrs, int *dims, int level, unsigned used)
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. */
@@ -4,7 +4,7 @@
#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;
@@ -17,7 +17,7 @@ routine1 (int n)
#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;
@@ -15,6 +15,8 @@ extern unsigned int *__restrict c;
void SERIAL ()
{
#pragma acc serial loop copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message "optimized: assigned OpenACC gang vector loop parallelism" } */
+/* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-1 } */
+/* { dg-warning "region contains vector partitioned code but is not vector partitioned" "" { target *-*-* } .-2 } */
for (unsigned int i = 0; i < N; i++)
c[i] = a[i] + b[i];
}
@@ -121,6 +121,7 @@ main ()
/*TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
{ dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part } */
/* { dg-optimized "assigned OpenACC gang worker vector loop parallelism" "" { target *-*-* } l_part$c_part } */
+ /* { dg-warning "region contains gang partitioned code but is not gang partitioned" "" { target *-*-* } .-6 } */
#pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */
/* { dg-message "note: parallelized loop nest in OpenACC 'kernels' region" "" { target *-*-* } l_loop_j$c_loop_j } */
@@ -3,9 +3,11 @@
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)
;
}
@@ -6,7 +6,7 @@ main ()
#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;
@@ -21,7 +21,7 @@ parallel_reduction ()
#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;
@@ -36,11 +36,11 @@ main ()
{
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;
@@ -1,16 +1,16 @@
#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 } */
{
}
@@ -4,6 +4,9 @@
#pragma acc routine gang
void g_1 (void)
+/* { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } */
+/* { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-2 } */
+/* { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-3 } */
{
}
#pragma acc routine (g_1) gang
@@ -5,12 +5,15 @@ void acc_parallel()
int i, j, k;
#pragma acc parallel num_gangs(i) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 } */
;
#pragma acc parallel num_workers(j) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-1 } */
;
#pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized" } */
+ /* { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-1 } */
;
}
@@ -15,6 +15,8 @@ program main
call setup(a, b)
!$acc serial loop copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) ! { dg-message "optimized: assigned OpenACC gang vector loop parallelism" }
+ ! { dg-warning "not gang partitioned" "" { target *-*-* } .-1 }
+ ! { dg-warning "not vector partitioned" "" { target *-*-* } .-2 }
do i = 0, n - 1
c(i) = a(i) + b(i)
end do
@@ -119,6 +119,7 @@ program main
!$acc end kernels
!$acc kernels
+ ! { dg-warning "not gang partitioned" "" { target *-*-* } .-1 }
y = f_g (a(5)) ! { dg-line l_part[incr c_part] }
!TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
! { dg-message "note: beginning 'gang-single' part in OpenACC 'kernels' region" "" { target *-*-* } l_part$c_part }
@@ -12,6 +12,9 @@ program test
!$acc no_create(n) &
!$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
@@ -123,6 +123,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop gang worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -133,6 +134,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop worker vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -143,6 +145,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop vector
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -153,6 +156,7 @@ contains
integer, intent (inout) :: a(N)
integer :: i
+ !$acc loop seq
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -2,8 +2,10 @@
! with the OpenACC routine directive. The C/C++ 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
@@ -50,6 +50,7 @@ contains
end do
end subroutine w_1
+ ! { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .+1 }
subroutine g_1
implicit none
!$acc routine gang
@@ -12,14 +12,14 @@
!$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
@@ -5,12 +5,15 @@ subroutine acc_parallel
integer :: i, j, k
!$acc parallel num_gangs(i) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "does not contain gang partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
!$acc parallel num_workers(j) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "does not contain worker partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
!$acc parallel vector_length(k) ! { dg-warning "is used uninitialized" }
+ ! { dg-warning "does not contain vector partitioned code" "" { target *-*-* } .-1 }
!$acc end parallel
end subroutine acc_parallel
@@ -117,6 +117,8 @@ void t4 ()
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++)
@@ -150,7 +150,7 @@ int gang_1 (int *ary, int size)
{
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++)
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
#include <stdio.h>
#include <openacc.h>
#include <gomp-constants.h>
@@ -16,6 +16,7 @@ int main ()
#pragma acc parallel num_workers(32) vector_length(32) copy(ary) copy(ondev) \
copyout(workersize)
+ /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 17 } */
{
#pragma acc loop worker
for (unsigned ix = 0; ix < N; ix++)
new file mode 100644
@@ -0,0 +1,37 @@
+
+/* 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;
+}
@@ -287,6 +287,7 @@ void t7()
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++;
}
@@ -310,6 +311,7 @@ void t8()
#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
@@ -339,6 +341,7 @@ void t9()
#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
@@ -371,6 +374,7 @@ void t10()
#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
@@ -404,6 +408,7 @@ void t11()
#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;
@@ -442,6 +447,7 @@ void t12()
#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;
@@ -488,6 +494,7 @@ void t13()
#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
@@ -613,6 +620,7 @@ void t16()
#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
@@ -665,6 +673,7 @@ void t17()
#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
@@ -882,6 +891,8 @@ void t21()
#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
@@ -905,6 +916,8 @@ void t22()
#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
@@ -931,6 +944,8 @@ void t23()
#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
@@ -957,6 +972,8 @@ void t24()
#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
@@ -988,6 +1005,7 @@ void t25()
#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
@@ -1020,6 +1038,7 @@ void t26()
#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
@@ -1070,6 +1089,8 @@ void t27()
#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;
@@ -103,7 +103,7 @@ int main ()
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). */
@@ -132,7 +132,7 @@ int main ()
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). */
@@ -161,7 +161,8 @@ int main ()
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 } 164 } */ \
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
@@ -205,7 +206,7 @@ int main ()
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)
{
@@ -617,6 +618,9 @@ int main ()
gangs_max = workers_max = vectors_max = INT_MIN;
#pragma acc serial copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max)
+/* { dg-warning "not gang partitioned" "" { target *-*-* } 619 } */
+/* { dg-warning "not worker partitioned" "" { target *-*-* } 619 } */
+/* { dg-warning "not vector partitioned" "" { target *-*-* } 619 } */
{
if (acc_on_device (acc_device_nvidia))
{
@@ -1,6 +1,6 @@
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
-/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach -w" } */
int a;
#pragma acc declare create(a)
@@ -22,6 +22,8 @@ void local_g_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;
@@ -295,6 +297,8 @@ void loop_g_1()
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++)
@@ -320,6 +324,7 @@ void loop_g_2()
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++)
@@ -348,6 +353,7 @@ void loop_g_3()
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++)
@@ -376,6 +382,7 @@ void loop_g_4()
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++)
@@ -408,6 +415,7 @@ void loop_g_5()
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++)
@@ -438,6 +446,7 @@ void loop_g_6()
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++)
@@ -559,6 +568,7 @@ void loop_w_1()
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;
@@ -875,6 +885,8 @@ void parallel_g_1()
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++)
@@ -904,6 +916,7 @@ void parallel_g_2()
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++)
@@ -14,6 +14,8 @@ void g_np_1()
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++)
@@ -28,6 +30,8 @@ void g_np_1()
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++)
@@ -52,6 +56,7 @@ void gv_np_1()
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++)
@@ -76,6 +81,7 @@ void gw_np_1()
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++)
@@ -239,6 +245,7 @@ void v_p_1()
#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++)
@@ -315,6 +322,7 @@ void w_p_1()
#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++)
@@ -6,6 +6,8 @@
#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++)
@@ -6,6 +6,7 @@
#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++)
@@ -42,7 +42,8 @@ int DoWorkVec (int nw)
ary[ix][jx] = 0xdeadbeef;
printf ("spawning %d ...", nw); fflush (stdout);
-
+
+ /* { dg-warning "region contains vector partitioned code but is not vector partitioned" "vector" { target openacc_radeon_accel_selected } 47 } */
#pragma acc parallel num_workers(nw) vector_length (NUM_VECTORS) copy (ary)
{
WorkVec ((int *)ary, WIDTH, HEIGHT, nw, NUM_VECTORS);
@@ -30,6 +30,8 @@ contains
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 *-*-* } 32 }
+ ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 32 }
!$acc loop gang private(x)
do i = 1, 32
x = i * 2;
@@ -55,6 +57,7 @@ contains
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 *-*-* } 59 }
!$acc loop gang private(pt)
do i = 0, 31
pt%x = i
@@ -14,7 +14,7 @@
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
@@ -36,7 +36,7 @@
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
@@ -14,7 +14,7 @@
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
@@ -36,7 +36,7 @@
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
@@ -84,6 +84,9 @@ program main
vectors_max = -huge(gangs_max) - 1 ! INT_MIN
!$acc serial copy (vectors_actual) &
!$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
+ ! { dg-warning "not gang partitioned" "" { target *-*-* } 86 }
+ ! { dg-warning "not worker partitioned" "" { target *-*-* } 86 }
+ ! { dg-warning "not vector partitioned" "" { target *-*-* } 86 }
if (acc_on_device (acc_device_nvidia)) then
! The GCC nvptx back end enforces vector_length (32).
! It's unclear if that's actually permissible here;
@@ -5,7 +5,7 @@ program foo
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
@@ -13,6 +13,8 @@ subroutine t1()
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;
@@ -37,6 +39,7 @@ subroutine t2()
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;
@@ -65,6 +68,7 @@ subroutine t3()
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;
@@ -98,6 +102,7 @@ subroutine t4()
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
@@ -208,6 +213,7 @@ subroutine t7()
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)
@@ -507,6 +513,8 @@ subroutine t14()
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;
@@ -100,7 +100,7 @@ subroutine gang (a)
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