This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[gomp4] Additional tests for declare directive and fixes.


Hi!

The following patch adds additional testing of the declare directive
and fixes for issues that arose from the testing.


Committed to gomp-4_0-branch.

Jim
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index e7df751..bcbd163 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1767,12 +1767,15 @@ finish_oacc_declare (tree fnbody, tree decls)
 	break;
     }
 
-  stmt = make_node (OACC_DECLARE);
-  TREE_TYPE (stmt) = void_type_node;
-  OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
-  SET_EXPR_LOCATION (stmt, loc);
+  if (ret_clauses)
+    {
+      stmt = make_node (OACC_DECLARE);
+      TREE_TYPE (stmt) = void_type_node;
+      OACC_DECLARE_CLAUSES (stmt) = ret_clauses;
+      SET_EXPR_LOCATION (stmt, loc);
 
-  tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+      tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+    }
 
   DECL_ATTRIBUTES (fndecl)
 	  = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
@@ -12812,6 +12815,14 @@ c_parser_oacc_declare (c_parser *parser)
 	      error = true;
 	      continue;
 	    }
+	  else if (TREE_PUBLIC (decl))
+	    {
+	      error_at (loc,
+			"invalid use of %<global%> variable %qD "
+			"in %<#pragma acc declare%>", decl);
+	      error = true;
+	      continue;
+	    }
 	  break;
 	}
 
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 15da51e..a35f599 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -14343,7 +14343,17 @@ finish_oacc_declare (tree fndecl, tree decls)
      {
 	t = tsi_stmt (i);
 	if (TREE_CODE (t) == BIND_EXPR)
-	  list = BIND_EXPR_BODY (t);
+	  {
+	    list = BIND_EXPR_BODY (t);
+	    if (TREE_CODE (list) != STATEMENT_LIST)
+	      {
+		stmt = list;
+		list = alloc_stmt_list ();
+		BIND_EXPR_BODY (t) = list;
+		i = tsi_start (list);
+		tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
+	      }
+	  }
       }
 
   if (clauses)
@@ -14371,11 +14381,11 @@ finish_oacc_declare (tree fndecl, tree decls)
 	    }
 	}
 
-	if (!found)
-	  {
-	    i = tsi_start (list);
-	    tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
-	  }
+      if (!found)
+	{
+	  i = tsi_start (list);
+	  tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+	}
     }
 
     while (oacc_returns)
@@ -14405,18 +14415,21 @@ finish_oacc_declare (tree fndecl, tree decls)
 	free (r);
      }
 
-  for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
+  if (ret_clauses)
     {
-      if (tsi_end_p (i))
-	break;
-    }
+      for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
+	{
+	  if (tsi_end_p (i))
+	    break;
+	}
 
-  stmt = make_node (OACC_DECLARE);
-  TREE_TYPE (stmt) = void_type_node;
-  OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
-  SET_EXPR_LOCATION (stmt, loc);
+      stmt = make_node (OACC_DECLARE);
+      TREE_TYPE (stmt) = void_type_node;
+      OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
+      SET_EXPR_LOCATION (stmt, loc);
 
-  tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+      tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+    }
 
   DECL_ATTRIBUTES (fndecl)
 	  = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 78bcb0a1..41fb35e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -32123,6 +32123,14 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
 	      error = true;
 	      continue;
 	    }
+	  else if (TREE_PUBLIC (decl))
+	    {
+	      error_at (loc,
+			"invalid use of %<global%> variable %qD "
+			"in %<#pragma acc declare%>", decl);
+	      error = true;
+	      continue;
+	    }
 	  break;
 	}
 
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c
index ce12463..7979f0c 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -63,4 +63,6 @@ f (void)
 
   extern int ve6;
 #pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */
+
+#pragma acc declare present (v9) /* { dg-error "invalid use of" } */
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
index 59cfe51..584b921 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
@@ -4,6 +4,26 @@
 #include <stdlib.h>
 #include <stdio.h>
 
+#define N 8
+
+void
+subr1 (int *a)
+{
+  int f[N];
+#pragma acc declare copy (f)
+
+#pragma acc parallel copy (a[0:N])
+  {
+    int i;
+
+    for (i = 0; i < N; i++)
+      {
+	f[i] = a[i];
+	a[i] = f[i] + f[i];
+      }
+  }
+}
+
 int b[8];
 #pragma acc declare create (b)
 
@@ -13,7 +33,6 @@ int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
 int
 main (int argc, char **argv)
 {
-  const int N = 8;
   int a[N];
   int e[N];
 #pragma acc declare create (e)
@@ -61,5 +80,18 @@ main (int argc, char **argv)
 	abort ();
     }
 
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 1234;
+    }
+
+  subr1 (&a[0]);
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 1234 * 2)
+	abort ();
+    }
+
   return 0;
 }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
new file mode 100644
index 0000000..8bb7dec
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
@@ -0,0 +1,16 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 8
+
+int
+main (int argc, char **argv)
+{
+  int a[N] __attribute__((unused));
+#pragma acc declare present (a)
+}
+
+/* { dg-shouldfail "" } */

Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]