diff mbox series

[committed] openmp: Fix up declare target handling for vars with DECL_LOCAL_DECL_ALIAS [PR102640]

Message ID 20211008090317.GE304296@tucnak
State New
Headers show
Series [committed] openmp: Fix up declare target handling for vars with DECL_LOCAL_DECL_ALIAS [PR102640] | expand

Commit Message

Jakub Jelinek Oct. 8, 2021, 9:03 a.m. UTC 
Hi!

The introduction of DECL_LOCAL_DECL_ALIAS and push_local_extern_decl_alias
in r11-3699-g4e62aca0e0520e4ed2532f2d8153581190621c1a broke the following
testcase.  The following patch fixes it by treating similarly not just
the variable to or link clause is put on, but also its DECL_LOCAL_DECL_ALIAS
if any.  If it hasn't been created yet, when it is created it will copy
attributes and therefore should get it for free, and as it is an extern,
nothing more than attributes is needed for it.

Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.

2021-10-08  Jakub Jelinek  <jakub@redhat.com>

	PR c++/102640
gcc/cp/
	* parser.c (handle_omp_declare_target_clause): New function.
	(cp_parser_omp_declare_target): Use it.
gcc/testsuite/
	* c-c++-common/gomp/pr102640.c: New test.


	Jakub
diff mbox series

Patch

--- gcc/cp/parser.c.jj	2021-10-07 12:52:34.986912260 +0200
+++ gcc/cp/parser.c	2021-10-07 16:36:12.748015996 +0200
@@ -45505,6 +45505,71 @@  cp_parser_late_parsing_omp_declare_simd
   return attrs;
 }
 
+/* Helper for cp_parser_omp_declare_target, handle one to or link clause
+   on #pragma omp declare target.  Return false if errors were reported.  */
+
+static bool
+handle_omp_declare_target_clause (tree c, tree t, int device_type)
+{
+  tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t));
+  tree at2 = lookup_attribute ("omp declare target link", DECL_ATTRIBUTES (t));
+  tree id;
+  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK)
+    {
+      id = get_identifier ("omp declare target link");
+      std::swap (at1, at2);
+    }
+  else
+    id = get_identifier ("omp declare target");
+  if (at2)
+    {
+      error_at (OMP_CLAUSE_LOCATION (c),
+		"%qD specified both in declare target %<link%> and %<to%>"
+		" clauses", t);
+      return false;
+    }
+  if (!at1)
+    {
+      DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+      if (TREE_CODE (t) != FUNCTION_DECL && !is_global_var (t))
+	return true;
+
+      symtab_node *node = symtab_node::get (t);
+      if (node != NULL)
+	{
+	  node->offloadable = 1;
+	  if (ENABLE_OFFLOADING)
+	    {
+	      g->have_offload = true;
+	      if (is_a <varpool_node *> (node))
+		vec_safe_push (offload_vars, t);
+	    }
+	}
+    }
+  if (TREE_CODE (t) != FUNCTION_DECL)
+    return true;
+  if ((device_type & OMP_CLAUSE_DEVICE_TYPE_HOST) != 0)
+    {
+      tree at3 = lookup_attribute ("omp declare target host",
+				   DECL_ATTRIBUTES (t));
+      if (at3 == NULL_TREE)
+	{
+	  id = get_identifier ("omp declare target host");
+	  DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+	}
+    }
+  if ((device_type & OMP_CLAUSE_DEVICE_TYPE_NOHOST) != 0)
+    {
+      tree at3 = lookup_attribute ("omp declare target nohost",
+				   DECL_ATTRIBUTES (t));
+      if (at3 == NULL_TREE)
+	{
+	  id = get_identifier ("omp declare target nohost");
+	  DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
+	}
+    }
+  return true;
+}
 
 /* OpenMP 4.0:
    # pragma omp declare target new-line
@@ -45557,67 +45622,16 @@  cp_parser_omp_declare_target (cp_parser
     {
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE_TYPE)
 	continue;
-      tree t = OMP_CLAUSE_DECL (c), id;
-      tree at1 = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (t));
-      tree at2 = lookup_attribute ("omp declare target link",
-				   DECL_ATTRIBUTES (t));
+      tree t = OMP_CLAUSE_DECL (c);
       only_device_type = false;
-      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LINK)
-	{
-	  id = get_identifier ("omp declare target link");
-	  std::swap (at1, at2);
-	}
-      else
-	id = get_identifier ("omp declare target");
-      if (at2)
-	{
-	  error_at (OMP_CLAUSE_LOCATION (c),
-		    "%qD specified both in declare target %<link%> and %<to%>"
-		    " clauses", t);
-	  continue;
-	}
-      if (!at1)
-	{
-	  DECL_ATTRIBUTES (t) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
-	  if (TREE_CODE (t) != FUNCTION_DECL && !is_global_var (t))
-	    continue;
-
-	  symtab_node *node = symtab_node::get (t);
-	  if (node != NULL)
-	    {
-	      node->offloadable = 1;
-	      if (ENABLE_OFFLOADING)
-		{
-		  g->have_offload = true;
-		  if (is_a <varpool_node *> (node))
-		    vec_safe_push (offload_vars, t);
-		}
-	    }
-	}
-      if (TREE_CODE (t) != FUNCTION_DECL)
+      if (!handle_omp_declare_target_clause (c, t, device_type))
 	continue;
-      if ((device_type & OMP_CLAUSE_DEVICE_TYPE_HOST) != 0)
-	{
-	  tree at3 = lookup_attribute ("omp declare target host",
-				       DECL_ATTRIBUTES (t));
-	  if (at3 == NULL_TREE)
-	    {
-	      id = get_identifier ("omp declare target host");
-	      DECL_ATTRIBUTES (t)
-		= tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
-	    }
-	}
-      if ((device_type & OMP_CLAUSE_DEVICE_TYPE_NOHOST) != 0)
-	{
-	  tree at3 = lookup_attribute ("omp declare target nohost",
-				       DECL_ATTRIBUTES (t));
-	  if (at3 == NULL_TREE)
-	    {
-	      id = get_identifier ("omp declare target nohost");
-	      DECL_ATTRIBUTES (t)
-		= tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (t));
-	    }
-	}
+      if (VAR_OR_FUNCTION_DECL_P (t)
+	  && DECL_LOCAL_DECL_P (t)
+	  && DECL_LANG_SPECIFIC (t)
+	  && DECL_LOCAL_DECL_ALIAS (t))
+	handle_omp_declare_target_clause (c, DECL_LOCAL_DECL_ALIAS (t),
+					  device_type);
     }
   if (device_type && only_device_type)
     warning_at (OMP_CLAUSE_LOCATION (clauses), 0,
--- gcc/testsuite/c-c++-common/gomp/pr102640.c.jj	2021-10-07 17:35:37.809208766 +0200
+++ gcc/testsuite/c-c++-common/gomp/pr102640.c	2021-10-07 17:13:16.750954087 +0200
@@ -0,0 +1,44 @@ 
+/* PR c++/102640 */
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-omplower" } */
+/* Verify var[123] are mapped without any copying, because they are
+   mentioned in declare target directive to clauses.  */
+/* { dg-final { scan-tree-dump-not "firstprivate\\\(var\[123]\\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump-not ".omp_data_arr.\[0-9]*.var" "omplower" } } */
+/* { dg-final { scan-tree-dump-not ".omp_data_i->var" "omplower" } } */
+
+void
+foo (void)
+{
+  extern int var1;
+  #pragma omp declare target to (var1)
+
+  #pragma omp target
+  var1++;
+}
+
+int
+bar (int x)
+{
+  extern int var2;
+  #pragma omp declare target to (var2)
+  if (x)
+    return var2;
+  #pragma omp target
+  var2++;
+  return -1;
+}
+#pragma omp declare target to (bar)
+
+#pragma omp declare target
+int
+baz (int x)
+{
+  extern int var3;
+  if (x)
+    return var3;
+  #pragma omp target
+  var3++;
+  return -1;
+}
+#pragma omp end declare target