diff mbox

[gomp4] filter out unrecognized device_type targets

Message ID 5578AD9E.2090300@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis June 10, 2015, 9:35 p.m. UTC 
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

Comments

Thomas Schwinge June 11, 2015, 7:19 a.m. UTC | #1 
Hi Cesar!

On Wed, 10 Jun 2015 14:35:26 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch addresses a bug where the c and c++ front ends would ICE [...]

Thanks!


The following also doesn't look quite appropriate:

    $ cat < ../d.f
          IMPLICIT NONE
    
          INTEGER X
    
    !$ACC PARALLEL DTYPE (NVIDIA) NUM_GANGS(10) COPYOUT(X)
          X = 0
    !$ACC END PARALLEL
    
          END
    $ build-gcc/gcc/xgcc -Bbuild-gcc/gcc/ ../d.f -Wall -fopenacc../d.f:5:29:
    
     !$ACC PARALLEL DTYPE (NVIDIA) NUM_GANGS(10) COPYOUT(X)
                                 1
    Error: Invalid character in name at (1)
    ../d.f:7:72: Error: Unexpected !$ACC END PARALLEL statement at (1)

(It should fail with a diagnostic that data clauses not allowed after
device_type.)


Grüße,
 Thomas
diff mbox

Patch

2015-06-10  Cesar Philippidis  <cesar@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;
+}