[gomp4] filter out unrecognized device_type targets

[Cesar Philippidis]

This patch addresses a bug where the c and c++ front ends would ICE in
the event that an unrecognized target is specified as a device_type,
e.g. device_type (foo). The device_type clause is still highly specific
to nvidia targets, so it's likely something to be be merged in gcc 7.1
rather than 6.1. In the meantime, I've applied this patch to
gomp-4_0-branch.

Cesar


2015-06-10  Cesar Philippidis  <ce...@codesourcery.com>

	gcc/
	* c-family/c-omp.c (oacc_filter_device_types): Filter out
	device_type clauses even when the device_type isn't recognized.

	gcc/testsuite
	* c-c++-common/goacc/dtype-3.c: New test.

diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 1c82bf5..c30e0d8 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -1109,6 +1109,7 @@ oacc_filter_device_types (tree clauses)
   tree dtype = NULL_TREE;
   tree seen_nvidia = NULL_TREE;
   tree seen_default = NULL_TREE;
+  int device = 0;
 
   /* First scan for all device_type clauses.  */
   for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
@@ -1119,33 +1120,43 @@ oacc_filter_device_types (tree clauses)
 
 	  if (code == GOMP_DEVICE_DEFAULT)
 	    {
-	      if (seen_default)
+	      if (device & (1 << GOMP_DEVICE_DEFAULT))
 		{
 		  seen_default = NULL_TREE;
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "duplicate device_type (*)");
 		  goto filter_error;
 		}
-	      else
-		seen_default = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
+
+	      seen_default = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
 	    }
-	  if (code & (1 << GOMP_DEVICE_NVIDIA_PTX))
+	  else if (code & (1 << GOMP_DEVICE_NVIDIA_PTX))
 	    {
-	      if (seen_nvidia)
+	      if (device & code)
 		{
 		  seen_nvidia = NULL_TREE;
 		  error_at (OMP_CLAUSE_LOCATION (c),
 			    "duplicate device_type (nvidia)");
 		  goto filter_error;
 		}
-	      else
-		seen_nvidia = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
+
+	      seen_nvidia = OMP_CLAUSE_DEVICE_TYPE_CLAUSES (c);
+	    }
+	  else
+	    {
+	      if (device & (1 << code))
+		{
+		  error_at (OMP_CLAUSE_LOCATION (c),
+			    "duplicate device_type");
+		  goto filter_error;
+		}
 	    }
+	  device |= (1 << code);
 	}
     }
 
   /* Don't do anything if there aren't any device_type clauses.  */
-  if (seen_nvidia == NULL_TREE && seen_default == NULL_TREE)
+  if (device == 0)
     return clauses;
 
   dtype = seen_nvidia ? seen_nvidia : seen_default;
diff --git a/gcc/testsuite/c-c++-common/goacc/dtype-3.c b/gcc/testsuite/c-c++-common/goacc/dtype-3.c
new file mode 100644
index 0000000..53ab94c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/dtype-3.c
@@ -0,0 +1,34 @@
+/* { dg-do compile } */
+
+float b;
+#pragma acc declare link (b)
+
+int
+main (int argc, char **argv)
+{
+  float a;
+
+  a = 2.0;
+
+#pragma acc parallel device_type (*) copy (a) /* { dg-error "not valid" } */
+  {
+  }
+
+#pragma acc parallel device_type (acc_device_nvidia) num_gangs (1)
+  {
+  }
+
+#pragma acc parallel device_type (acc_device_host, acc_device_nvidia) num_gangs (1)
+  {
+  }
+
+#pragma acc parallel device_type (acc_device_host) num_gangs (1)
+  {
+  }
+
+#pragma acc parallel device_type (foo)
+  {
+  }
+
+  return 0;
+}