From patchwork Thu Nov  3 16:13:02 2016
Content-Type: text/plain; charset="utf-8"
MIME-Version: 1.0
Content-Transfer-Encoding: 7bit
X-Patchwork-Submitter: Cesar Philippidis <cesar@codesourcery.com>
X-Patchwork-Id: 80695
Delivered-To: patch@linaro.org
Received: by 10.140.97.247 with SMTP id m110csp714751qge;
 Thu, 3 Nov 2016 09:13:41 -0700 (PDT)
X-Received: by 10.99.170.13 with SMTP id e13mr14882055pgf.68.1478189621552; 
 Thu, 03 Nov 2016 09:13:41 -0700 (PDT)
Return-Path: <gcc-patches-return-440361-patch=linaro.org@gcc.gnu.org>
Received: from sourceware.org (server1.sourceware.org. [209.132.180.131])
 by mx.google.com with ESMTPS id
 t129si10422647pgc.171.2016.11.03.09.13.41 for <patch@linaro.org>
 (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128);
 Thu, 03 Nov 2016 09:13:41 -0700 (PDT)
Received-SPF: pass (google.com: domain of
 gcc-patches-return-440361-patch=linaro.org@gcc.gnu.org
 designates 209.132.180.131 as permitted sender)
 client-ip=209.132.180.131; 
Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org;
 spf=pass (google.com: domain of
 gcc-patches-return-440361-patch=linaro.org@gcc.gnu.org
 designates 209.132.180.131 as permitted sender)
 smtp.mailfrom=gcc-patches-return-440361-patch=linaro.org@gcc.gnu.org
DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id
 :list-unsubscribe:list-archive:list-post:list-help:sender:from
 :subject:to:message-id:date:mime-version:content-type; q=dns; s=
 default; b=UvjeuK/wcqVftEZPGguWykObuca8bg6bQwVuWsCmsow3B1khd6ttO
 N4w8Ji7XF42XjviFgSIypvqGfqlUHKIXuaXtykYJegqOGpM/+wvkFV/4i2a0ysGi
 APiUA2SVy9g/O+qwPhazs+bFtD8uvxR550rgprD4r+HdbxG48dcBxg=
DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id
 :list-unsubscribe:list-archive:list-post:list-help:sender:from
 :subject:to:message-id:date:mime-version:content-type; s=
 default; bh=1kee96aKiIfEjBTrhmvV9OAA8O8=; b=bPoii8Vxe7JUuJ4xq/Yq
 jMjj0OD4teLt5KMFp2ZYbRQmdAeJbLfm67ApNt3RkSIHyyt3ljIdyn/o8lDheiRU
 sjRR+X/mcC143TKSNvJ7HL2piQFEuAWQGDLOWc352s7g58AVpEx1g3Gr0LJHN5I7
 /1r9urI768n9h3bw2fKhVwQ=
Received: (qmail 63620 invoked by alias); 3 Nov 2016 16:13:23 -0000
Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm
Precedence: bulk
List-Id: <gcc-patches.gcc.gnu.org>
List-Unsubscribe: <mailto:gcc-patches-unsubscribe-patch=linaro.org@gcc.gnu.org>
List-Archive: <http://gcc.gnu.org/ml/gcc-patches/>
List-Post: <mailto:gcc-patches@gcc.gnu.org>
List-Help: <mailto:gcc-patches-help@gcc.gnu.org>
Sender: gcc-patches-owner@gcc.gnu.org
Delivered-To: mailing list gcc-patches@gcc.gnu.org
Received: (qmail 63599 invoked by uid 89); 3 Nov 2016 16:13:22 -0000
Authentication-Results: sourceware.org; auth=none
X-Virus-Found: No
X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00,
 RCVD_IN_DNSWL_NONE, SPF_PASS,
 URIBL_RED autolearn=ham version=3.3.2 spammy=3707, 1213, 7548,
 3247
X-HELO: relay1.mentorg.com
Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131)
 by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with
 ESMTP; Thu, 03 Nov 2016 16:13:11 +0000
Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204])	by
 relay1.mentorg.com with esmtp id 1c2KdY-000155-M8 from
 Cesar_Philippidis@mentor.com ; Thu, 03 Nov 2016 09:13:08 -0700
Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com
 (147.34.90.204) with Microsoft SMTP Server (TLS) id
 15.0.1210.3; Thu, 3 Nov 2016 09:13:03 -0700
From: Cesar Philippidis <cesar@codesourcery.com>
Subject: [openacc] add warnings for unused parallelism
To: "gcc-patches@gcc.gnu.org" <gcc-patches@gcc.gnu.org>,
 Jakub Jelinek	<jakub@redhat.com>
Message-ID: <3f6a3cec-132b-aadc-886c-be20888f56e8@codesourcery.com>
Date: Thu, 3 Nov 2016 09:13:02 -0700
User-Agent: Mozilla/5.0 (X11; Linux x86_64;
 rv:45.0) Gecko/20100101 Thunderbird/45.3.0
MIME-Version: 1.0
X-ClientProxiedBy: svr-orw-mbx-02.mgc.mentorg.com (147.34.90.202) To
 SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204)

OpenACC permits the user to request more gang, worker and vector level
parallelism than what the compiler can utilize. For instance, if the
user writes worker routine without including a worker-partitioned loop,
the compiler will not generate worker-partitioned code for that function.

The intent behind this patch is to warn the user of any potentially
unutilized parallelism as a debugging aid. Users often find it
disconcerting when their code doesn't speed up despite explicitly
setting num_gangs, num_workers and vector_length. This patch at least
warns them not to expect parallelism across a specific axis.

Is this patch OK for trunk? This patch was originally posted by Nathan
here <https://gcc.gnu.org/ml/gcc-patches/2016-01/msg00210.html>. Most of
the changes in that patch are already in trunk.

Cesar

2016-11-03  Cesar Philippidis  <cesar@codesourcery.com>
	    Nathan Sidwell <nathan@acm.org>

	gcc/
	* omp-low.c (oacc_validate_dims): Emit warnings about strange
	partitioning choices.

	gcc/testsuite/
	* c-c++-common/goacc/pr70688.c (parallel_reduction): Adjust expected
	warnings.
	* c-c++-common/goacc/routine-1.c: Likewise.
	* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
	* gcc.dg/goacc/loop-processing-1.c (void vector_1): Likewise.
	* gfortran.dg/goacc/parallel-tree.f95: Likewise.
	* gfortran.dg/goacc/routine-4.f90: Likewise.
	* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
	* gfortran.dg/goacc/vector_length.f90: Likewise.

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/crash-1.c: Adjust to account
	for insufficient parallelism warnings.
	* testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: Likewise.
	* 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-w-1.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/mode-transitions.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/reduction-initial-1.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/private-variables.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise.

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e5b9e4c..cbf4f3e 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -18882,6 +18882,36 @@ oacc_validate_dims (tree fn, tree attrs, int *dims, int level, unsigned used)
       pos = TREE_CHAIN (pos);
     }
 
+  bool check = true;
+#ifdef ACCEL_COMPILER
+  /* When device_type is implemented, we should also check on the
+     target, if device_type has been used to affect the partitioning
+     and/or dimensions.  */
+  check = false;
+#endif
+  if (!is_kernel && check)
+    {
+      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 partitoned 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);
 
   /* Default anything left to 1 or a partitioned default.  */
diff --git a/gcc/testsuite/c-c++-common/goacc/pr70688.c b/gcc/testsuite/c-c++-common/goacc/pr70688.c
index 5a23665..37c3885 100644
--- a/gcc/testsuite/c-c++-common/goacc/pr70688.c
+++ b/gcc/testsuite/c-c++-common/goacc/pr70688.c
@@ -1,3 +1,5 @@
+/* { dg-compile } */
+
 const int n = 100;
 
 int
@@ -21,7 +23,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 +38,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;
 
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
index a5e0d69..95e9729 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -1,16 +1,17 @@
+/* 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 } */
 {
 }
 
diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
index 0a006e3..8bfb47c 100644
--- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -8,12 +8,17 @@ main (void)
 {
   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++)
+    ;
+
+  return 0;
 }
diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
index 619576a..c00dd9e 100644
--- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
+++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
@@ -5,7 +5,7 @@ extern int place ();
 
 void vector_1 (int *ary, int size)
 {
-#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+#pragma acc parallel num_workers (32) vector_length (32) copy(ary[0:size]) firstprivate (size) /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" } */
   {
 #pragma acc loop gang
     for (int jx = 0; jx < 1; jx++)
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 5b2e01d..763436e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -12,9 +12,13 @@ program test
   !$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 *-*-* } 14 }
+  ! { dg-warning "region is worker partitioned but does not contain worker partitioned code" "" { target *-*-* } 14 }
+  ! { dg-warning "region is vector partitioned but does not contain vector partitioned code" "" { target *-*-* } 14 }
   !$acc end parallel
 
 end program test
+
 ! { dg-final { scan-tree-dump-times "pragma acc parallel" 1 "original" } } 
 
 ! { dg-final { scan-tree-dump-times "if" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
index 6714c7b..3fb60e71 100644
--- a/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-4.f90
@@ -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
diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
index b87d26f..dd3723a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
@@ -5,13 +5,18 @@ program test
   implicit none
   integer :: i, j, k
 
-  !$acc parallel num_gangs(i) ! { 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 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 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 program test
diff --git a/gcc/testsuite/gfortran.dg/goacc/vector_length.f90 b/gcc/testsuite/gfortran.dg/goacc/vector_length.f90
index ddab9cf..f335f42 100644
--- a/gcc/testsuite/gfortran.dg/goacc/vector_length.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/vector_length.f90
@@ -3,7 +3,8 @@ program t
   integer, parameter :: n = 100
   integer a(n), i
 
-  !$acc parallel loop num_gangs(100) num_workers(1) vector_length(32)
+  !$acc parallel loop num_gangs(100) num_workers(1) vector_length(32) &
+  !$acc& gang vector
   do i = 1, n
      a(i) = i
   enddo
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
index dcf1485..100e4b1 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/crash-1.c
@@ -3,7 +3,7 @@
 /* For -O0, ICEd in nvptx backend due to unexpected frame size.  */
 #pragma acc routine worker
 void
-worker_matmul (int *c, int i)
+worker_matmul (int *c, int i) /* { dg-warning "region is vector partitioned" } */
 {
   int j;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
index 689a443..14bc3af 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c
@@ -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++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
index 3ca9388..0df2969 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c
@@ -103,7 +103,7 @@ int vector_1 (int *ary, int size)
 {
   clear (ary, size);
   
-#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size) /* { dg-warning "region is worker partitioned" } */
   {
 #pragma acc loop gang
     for (int jx = 0; jx < 1; jx++)
@@ -153,7 +153,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++)
@@ -187,7 +187,7 @@ int gang_3 (int *ary, int size)
 {
   clear (ary, size);
   
-#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size)
+#pragma acc parallel num_workers (32) vector_length(32) copy(ary[0:size]) firstprivate (size) /* { dg-warning "region is worker partitioned" } */
   {
 #pragma acc loop auto
     for (int jx = 0; jx <  size  / 64; jx++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
index 7bff6cd..c7b43b7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-1.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
 /* 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" } { "" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
index 92b82a0..6c4c0d2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-g-2.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
 /* 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" } { "" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
index d241d41..e5c302e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-g-1.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
 /* 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" } { "" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
index 19021d9..efda662 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
 /* 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" } { "" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
index f0c9d81..174a3ff 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c
@@ -1,3 +1,4 @@
+/* { dg-additional-options "-w" } */
 /* 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" } { "" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
index 33b6eae..30e8e78 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-w-1.c
@@ -16,6 +16,7 @@ int main ()
     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 *-*-* } 18 } */
   {
 #pragma acc loop worker
     for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
index 2394ac8..16a1e0d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/mode-transitions.c
@@ -163,6 +163,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 *-*-* } 164 } */
   {
     n++;
   }
@@ -186,6 +187,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 *-*-* } 188 } */
       {
 	int j;
 	#pragma acc loop gang
@@ -215,6 +217,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 *-*-* } 218 } */
       {
 	int j;
 	#pragma acc loop gang
@@ -247,6 +250,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 *-*-* } 251 } */
       {
 	int j;
 	#pragma acc loop gang
@@ -280,6 +284,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 *-*-* } 285 } */
   {
     int j;
 
@@ -318,6 +323,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 *-*-* } 324 } */
   {
     int j;
     
@@ -364,6 +370,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 *-*-* } 371 } */
   {
     int j;
     #pragma acc loop gang
@@ -395,6 +402,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 *-*-* } 403 } */
   {
     int j;
     #pragma acc loop gang
@@ -447,6 +455,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 *-*-* } 456 } */
 	{
 	  int j;
 	  #pragma acc loop gang
@@ -664,6 +673,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 *-*-* } 674 } */
+  /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 674 } */
   {
     int j;
     #pragma acc loop gang
@@ -687,6 +698,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 *-*-* } 699 } */
+  /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 699 } */
   {
     int j;
     #pragma acc loop gang
@@ -713,6 +726,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 *-*-* } 727 } */
+  /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 727 } */
   {
     int j;
     #pragma acc loop gang
@@ -739,6 +754,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 *-*-* } 755 } */
+  /* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 755 } */
   {
     int j;
     #pragma acc loop gang
@@ -770,6 +787,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 *-*-* } 788 } */
   {
     int j;
     #pragma acc loop gang
@@ -805,6 +823,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 *-*-* } 824 } */
+  /* { dg-warning "region is worker partitioned but does not contain worker partitioned code" "worker" { target *-*-* } 824 } */
   {
     int j;
 
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
index 53f03d1..f0c3447 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-variables.c
@@ -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++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
index cc3cd07..61b548e 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-7.c
@@ -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++)
@@ -236,6 +242,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 *-*-* } 243 } */
   {
     #pragma acc loop gang
     for (j = 0; j < 32; j++)
@@ -312,6 +319,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 *-*-* } 320 } */
   {
     #pragma acc loop gang
     for (j = 0; j < 32; j++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c
index 9991db0..81214e8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-initial-1.c
@@ -13,7 +13,7 @@ main(void)
 #pragma acc parallel vector_length(N) copy(s)
   {
     int i;
-#pragma acc loop reduction(+:s)
+#pragma acc loop reduction(+:s) vector
     for (i = 0; i < N; ++i)
       s += a;
   }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
index 9d14c3b..2ef5a55 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-g-1.c
@@ -8,6 +8,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 *-*-* } 10 } */
+/* { dg-warning "region is vector partitioned but does not contain vector partitioned code" "vector" { target *-*-* } 10 } */
 {
 #pragma acc loop gang
     for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
index 80cd462..0b03a01 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/routine-w-1.c
@@ -8,6 +8,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 *-*-* } 10 } */
 {
 #pragma acc loop worker
   for (unsigned ix = 0; ix < N; ix++)
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90 b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
index 3c1940b..2535eb8 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/private-variables.f90
@@ -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;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
index 200188e..ef2ff04 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/routine-7.f90
@@ -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