Re: [gomp4] declare directive [2/5]

[James Norris <jnorris at codesourcery dot com>]

diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 261a12d..15da51e 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -78,6 +78,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "cilk.h"
 #include "wide-int.h"
 #include "builtins.h"
+#include "gomp-constants.h"
 
 /* Possible cases of bad specifiers type used by bad_specifiers. */
 enum bad_spec_place {
@@ -14113,6 +14114,314 @@ maybe_save_function_definition (tree fun)
     register_constexpr_fundef (fun, DECL_SAVED_TREE (fun));
 }
 
+static tree
+check_oacc_vars_1 (tree *tp, int *, void *l)
+{
+  if (TREE_CODE (*tp) == VAR_DECL && TREE_PUBLIC (*tp))
+    {
+      location_t loc = DECL_SOURCE_LOCATION (*tp);
+      tree attrs;
+      attrs = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (*tp));
+      if (attrs)
+	{
+	  tree t;
+
+	  for (t = TREE_VALUE (attrs); t; t = TREE_CHAIN (t))
+	    {
+	      loc = EXPR_LOCATION ((tree) l);
+
+	      if (OMP_CLAUSE_MAP_KIND (TREE_VALUE (t)) == GOMP_MAP_LINK)
+		{
+		  error_at (loc, "%<link%> clause cannot be used with %qE",
+			    *tp);
+		  break;
+		}
+	    }
+	}
+      else
+	error_at (loc, "no %<#pragma acc declare%> for %qE", *tp);
+    }
+  return NULL_TREE;
+}
+
+static tree
+check_oacc_vars (tree *tp, int *, void *)
+{
+  if (TREE_CODE (*tp) == STATEMENT_LIST)
+    {
+      tree_stmt_iterator i;
+
+      for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
+	{
+	  tree t = tsi_stmt (i);
+	  walk_tree_without_duplicates (&t, check_oacc_vars_1, t);
+	}
+    }
+
+  return NULL_TREE;
+}
+
+static struct oacc_return
+{
+  tree_stmt_iterator iter;
+  tree stmt;
+  int op;
+  struct oacc_return *next;
+} *oacc_returns;
+
+static tree
+find_oacc_return (tree *tp, int *, void *)
+{
+  if (TREE_CODE (*tp) == STATEMENT_LIST)
+    {
+      tree_stmt_iterator i;
+
+      for (i = tsi_start (*tp); !tsi_end_p (i); tsi_next (&i))
+	{
+	  tree t;
+	  struct oacc_return *r;
+
+	  t = tsi_stmt (i);
+
+	  if (TREE_CODE (t) == RETURN_EXPR)
+	    {
+	      r = XNEW (struct oacc_return);
+	      r->iter = i;
+	      r->stmt = NULL_TREE;
+	      r->op = 1;
+	      r->next = NULL;
+
+	      if (oacc_returns)
+		r->next = oacc_returns;
+
+	      oacc_returns = r;
+	    }
+	  else if (TREE_CODE (t) == IF_STMT)
+	    {
+	       bool op1, op2;
+	       tree op;
+
+	       op1 = op2 = false;
+
+	       op = TREE_OPERAND (t, 1);
+	       op1 = (op && TREE_CODE (op) == RETURN_EXPR);
+
+	       op = TREE_OPERAND (t, 2);
+	       op2 = (op && TREE_CODE (op) == RETURN_EXPR);
+
+	       if (op1 || op2)
+		{
+		  r = XNEW (struct oacc_return);
+		  r->stmt = t;
+		  r->op = op1 ? 1 : 2;
+		  r->next = NULL;
+
+		  if (oacc_returns)
+		    r->next = oacc_returns;
+
+		  oacc_returns = r;
+		}
+	    }
+	}
+    }
+
+  return NULL_TREE;
+}
+
+static void
+finish_oacc_declare (tree fndecl, tree decls)
+{
+  tree t, stmt, list, c, ret_clauses, clauses;
+  location_t loc;
+  tree_stmt_iterator i;
+
+  list = cur_stmt_list;
+
+  if (lookup_attribute ("oacc function", DECL_ATTRIBUTES (fndecl)))
+    {
+      if (lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl)))
+	{
+	  location_t loc = DECL_SOURCE_LOCATION (fndecl);
+	  error_at (loc, "%<#pragma acc declare%> not allowed in %qE", fndecl);
+	}
+
+      walk_tree_without_duplicates (&list, check_oacc_vars, NULL);
+      return;
+    }
+
+  if (!decls)
+    return;
+
+  walk_tree_without_duplicates (&list, find_oacc_return, NULL);
+
+  clauses = NULL_TREE;
+
+  for (t = decls; t; t = TREE_CHAIN (t))
+    {
+      c = TREE_VALUE (TREE_VALUE (t));
+
+      if (clauses)
+	OMP_CLAUSE_CHAIN (c) = clauses;
+      else
+	loc = OMP_CLAUSE_LOCATION (c);
+
+      clauses = c;
+    }
+
+  ret_clauses = NULL_TREE;
+
+  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    {
+      bool ret = false;
+      HOST_WIDE_INT kind, new_op;
+
+      kind = OMP_CLAUSE_MAP_KIND (c);
+
+      switch (kind)
+	{
+	  case GOMP_MAP_ALLOC:
+	  case GOMP_MAP_FORCE_ALLOC:
+	  case GOMP_MAP_FORCE_TO:
+	    new_op = GOMP_MAP_FORCE_DEALLOC;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FORCE_FROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+	    new_op = GOMP_MAP_FORCE_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FORCE_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_TO);
+	    new_op = GOMP_MAP_FORCE_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_FROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_ALLOC);
+	    new_op = GOMP_MAP_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_TOFROM:
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+	    new_op = GOMP_MAP_FROM;
+	    ret = true;
+	    break;
+
+	  case GOMP_MAP_DEVICE_RESIDENT:
+	  case GOMP_MAP_FORCE_DEVICEPTR:
+	  case GOMP_MAP_FORCE_PRESENT:
+	  case GOMP_MAP_POINTER:
+	  case GOMP_MAP_TO:
+	    break;
+
+	  case GOMP_MAP_LINK:
+	    continue;
+
+	  default:
+	    gcc_unreachable ();
+	    break;
+	}
+
+      if (ret)
+	{
+	  t = copy_node (c);
+
+	  OMP_CLAUSE_SET_MAP_KIND (t, new_op);
+
+	  if (ret_clauses)
+	    OMP_CLAUSE_CHAIN (t) = ret_clauses;
+
+	  ret_clauses = t;
+	}
+    }
+
+   i = tsi_start (list);
+   if (!tsi_end_p (i))
+     {
+	t = tsi_stmt (i);
+	if (TREE_CODE (t) == BIND_EXPR)
+	  list = BIND_EXPR_BODY (t);
+      }
+
+  if (clauses)
+    {
+      bool found = false;
+
+      stmt = make_node (OACC_DECLARE);
+      TREE_TYPE (stmt) = void_type_node;
+      OMP_STANDALONE_CLAUSES (stmt) = clauses;
+      SET_EXPR_LOCATION (stmt, loc);
+
+      c = OMP_CLAUSE_DECL (TREE_VALUE (TREE_VALUE (decls)));
+
+      for (i = tsi_start (list); !tsi_end_p (i); tsi_next (&i))
+	{
+	  tree it;
+
+	  it = tsi_stmt (i);
+
+	  if ((TREE_CODE (it) == DECL_EXPR) && (DECL_EXPR_DECL (it) == c))
+	    {
+	      tsi_link_after (&i, stmt, TSI_CONTINUE_LINKING);
+	      found = true;
+	      break;
+	    }
+	}
+
+	if (!found)
+	  {
+	    i = tsi_start (list);
+	    tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+	  }
+    }
+
+    while (oacc_returns)
+      {
+	struct oacc_return *r;
+
+	stmt = make_node (OACC_DECLARE);
+	TREE_TYPE (stmt) = void_type_node;
+	OMP_STANDALONE_CLAUSES (stmt) = ret_clauses;
+	SET_EXPR_LOCATION (stmt, loc);
+
+	r = oacc_returns;
+	if (r->stmt)
+	  {
+	    tree l;
+
+	    l = alloc_stmt_list ();
+	    append_to_statement_list (stmt, &l);
+	    stmt = TREE_OPERAND (r->stmt, r->op);
+	    append_to_statement_list (stmt, &l);
+	    TREE_OPERAND (r->stmt, r->op) = l;
+	  }
+	else
+	  tsi_link_before (&r->iter, stmt, TSI_CONTINUE_LINKING);
+
+	oacc_returns = r->next;
+	free (r);
+     }
+
+  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);
+
+  tsi_link_before (&i, stmt, TSI_CONTINUE_LINKING);
+
+  DECL_ATTRIBUTES (fndecl)
+	  = remove_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
+}
+
 /* Finish up a function declaration and compile that function
    all the way to assembler language output.  The free the storage
    for the function definition.
@@ -14141,6 +14450,9 @@ finish_function (int flags)
   gcc_assert (!defer_mark_used_calls);
   defer_mark_used_calls = true;
 
+  tree decls = lookup_attribute ("oacc declare", DECL_ATTRIBUTES (fndecl));
+  finish_oacc_declare (fndecl, decls);
+
   record_key_method_defined (fndecl);
 
   fntype = TREE_TYPE (fndecl);
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 2947bf4..fb6b7ed 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -61,6 +61,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "type-utils.h"
 #include "omp-low.h"
 #include "gomp-constants.h"
+#include "tree-iterator.h"
 
 
 /* The lexer.  */
@@ -32035,6 +32036,221 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
   return stmt;
 }
 
+/* OpenACC 2.0:
+   # pragma acc declare oacc-data-clause[optseq] new-line
+*/
+
+static int oacc_dcl_idx = 0;
+
+#define OACC_DECLARE_CLAUSE_MASK					\
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+
+static tree
+cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
+{
+  tree clauses;
+  bool error = false;
+
+  clauses = cp_parser_oacc_all_clauses (parser, OACC_DECLARE_CLAUSE_MASK,
+					"#pragma acc declare", pragma_tok);
+
+  if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+    {
+      error_at (pragma_tok->location,
+		"no valid clauses specified in %<#pragma acc declare%>");
+      return NULL_TREE;
+    }
+
+  for (tree t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+    {
+      location_t loc = OMP_CLAUSE_LOCATION (t);
+      tree decl = OMP_CLAUSE_DECL (t);
+      tree devres = NULL_TREE;
+      if (!DECL_P (decl))
+	{
+	  error_at (loc, "subarray in %<#pragma acc declare%>");
+	  error = true;
+	  continue;
+	}
+      gcc_assert (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_MAP);
+      switch (OMP_CLAUSE_MAP_KIND (t))
+	{
+	case GOMP_MAP_FORCE_ALLOC:
+	case GOMP_MAP_FORCE_TO:
+	case GOMP_MAP_FORCE_DEVICEPTR:
+	  break;
+
+	case GOMP_MAP_DEVICE_RESIDENT:
+	  devres = t;
+	  break;
+
+	case GOMP_MAP_POINTER:
+	  /* Generated by c_finish_omp_clauses from array sections;
+	     avoid spurious diagnostics.  */
+	  break;
+
+	case GOMP_MAP_LINK:
+	  if (!global_bindings_p () && !DECL_EXTERNAL (decl))
+	    {
+	      error_at (loc,
+			"%qD must be a global variable in"
+			"%<#pragma acc declare link%>",
+			decl);
+	      error = true;
+	      continue;
+	    }
+	  break;
+
+	default:
+	  if (global_bindings_p ())
+	    {
+	      error_at (loc, "invalid OpenACC clause at file scope");
+	      error = true;
+	      continue;
+	    }
+	  if (DECL_EXTERNAL (decl))
+	    {
+	      error_at (loc,
+			"invalid use of %<extern%> variable %qD "
+			"in %<#pragma acc declare%>", decl);
+	      error = true;
+	      continue;
+	    }
+	  break;
+	}
+
+      /* Store the clause in an attribute on the variable, at file
+	 scope, or the function, at block scope.  */
+      tree decl_for_attr;
+      if (global_bindings_p ())
+	{
+	  decl_for_attr = decl;
+	  tree prev_attr = lookup_attribute ("oacc declare",
+					     DECL_ATTRIBUTES (decl));
+	  if (prev_attr)
+	    {
+	      tree p = TREE_VALUE (prev_attr);
+	      tree cl = TREE_VALUE (p);
+
+	      if (!devres
+		  && OMP_CLAUSE_MAP_KIND (cl) != GOMP_MAP_DEVICE_RESIDENT)
+		{
+		  error_at (loc,
+			    "variable %qD used more than once with "
+			    "%<#pragma acc declare%>", decl);
+		  inform (OMP_CLAUSE_LOCATION (TREE_VALUE (p)),
+			  "previous directive was here");
+		  error = true;
+		  continue;
+		}
+	    }
+	}
+      else
+	{
+	  decl_for_attr = current_function_decl;
+	  tree prev_attr = lookup_attribute ("oacc declare",
+					     DECL_ATTRIBUTES (decl_for_attr));
+	  for (;
+	       prev_attr;
+	       prev_attr = lookup_attribute ("oacc declare",
+					     TREE_CHAIN (prev_attr)))
+	    {
+	      tree p = TREE_VALUE (prev_attr);
+	      tree cl = TREE_VALUE (p);
+	      if (OMP_CLAUSE_DECL (cl) == decl)
+		{
+		  error_at (loc,
+			    "variable %qD used more than once with "
+			    "%<#pragma acc declare%>", decl);
+		  inform (OMP_CLAUSE_LOCATION (cl),
+			  "previous directive was here");
+		  error = true;
+		  break;
+		}
+	    }
+	}
+
+      if (!error)
+	{
+	  tree attr = tree_cons (NULL_TREE, t, NULL_TREE);
+	  tree attrs = tree_cons (get_identifier ("oacc declare"),
+				  attr, NULL_TREE);
+	  decl_attributes (&decl_for_attr, attrs, 0);
+	}
+    }
+
+  if (error)
+    return NULL_TREE;
+
+  if (global_bindings_p ())
+    {
+      char buf[128];
+      cp_decl_specifier_seq decl_specifiers;
+      cp_declarator *declarator;
+      tree attrs, parms;
+      tree f, t, call_fn, stmt;
+      location_t loc = UNKNOWN_LOCATION;
+      void *p;
+
+      p = obstack_alloc (&declarator_obstack, 0);
+      clear_decl_specs (&decl_specifiers);
+      decl_specifiers.type = void_type_node;
+      sprintf (buf, "__openacc_cp_constructor__%d", oacc_dcl_idx++);
+
+      declarator = make_id_declarator (NULL_TREE, get_identifier (buf),
+				       sfk_none);
+      parms = void_list_node;
+      declarator = make_call_declarator (declarator, parms,
+		      TYPE_UNQUALIFIED,
+		      VIRT_SPEC_UNSPECIFIED,
+		      REF_QUAL_NONE,
+		      NULL_TREE,
+		      NULL_TREE);
+      attrs = tree_cons (get_identifier ("constructor") , NULL_TREE, NULL_TREE);
+      start_function (&decl_specifiers, declarator, attrs);
+      f = begin_compound_stmt (0);
+      call_fn = builtin_decl_explicit (BUILT_IN_GOACC_STATIC);
+      TREE_SIDE_EFFECTS (call_fn) = 1;
+
+      for (t = clauses; t; t = OMP_CLAUSE_CHAIN (t))
+	{
+	  tree d, a1, a2, a3;
+	  vec<tree, va_gc> *args;
+	  vec_alloc (args, 3);
+
+	  d = OMP_CLAUSE_DECL (t);
+
+	  a1 = build_unary_op (loc, ADDR_EXPR, d, 0);
+	  a2 = DECL_SIZE_UNIT (d);
+	  a3 = build_int_cst (unsigned_type_node, OMP_CLAUSE_MAP_KIND (t));
+
+	  args->quick_push (a1);
+	  args->quick_push (a2);
+	  args->quick_push (a3);
+
+	  stmt = build_function_call_vec (loc, vNULL, call_fn, args, NULL);
+	  finish_expr_stmt (stmt);
+	}
+
+      finish_compound_stmt (f);
+      expand_or_defer_fn (finish_function (0));
+      obstack_free (&declarator_obstack, p);
+    }
+
+  return NULL_TREE;
+}
+
 #define OACC_HOST_DATA_CLAUSE_MASK					\
   ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
 
@@ -33903,6 +34119,10 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
       cp_parser_omp_declare (parser, pragma_tok, context);
       return false;
 
+    case PRAGMA_OACC_DECLARE:
+      cp_parser_oacc_declare (parser, pragma_tok);
+      return false;
+
     case PRAGMA_OACC_ENTER_DATA:
       if (context == pragma_stmt)
 	{
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index caafb43..f6e5c3b 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -14230,6 +14230,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl,
       break;
 
     case OMP_TARGET_UPDATE:
+    case OACC_DECLARE:
     case OACC_ENTER_DATA:
     case OACC_EXIT_DATA:
     case OACC_UPDATE: