[PATCH] openmp: Implicit 'declare target' for C++ static initializers

[PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Kwok Cheung Yeung]

[-- Attachment #1: Type: text/plain, Size: 1247 bytes --]

Hello

OpenMP 5.0 has a new feature for implicitly marking variables and functions that 
are referenced in the initializers of static variables and functions that are 
already marked 'declare target'. Support was added in the commit 'openmp: 
Implement discovery of implicit declare target to clauses' 
(dc703151d4f4560e647649506d5b4ceb0ee11e90). However, this does not work with 
non-constant C++ initializers, where the initializers can contain references to 
other (non-constant) variables and function calls.

The C++ front-end stores the initialization information in the static_aggregates 
list (with the variable decl in the TREE_VALUE of an entry and the 
initialization in TREE_PURPOSE) rather than in TREE_INITIAL(var_decl). I have 
added an extra function in omp-offload.cpp to walk the variable initialiser 
trees in static_aggregates, and added a call to it from the FE shortly before 
the initializations are emitted. I have also added a testcase to ensure that the 
implicitly marked variables/functions can be referenced in offloaded code.

The libgomp tests have been run with offloading to a Nvidia card with no 
regressions, and I have also bootstrapped the compiler with no offloading on 
x86-64. Okay for trunk?

Thanks

Kwok

[-- Attachment #2: c++_init_implicit_target.patch --]
[-- Type: text/plain, Size: 4741 bytes --]

commit d2c8c5bd2826851b727e93a8ea2141596e50a621
Author: Kwok Cheung Yeung <kcy@codesourcery.com>
Date:   Wed Oct 28 07:13:14 2020 -0700

    openmp: Implicitly add 'declare target' directives for dynamic static initializers in C++
    
    2020-10-28  Kwok Cheung Yeung  <kcy@codesourcery.com>
    
    	cp/
    	* decl2.c: Include omp-offload.h
    	(c_parse_final_cleanups): Call omp_mark_target_static_initializers.
    	* omp-offload.c (omp_discover_declare_target_var_r): Add all static
    	variables to worklist.
    	(omp_discover_implicit_declare_target): Check that worklist items
    	that are variable declarations have an initialization expression
    	before walking.
    	(omp_mark_target_static_initializers): New.
    	* omp-offload.h	(omp_mark_target_static_initializers): New prototype.
    
    	libgomp/
    	* testsuite/libgomp.c++/declare_target-3.C: New.

diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c
index 2f0d637..b207d58 100644
--- a/gcc/cp/decl2.c
+++ b/gcc/cp/decl2.c
@@ -48,6 +48,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "intl.h"
 #include "c-family/c-ada-spec.h"
 #include "asan.h"
+#include "omp-offload.h"
 
 /* Id for dumping the raw trees.  */
 int raw_dump_id;
@@ -4970,6 +4971,12 @@ c_parse_final_cleanups (void)
 	  /* Make sure the back end knows about all the variables.  */
 	  write_out_vars (vars);
 
+	  /* Mark functions and variables in static initializers as
+	     'omp declare target' if the initialized variable is marked
+	     as such.  */
+	  if (flag_openmp)
+	    omp_mark_target_static_initializers (vars);
+
 	  /* Set the line and file, so that it is obviously not from
 	     the source file.  */
 	  input_location = locus_at_end_of_parsing;
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 3e9c31d..8ecc181 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -296,7 +296,7 @@ omp_discover_declare_target_var_r (tree *tp, int *walk_subtrees, void *data)
 	  DECL_ATTRIBUTES (*tp)
 	    = remove_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp));
 	}
-      if (TREE_STATIC (*tp) && DECL_INITIAL (*tp))
+      if (TREE_STATIC (*tp))
 	((vec<tree> *) data)->safe_push (*tp);
       DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp));
       symtab_node *node = symtab_node::get (*tp);
@@ -348,7 +348,7 @@ omp_discover_implicit_declare_target (void)
   while (!worklist.is_empty ())
     {
       tree decl = worklist.pop ();
-      if (VAR_P (decl))
+      if (VAR_P (decl) && DECL_INITIAL (decl))
 	walk_tree_without_duplicates (&DECL_INITIAL (decl),
 				      omp_discover_declare_target_var_r,
 				      &worklist);
@@ -363,6 +363,33 @@ omp_discover_implicit_declare_target (void)
     }
 }
 
+void
+omp_mark_target_static_initializers (tree vars)
+{
+  tree node;
+  auto_vec<tree> worklist;
+
+  for (node = vars; node; node = TREE_CHAIN (node))
+    if (omp_declare_target_var_p (TREE_VALUE (node)))
+	worklist.safe_push (TREE_VALUE (node));
+
+  while (!worklist.is_empty ())
+    {
+      tree decl = worklist.pop ();
+
+      if (!VAR_P (decl) || !TREE_STATIC (decl))
+	continue;
+
+      for (node = vars; node; node = TREE_CHAIN (node))
+	if (TREE_VALUE (node) == decl)
+	  {
+	    walk_tree_without_duplicates (&TREE_PURPOSE (node),
+					  omp_discover_declare_target_var_r,
+					  &worklist);
+	    break;
+	  }
+    }
+}
 
 /* Create new symbols containing (address, size) pairs for global variables,
    marked with "omp declare target" attribute, as well as addresses for the
diff --git a/gcc/omp-offload.h b/gcc/omp-offload.h
index 0809189..a23a4d3 100644
--- a/gcc/omp-offload.h
+++ b/gcc/omp-offload.h
@@ -31,5 +31,6 @@ extern GTY(()) vec<tree, va_gc> *offload_vars;
 
 extern void omp_finish_file (void);
 extern void omp_discover_implicit_declare_target (void);
+extern void omp_mark_target_static_initializers (tree vars);
 
 #endif /* GCC_OMP_DEVICE_H */
diff --git a/libgomp/testsuite/libgomp.c++/declare_target-3.C b/libgomp/testsuite/libgomp.c++/declare_target-3.C
new file mode 100644
index 0000000..c545613
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/declare_target-3.C
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#include <stdlib.h>
+
+int f()
+{
+  return 54321;
+}
+
+int g()
+{
+  return 3333;
+}
+
+static int x = f() + 7777;
+static int y = g() + x + 12345;
+
+#pragma omp declare target(y)
+
+int main()
+{
+  int err = 0;
+  #pragma omp target map(from:err)
+  {
+    err |= x != 62098;
+    err |= y != 77776;
+    err |= f() != 54321;
+    err |= g() != 3333;
+  }
+  if (err)
+    abort ();
+  return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "__attribute__\\(\\(omp declare target\\)\\)" 2 "gimple" } } */

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Jakub Jelinek]

On Wed, Oct 28, 2020 at 02:20:29PM +0000, Kwok Cheung Yeung wrote:
> OpenMP 5.0 has a new feature for implicitly marking variables and functions
> that are referenced in the initializers of static variables and functions
> that are already marked 'declare target'. Support was added in the commit
> 'openmp: Implement discovery of implicit declare target to clauses'
> (dc703151d4f4560e647649506d5b4ceb0ee11e90). However, this does not work with
> non-constant C++ initializers, where the initializers can contain references
> to other (non-constant) variables and function calls.
> 
> The C++ front-end stores the initialization information in the
> static_aggregates list (with the variable decl in the TREE_VALUE of an entry
> and the initialization in TREE_PURPOSE) rather than in
> TREE_INITIAL(var_decl). I have added an extra function in omp-offload.cpp to
> walk the variable initialiser trees in static_aggregates, and added a call
> to it from the FE shortly before the initializations are emitted. I have
> also added a testcase to ensure that the implicitly marked
> variables/functions can be referenced in offloaded code.

I'm actually not sure how this can work correctly.
Let's say we have 
int foo () { return 1; }
int bar () { return 2; }
int baz () { return 3; }
int qux () { return 4; }
int a = foo ();
int b = bar ();
int c = baz ();
int *d = &c;
int e = qux ();
int f = e + 1;
int *g = &f;
#pragma omp declare target to (b, d, g)
So, for the implicit declare target discovery, a is not declare target to,
nor is foo, and everything else is; b, d, g explicitly, c because it is
referenced in initializer of b, f because it is mentioned in initializer of
g and e because it is mentioned in initializer of f.
Haven't checked if the new function you've added is called before or after
analyze_function calls omp_discover_implicit_declare_target, but I don't
really see how it can work when it is not inside of that function, so that
discovery of new static vars that are implicitly declare target to doesn't
result in marking of its dynamic initializers too.  Perhaps we need a
langhook for that.  But if it is a separate function, either it is called
before the other discovery and will ignore static initializers for vars
that will only be marked as implicit declare target to later, or it is done
afterwards, but then it would really need to duplicate everything what the
other function does, otherwise it woiuldn't discover everything.

Anyway, that is one thing, the other is even if the implicit declare target
discovery handles those correctly, the question is what should we do
afterwards.  Because the C++ FE normally creates a single function that
performs the dynamic initialization of the TUs variables.  But that function
shouldn't be really declare target to, it initializes not only (explicit or
implicit) declare target to variables, but also host only variables.
So we'll probably need to create next to that host only TU constructor
also a device only constructor function that will only initialize the
declare target to variables.

	Jakub

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Kwok Cheung Yeung]

[-- Attachment #1: Type: text/plain, Size: 4012 bytes --]

On 29/10/2020 10:03 am, Jakub Jelinek wrote:
> I'm actually not sure how this can work correctly.
> Let's say we have
> int foo () { return 1; }
> int bar () { return 2; }
> int baz () { return 3; }
> int qux () { return 4; }
> int a = foo ();
> int b = bar ();
> int c = baz ();
> int *d = &c;
> int e = qux ();
> int f = e + 1;
> int *g = &f;
> #pragma omp declare target to (b, d, g)
> So, for the implicit declare target discovery, a is not declare target to,
> nor is foo, and everything else is; b, d, g explicitly, c because it is
> referenced in initializer of b, f because it is mentioned in initializer of
> g and e because it is mentioned in initializer of f.
> Haven't checked if the new function you've added is called before or after
> analyze_function calls omp_discover_implicit_declare_target, but I don't
> really see how it can work when it is not inside of that function, so that
> discovery of new static vars that are implicitly declare target to doesn't
> result in marking of its dynamic initializers too.  Perhaps we need a
> langhook for that.  But if it is a separate function, either it is called
> before the other discovery and will ignore static initializers for vars
> that will only be marked as implicit declare target to later, or it is done
> afterwards, but then it would really need to duplicate everything what the
> other function does, otherwise it woiuldn't discover everything.
>

I have added a new langhook GET_DECL_INIT that by default returns the 
DECL_INITIAL of a variable declaration, but for C++ can also return the dynamic 
initializer if present. omp_discover_implicit_declare_target and 
omp_discover_declare_target_var_r have been changed to use the new langhook 
instead of using DECL_INITIAL.

The dynamic initializer information is stored in a new variable 
dynamic_initializers. The information is originally stored in static_aggregates, 
but this is nulled by calling prune_vars_needing_no_initialization in 
c_parse_final_cleanups. I copy the information into a separate variable before 
it is discarded - this avoids any potential problems that may be caused by 
trying to change the way that static_aggregates currently works.

With this, all the functions and variables in your example are marked correctly:

foo ()
...

__attribute__((omp declare target))
bar ()
...

__attribute__((omp declare target))
baz ()
...

__attribute__((omp declare target))
qux ()
...

.offload_var_table:
         .quad   g
         .quad   8
         .quad   d
         .quad   8
         .quad   b
         .quad   4
         .quad   c
         .quad   4
         .quad   f
         .quad   4
         .quad   e
         .quad   4

Your example is now a compile test in g++.dg/gomp/.

> Anyway, that is one thing, the other is even if the implicit declare target
> discovery handles those correctly, the question is what should we do
> afterwards.  Because the C++ FE normally creates a single function that
> performs the dynamic initialization of the TUs variables.  But that function
> shouldn't be really declare target to, it initializes not only (explicit or
> implicit) declare target to variables, but also host only variables.
> So we'll probably need to create next to that host only TU constructor
> also a device only constructor function that will only initialize the
> declare target to variables.

Even without this patch, G++ currently accepts something like

int foo() { return 1; }
int x = foo();
#pragma omp declare target to(x)

but will not generate the device-side initializer for x, even though x is now 
present on the device. So this part of the implementation is broken with or 
without the patch.

Given that my patch doesn't make the current situation any worse, can I commit 
this portion of it to trunk for now, and leave device-side dynamic 
initialization for later?

Bootstrapped on x86_64 with no offloading, G++ testsuite ran with no 
regressions, and no regressions in the libgomp testsuite with Nvidia offloading.

Thanks,

Kwok

[-- Attachment #2: 0001-openmp-Implicitly-add-declare-target-directives-for-.patch --]
[-- Type: text/plain, Size: 10254 bytes --]

From 0348b149474d0922d79209705e6777e7af271e0d Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Wed, 18 Nov 2020 13:54:01 -0800
Subject: [PATCH] openmp: Implicitly add 'declare target' directives for
 dynamic initializers in C++

2020-11-18  Kwok Cheung Yeung  <kcy@codesourcery.com>

	gcc/
	* langhooks-def.h (lhd_get_decl_init): New.
	(LANG_HOOKS_GET_DECL_INIT): New.
	(LANG_HOOKS_DECLS): Add LANG_HOOKS_GET_DECL_INIT.
	* langhooks.h (struct lang_hooks_for_decls): Add get_decl_init.
	* omp-offload.c (omp_discover_declare_target_var_r): Use
	get_decl_init langhook in place of DECL_INITIAL.

	gcc/cp/
	* cp-lang.c (cxx_get_decl_init): New.
	(LANG_HOOKS_GET_DECL_INIT): New.
	* cp-tree.h (dynamic_initializers): New.
	* decl.c (dynamic_initializers): New.
	* decl2.c (c_parse_final_cleanups): Copy vars into
	dynamic_initializers.

	gcc/testsuite/
	* g++.dg/gomp/declare-target-3.C: New.
---
 gcc/cp/cp-lang.c                             | 24 +++++++++++++++++++++
 gcc/cp/cp-tree.h                             |  2 ++
 gcc/cp/decl.c                                |  6 ++++++
 gcc/cp/decl2.c                               |  5 +++++
 gcc/langhooks-def.h                          |  5 ++++-
 gcc/langhooks.c                              |  8 +++++++
 gcc/langhooks.h                              |  5 +++++
 gcc/omp-offload.c                            |  9 ++++----
 gcc/testsuite/g++.dg/gomp/declare-target-3.C | 31 ++++++++++++++++++++++++++++
 9 files changed, 90 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/g++.dg/gomp/declare-target-3.C

diff --git a/gcc/cp/cp-lang.c b/gcc/cp/cp-lang.c
index 9e980bc..44bf847 100644
--- a/gcc/cp/cp-lang.c
+++ b/gcc/cp/cp-lang.c
@@ -34,6 +34,7 @@ static tree cp_eh_personality (void);
 static tree get_template_innermost_arguments_folded (const_tree);
 static tree get_template_argument_pack_elems_folded (const_tree);
 static tree cxx_enum_underlying_base_type (const_tree);
+static tree* cxx_get_decl_init (tree);
 
 /* Lang hooks common to C++ and ObjC++ are declared in cp/cp-objcp-common.h;
    consequently, there should be very few hooks below.  */
@@ -86,6 +87,9 @@ static tree cxx_enum_underlying_base_type (const_tree);
 #undef LANG_HOOKS_GET_SUBSTRING_LOCATION
 #define LANG_HOOKS_GET_SUBSTRING_LOCATION c_get_substring_location
 
+#undef LANG_HOOKS_GET_DECL_INIT
+#define LANG_HOOKS_GET_DECL_INIT cxx_get_decl_init
+
 /* Each front end provides its own lang hook initializer.  */
 struct lang_hooks lang_hooks = LANG_HOOKS_INITIALIZER;
 
@@ -227,6 +231,26 @@ tree cxx_enum_underlying_base_type (const_tree type)
   return underlying_type;
 }
 
+/* The C++ version of the get_decl_init langhook returns the static
+   initializer for a variable declaration if present, otherwise it
+   tries to find and return the dynamic initializer.  If not present,
+   it returns NULL.  */
+
+static tree*
+cxx_get_decl_init (tree decl)
+{
+  tree node;
+
+  if (DECL_INITIAL (decl))
+    return &DECL_INITIAL (decl);
+
+  for (node = dynamic_initializers; node; node = TREE_CHAIN (node))
+    if (TREE_VALUE (node) == decl)
+      return &TREE_PURPOSE (node);
+
+  return NULL;
+}
+
 #if CHECKING_P
 
 namespace selftest {
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 81485de..53277e1 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5509,6 +5509,8 @@ extern GTY(()) tree static_aggregates;
 /* Likewise, for thread local storage.  */
 extern GTY(()) tree tls_aggregates;
 
+extern GTY(()) tree dynamic_initializers;
+
 enum overload_flags { NO_SPECIAL = 0, DTOR_FLAG, TYPENAME_FLAG };
 
 /* These are uses as bits in flags passed to various functions to
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index d90e984..44558da 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -146,6 +146,12 @@ tree static_aggregates;
 /* Like static_aggregates, but for thread_local variables.  */
 tree tls_aggregates;
 
+/* A list of all objects with initializers.  Unlike static_aggegates,
+   this list is not cleared by the frontend.  The decl is stored in
+   the TREE_VALUE slot and the initializer is stored in the
+   TREE_PURPOSE slot.  */
+tree dynamic_initializers;
+
 /* -- end of C++ */
 
 /* A node for the integer constant 2.  */
diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c
index 1bc7b7e..4fcf82a 100644
--- a/gcc/cp/decl2.c
+++ b/gcc/cp/decl2.c
@@ -4940,6 +4940,11 @@ c_parse_final_cleanups (void)
 	 loop.  */
       vars = prune_vars_needing_no_initialization (&static_aggregates);
 
+      /* Copy the contents of VARS into DYNAMIC_INITIALIZERS.  */
+      for (t = vars; t; t = TREE_CHAIN (t))
+	dynamic_initializers = tree_cons (TREE_PURPOSE (t), TREE_VALUE (t),
+					  dynamic_initializers);
+
       if (vars)
 	{
 	  /* We need to start a new initialization function each time
diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h
index 2f66f5e..d19fbee 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -87,6 +87,7 @@ extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
 extern bool lhd_omp_mappable_type (tree);
 extern bool lhd_omp_scalar_p (tree);
+extern tree* lhd_get_decl_init (tree);
 
 extern const char *lhd_get_substring_location (const substring_loc &,
 					       location_t *out_loc);
@@ -265,6 +266,7 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
 #define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null
 #define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause
 #define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p
+#define LANG_HOOKS_GET_DECL_INIT lhd_get_decl_init
 
 #define LANG_HOOKS_DECLS { \
   LANG_HOOKS_GLOBAL_BINDINGS_P, \
@@ -293,7 +295,8 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
   LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR, \
   LANG_HOOKS_OMP_CLAUSE_DTOR, \
   LANG_HOOKS_OMP_FINISH_CLAUSE, \
-  LANG_HOOKS_OMP_SCALAR_P \
+  LANG_HOOKS_OMP_SCALAR_P, \
+  LANG_HOOKS_GET_DECL_INIT \
 }
 
 /* LTO hooks.  */
diff --git a/gcc/langhooks.c b/gcc/langhooks.c
index d82f542..6aa96bb 100644
--- a/gcc/langhooks.c
+++ b/gcc/langhooks.c
@@ -632,6 +632,14 @@ lhd_omp_scalar_p (tree decl)
   return false;
 }
 
+/* Return static initializer for DECL.  */
+
+tree*
+lhd_get_decl_init (tree decl)
+{
+  return &DECL_INITIAL (decl);
+}
+
 /* Register language specific type size variables as potentially OpenMP
    firstprivate variables.  */
 
diff --git a/gcc/langhooks.h b/gcc/langhooks.h
index f12589e..6cf2c59 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -299,6 +299,11 @@ struct lang_hooks_for_decls
   /* Return true if DECL is a scalar variable (for the purpose of
      implicit firstprivatization).  */
   bool (*omp_scalar_p) (tree decl);
+
+  /* Return a pointer to the tree representing the initializer
+     expression for the non-local variable DECL.  Return NULL if
+     DECL is not initialized.  */
+  tree* (*get_decl_init) (tree decl);
 };
 
 /* Language hooks related to LTO serialization.  */
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 9013961..d1c69ed 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -315,7 +315,7 @@ omp_discover_declare_target_var_r (tree *tp, int *walk_subtrees, void *data)
 	  DECL_ATTRIBUTES (*tp)
 	    = remove_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp));
 	}
-      if (TREE_STATIC (*tp) && DECL_INITIAL (*tp))
+      if (TREE_STATIC (*tp) && lang_hooks.decls.get_decl_init (*tp))
 	((vec<tree> *) data)->safe_push (*tp);
       DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp));
       symtab_node *node = symtab_node::get (*tp);
@@ -361,14 +361,15 @@ omp_discover_implicit_declare_target (void)
 		   && DECL_STRUCT_FUNCTION (cgn->decl)->has_omp_target)
 	    worklist.safe_push (cgn->decl);
       }
-  FOR_EACH_STATIC_INITIALIZER (vnode)
-    if (omp_declare_target_var_p (vnode->decl))
+  FOR_EACH_VARIABLE (vnode)
+    if (lang_hooks.decls.get_decl_init (vnode->decl)
+	&& omp_declare_target_var_p (vnode->decl))
       worklist.safe_push (vnode->decl);
   while (!worklist.is_empty ())
     {
       tree decl = worklist.pop ();
       if (VAR_P (decl))
-	walk_tree_without_duplicates (&DECL_INITIAL (decl),
+	walk_tree_without_duplicates (lang_hooks.decls.get_decl_init (decl),
 				      omp_discover_declare_target_var_r,
 				      &worklist);
       else if (omp_declare_target_fn_p (decl))
diff --git a/gcc/testsuite/g++.dg/gomp/declare-target-3.C b/gcc/testsuite/g++.dg/gomp/declare-target-3.C
new file mode 100644
index 0000000..8e9eafc
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-target-3.C
@@ -0,0 +1,31 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// Test implicit marking of declare target to.
+
+int foo () { return 1; }
+int bar () { return 2; }	// Implicitly marked (due to b)
+int baz () { return 3; }	// Implicitly marked (due to d via c)
+int qux () { return 4; }	// Implicitly marked (due to g via f and e)
+
+int a = foo ();
+int b = bar ();	// Explicitly marked
+int c = baz ();	// Implicitly marked (due to d)
+int *d = &c;	// Explicitly marked
+int e = qux ();	// Implicitly marked (due to g via f)
+int f = e + 1;	// Implicitly marked (due to g)
+int *g = &f;	// Explicitly marked
+
+#pragma omp declare target to(b, d, g)
+
+// { dg-final { scan-tree-dump-not "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nfoo \\\(\\\)" "gimple" } }
+// { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nbar \\\(\\\)" "gimple" } }
+// { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nbaz \\\(\\\)" "gimple" } }
+// { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nqux \\\(\\\)" "gimple" } }
+// { dg-final { scan-assembler-not "\\\.offload_var_table:\\n.+\\\.quad\\s+a" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+b" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+c" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+d" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+e" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+f" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+g" } }
-- 
2.8.1

PING Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Kwok Cheung Yeung]

Hello

This patch still needs review.

Thanks

Kwok

On 19/11/2020 6:07 pm, Kwok Cheung Yeung wrote:
> On 29/10/2020 10:03 am, Jakub Jelinek wrote:
>> I'm actually not sure how this can work correctly.
>> Let's say we have
>> int foo () { return 1; }
>> int bar () { return 2; }
>> int baz () { return 3; }
>> int qux () { return 4; }
>> int a = foo ();
>> int b = bar ();
>> int c = baz ();
>> int *d = &c;
>> int e = qux ();
>> int f = e + 1;
>> int *g = &f;
>> #pragma omp declare target to (b, d, g)
>> So, for the implicit declare target discovery, a is not declare target to,
>> nor is foo, and everything else is; b, d, g explicitly, c because it is
>> referenced in initializer of b, f because it is mentioned in initializer of
>> g and e because it is mentioned in initializer of f.
>> Haven't checked if the new function you've added is called before or after
>> analyze_function calls omp_discover_implicit_declare_target, but I don't
>> really see how it can work when it is not inside of that function, so that
>> discovery of new static vars that are implicitly declare target to doesn't
>> result in marking of its dynamic initializers too.  Perhaps we need a
>> langhook for that.  But if it is a separate function, either it is called
>> before the other discovery and will ignore static initializers for vars
>> that will only be marked as implicit declare target to later, or it is done
>> afterwards, but then it would really need to duplicate everything what the
>> other function does, otherwise it woiuldn't discover everything.
>>
> 
> I have added a new langhook GET_DECL_INIT that by default returns the 
> DECL_INITIAL of a variable declaration, but for C++ can also return the dynamic 
> initializer if present. omp_discover_implicit_declare_target and 
> omp_discover_declare_target_var_r have been changed to use the new langhook 
> instead of using DECL_INITIAL.
> 
> The dynamic initializer information is stored in a new variable 
> dynamic_initializers. The information is originally stored in static_aggregates, 
> but this is nulled by calling prune_vars_needing_no_initialization in 
> c_parse_final_cleanups. I copy the information into a separate variable before 
> it is discarded - this avoids any potential problems that may be caused by 
> trying to change the way that static_aggregates currently works.
> 
> With this, all the functions and variables in your example are marked correctly:
> 
> foo ()
> ...
> 
> __attribute__((omp declare target))
> bar ()
> ...
> 
> __attribute__((omp declare target))
> baz ()
> ...
> 
> __attribute__((omp declare target))
> qux ()
> ...
> 
> .offload_var_table:
>          .quad   g
>          .quad   8
>          .quad   d
>          .quad   8
>          .quad   b
>          .quad   4
>          .quad   c
>          .quad   4
>          .quad   f
>          .quad   4
>          .quad   e
>          .quad   4
> 
> Your example is now a compile test in g++.dg/gomp/.
> 
>> Anyway, that is one thing, the other is even if the implicit declare target
>> discovery handles those correctly, the question is what should we do
>> afterwards.  Because the C++ FE normally creates a single function that
>> performs the dynamic initialization of the TUs variables.  But that function
>> shouldn't be really declare target to, it initializes not only (explicit or
>> implicit) declare target to variables, but also host only variables.
>> So we'll probably need to create next to that host only TU constructor
>> also a device only constructor function that will only initialize the
>> declare target to variables.
> 
> Even without this patch, G++ currently accepts something like
> 
> int foo() { return 1; }
> int x = foo();
> #pragma omp declare target to(x)
> 
> but will not generate the device-side initializer for x, even though x is now 
> present on the device. So this part of the implementation is broken with or 
> without the patch.
> 
> Given that my patch doesn't make the current situation any worse, can I commit 
> this portion of it to trunk for now, and leave device-side dynamic 
> initialization for later?
> 
> Bootstrapped on x86_64 with no offloading, G++ testsuite ran with no 
> regressions, and no regressions in the libgomp testsuite with Nvidia offloading.
> 
> Thanks,
> 
> Kwok

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Jakub Jelinek]

On Thu, Nov 19, 2020 at 06:07:28PM +0000, Kwok Cheung Yeung wrote:
> Even without this patch, G++ currently accepts something like

Sorry for the delay.

> int foo() { return 1; }
> int x = foo();
> #pragma omp declare target to(x)
> 
> but will not generate the device-side initializer for x, even though x is
> now present on the device. So this part of the implementation is broken with
> or without the patch.
> 
> Given that my patch doesn't make the current situation any worse, can I
> commit this portion of it to trunk for now, and leave device-side dynamic
> initialization for later?

Ok, but for the patch I have a few nits:

> +/* The C++ version of the get_decl_init langhook returns the static
> +   initializer for a variable declaration if present, otherwise it
> +   tries to find and return the dynamic initializer.  If not present,
> +   it returns NULL.  */
> +
> +static tree*
> +cxx_get_decl_init (tree decl)

The GCC coding style (appart from libstdc++) is type * rather than type*,
occurs several times in the patch.

> +{
> +  tree node;
> +
> +  if (DECL_INITIAL (decl))
> +    return &DECL_INITIAL (decl);
> +
> +  for (node = dynamic_initializers; node; node = TREE_CHAIN (node))
> +    if (TREE_VALUE (node) == decl)
> +      return &TREE_PURPOSE (node);

I'm worried with many dynamic initializers this will be worst case
quadratic.  Can't you use instead a hash map?  Note, as this is in the
FE, we might need to worry about PCH and GC.
Thus the hash map needs to be indexed by DECL_UIDs rather than pointers,
so perhaps use decl_tree_map?
Also, I'm worried that nothing releases dynamic_initializers (or the
decl_tree_map replacement).  We need it only during the discovery and not
afterwards, so it would be nice if the omp declare target discovery at the
end called another lang hook that would free the decl_tree_map, so that GC
can take it all.
If trees would remain there afterwards, we'd need to worry about destructive
gimplifier too and would need to unshare the dynamic initializers or
something.

I think it would be best to use omp_ in the hook name(s), and:
> --- a/gcc/cp/decl2.c
> +++ b/gcc/cp/decl2.c
> @@ -4940,6 +4940,11 @@ c_parse_final_cleanups (void)
>  	 loop.  */
>        vars = prune_vars_needing_no_initialization (&static_aggregates);
>  
> +      /* Copy the contents of VARS into DYNAMIC_INITIALIZERS.  */
> +      for (t = vars; t; t = TREE_CHAIN (t))
> +	dynamic_initializers = tree_cons (TREE_PURPOSE (t), TREE_VALUE (t),
> +					  dynamic_initializers);

Not to add there anything if (!flag_openmp).  We don't need to waste memory
when nobody is going to look at it.

	Jakub

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Kwok Cheung Yeung]

[-- Attachment #1: Type: text/plain, Size: 2254 bytes --]

On 08/12/2020 4:24 pm, Jakub Jelinek wrote:
> The GCC coding style (appart from libstdc++) is type * rather than type*,
> occurs several times in the patch.

Fixed.

>> +{
>> +  tree node;
>> +
>> +  if (DECL_INITIAL (decl))
>> +    return &DECL_INITIAL (decl);
>> +
>> +  for (node = dynamic_initializers; node; node = TREE_CHAIN (node))
>> +    if (TREE_VALUE (node) == decl)
>> +      return &TREE_PURPOSE (node);
> 
> I'm worried with many dynamic initializers this will be worst case
> quadratic.  Can't you use instead a hash map?  Note, as this is in the
> FE, we might need to worry about PCH and GC.
> Thus the hash map needs to be indexed by DECL_UIDs rather than pointers,
> so perhaps use decl_tree_map?
> Also, I'm worried that nothing releases dynamic_initializers (or the
> decl_tree_map replacement).  We need it only during the discovery and not
> afterwards, so it would be nice if the omp declare target discovery at the
> end called another lang hook that would free the decl_tree_map, so that GC
> can take it all.
> If trees would remain there afterwards, we'd need to worry about destructive
> gimplifier too and would need to unshare the dynamic initializers or
> something.
> 
> I think it would be best to use omp_ in the hook name(s), and:

I have now changed dynamic_initializers to use a decl_tree_map instead. 
get_decl_init has been renamed to omp_get_decl_init, and I have added a hook 
omp_finish_decl_inits which is called at the end of 
omp_discover_implicit_declare_target to free the decl_tree_map for GC.

>> --- a/gcc/cp/decl2.c
>> +++ b/gcc/cp/decl2.c
>> @@ -4940,6 +4940,11 @@ c_parse_final_cleanups (void)
>>   	 loop.  */
>>         vars = prune_vars_needing_no_initialization (&static_aggregates);
>>   
>> +      /* Copy the contents of VARS into DYNAMIC_INITIALIZERS.  */
>> +      for (t = vars; t; t = TREE_CHAIN (t))
>> +	dynamic_initializers = tree_cons (TREE_PURPOSE (t), TREE_VALUE (t),
>> +					  dynamic_initializers);
> 
> Not to add there anything if (!flag_openmp).  We don't need to waste memory
> when nobody is going to look at it.

Done.

I have retested all the gomp tests in the main testsuite, retested libgomp, and 
checked bootstrapping. Is this version okay for trunk now?

Thanks

Kwok

[-- Attachment #2: 0001-openmp-Implicitly-add-declare-target-directives-for-.patch --]
[-- Type: text/plain, Size: 11950 bytes --]

From ef4a42c5174372dd0d72dc0efe2c608e693c7877 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Thu, 17 Dec 2020 12:10:18 -0800
Subject: [PATCH] openmp: Implicitly add 'declare target' directives for
 dynamic initializers in C++

2020-12-17  Kwok Cheung Yeung  <kcy@codesourcery.com>

	gcc/
	* langhooks-def.h (lhd_get_decl_init): New.
	(lhd_finish_decl_inits): New.
	(LANG_HOOKS_GET_DECL_INIT): New.
	(LANG_HOOKS_OMP_FINISH_DECL_INITS): New.
	(LANG_HOOKS_DECLS): Add LANG_HOOKS_GET_DECL_INIT and
	LANG_HOOKS_OMP_FINISH_DECL_INITS.
	* langhooks.c (lhd_omp_get_decl_init): New.
	(lhd_omp_finish_decl_inits): New.
	* langhooks.h (struct lang_hooks_for_decls): Add omp_get_decl_init
	and omp_finish_decl_inits.
	* omp-offload.c (omp_discover_declare_target_var_r): Use
	get_decl_init langhook in place of DECL_INITIAL.  Call
	omp_finish_decl_inits langhook at end of function.

	gcc/cp/
	* cp-lang.c (cxx_get_decl_init): New.
	(cxx_omp_finish_decl_inits): New.
	(LANG_HOOKS_GET_DECL_INIT): New.
	(LANG_HOOKS_OMP_FINISH_DECL_INITS): New.
	* cp-tree.h (dynamic_initializers): New.
	* decl.c (dynamic_initializers): New.
	* decl2.c (c_parse_final_cleanups): Add initializer entries
	from vars to dynamic_initializers.

	gcc/testsuite/
	* g++.dg/gomp/declare-target-3.C: New.
---
 gcc/cp/cp-lang.c                             | 32 ++++++++++++++++++++++++++++
 gcc/cp/cp-tree.h                             |  4 ++++
 gcc/cp/decl.c                                |  4 ++++
 gcc/cp/decl2.c                               |  7 ++++++
 gcc/langhooks-def.h                          |  8 ++++++-
 gcc/langhooks.c                              | 16 ++++++++++++++
 gcc/langhooks.h                              | 10 +++++++++
 gcc/omp-offload.c                            | 11 ++++++----
 gcc/testsuite/g++.dg/gomp/declare-target-3.C | 31 +++++++++++++++++++++++++++
 9 files changed, 118 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/g++.dg/gomp/declare-target-3.C

diff --git a/gcc/cp/cp-lang.c b/gcc/cp/cp-lang.c
index 5d2aef4..bde11db 100644
--- a/gcc/cp/cp-lang.c
+++ b/gcc/cp/cp-lang.c
@@ -34,6 +34,8 @@ static tree cp_eh_personality (void);
 static tree get_template_innermost_arguments_folded (const_tree);
 static tree get_template_argument_pack_elems_folded (const_tree);
 static tree cxx_enum_underlying_base_type (const_tree);
+static tree * cxx_omp_get_decl_init (tree);
+static void cxx_omp_finish_decl_inits (void);
 
 /* Lang hooks common to C++ and ObjC++ are declared in cp/cp-objcp-common.h;
    consequently, there should be very few hooks below.  */
@@ -92,6 +94,12 @@ static tree cxx_enum_underlying_base_type (const_tree);
 #undef LANG_HOOKS_GET_SUBSTRING_LOCATION
 #define LANG_HOOKS_GET_SUBSTRING_LOCATION c_get_substring_location
 
+#undef LANG_HOOKS_OMP_GET_DECL_INIT
+#define LANG_HOOKS_OMP_GET_DECL_INIT cxx_omp_get_decl_init
+
+#undef LANG_HOOKS_OMP_FINISH_DECL_INITS
+#define LANG_HOOKS_OMP_FINISH_DECL_INITS cxx_omp_finish_decl_inits
+
 /* Each front end provides its own lang hook initializer.  */
 struct lang_hooks lang_hooks = LANG_HOOKS_INITIALIZER;
 
@@ -233,6 +241,30 @@ tree cxx_enum_underlying_base_type (const_tree type)
   return underlying_type;
 }
 
+/* The C++ version of the omp_get_decl_init langhook returns the static
+   initializer for a variable declaration if present, otherwise it
+   tries to find and return the dynamic initializer.  If not present,
+   it returns NULL.  */
+
+static tree *
+cxx_omp_get_decl_init (tree decl)
+{
+  if (DECL_INITIAL (decl))
+    return &DECL_INITIAL (decl);
+
+  return hash_map_safe_get (dynamic_initializers, decl);
+}
+
+/* The C++ version of the omp_finish_decl_inits langhook allows GC to
+   reclaim the memory used by the hash-map used to hold dynamic initializer
+   information.  */
+
+static void
+cxx_omp_finish_decl_inits (void)
+{
+  dynamic_initializers = NULL;
+}
+
 #if CHECKING_P
 
 namespace selftest {
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index ef5baea..edaa594 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5631,6 +5631,10 @@ extern GTY(()) tree static_aggregates;
 /* Likewise, for thread local storage.  */
 extern GTY(()) tree tls_aggregates;
 
+/* A hash-map mapping from variable decls to the dynamic initializer for
+   the decl.  This is currently only used by OpenMP.  */
+extern GTY(()) decl_tree_map *dynamic_initializers;
+
 enum overload_flags { NO_SPECIAL = 0, DTOR_FLAG, TYPENAME_FLAG };
 
 /* These are uses as bits in flags passed to various functions to
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index b56eb11..c349716 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -146,6 +146,10 @@ tree static_aggregates;
 /* Like static_aggregates, but for thread_local variables.  */
 tree tls_aggregates;
 
+/* A hash-map mapping from variable decls to the dynamic initializer for
+   the decl.  This is currently only used by OpenMP.  */
+decl_tree_map *dynamic_initializers;
+
 /* -- end of C++ */
 
 /* A node for the integer constant 2.  */
diff --git a/gcc/cp/decl2.c b/gcc/cp/decl2.c
index 6d8158a..af88e7f 100644
--- a/gcc/cp/decl2.c
+++ b/gcc/cp/decl2.c
@@ -5006,6 +5006,13 @@ c_parse_final_cleanups (void)
 	 loop.  */
       if (tree vars = prune_vars_needing_no_initialization (&static_aggregates))
 	{
+	  if (flag_openmp)
+	    /* Add initializer information from VARS into
+	       DYNAMIC_INITIALIZERS.  */
+	    for (t = vars; t; t = TREE_CHAIN (t))
+	      hash_map_safe_put<hm_ggc> (dynamic_initializers,
+					 TREE_VALUE (t), TREE_PURPOSE (t));
+
 	  /* We need to start a new initialization function each time
 	     through the loop.  That's because we need to know which
 	     vtables have been referenced, and TREE_SYMBOL_REFERENCED
diff --git a/gcc/langhooks-def.h b/gcc/langhooks-def.h
index 2f66f5e..59178d2 100644
--- a/gcc/langhooks-def.h
+++ b/gcc/langhooks-def.h
@@ -87,6 +87,8 @@ extern void lhd_omp_firstprivatize_type_sizes (struct gimplify_omp_ctx *,
 					       tree);
 extern bool lhd_omp_mappable_type (tree);
 extern bool lhd_omp_scalar_p (tree);
+extern tree * lhd_omp_get_decl_init (tree);
+extern void lhd_omp_finish_decl_inits ();
 
 extern const char *lhd_get_substring_location (const substring_loc &,
 					       location_t *out_loc);
@@ -265,6 +267,8 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
 #define LANG_HOOKS_OMP_CLAUSE_DTOR hook_tree_tree_tree_null
 #define LANG_HOOKS_OMP_FINISH_CLAUSE lhd_omp_finish_clause
 #define LANG_HOOKS_OMP_SCALAR_P lhd_omp_scalar_p
+#define LANG_HOOKS_OMP_GET_DECL_INIT lhd_omp_get_decl_init
+#define LANG_HOOKS_OMP_FINISH_DECL_INITS lhd_omp_finish_decl_inits
 
 #define LANG_HOOKS_DECLS { \
   LANG_HOOKS_GLOBAL_BINDINGS_P, \
@@ -293,7 +297,9 @@ extern tree lhd_unit_size_without_reusable_padding (tree);
   LANG_HOOKS_OMP_CLAUSE_LINEAR_CTOR, \
   LANG_HOOKS_OMP_CLAUSE_DTOR, \
   LANG_HOOKS_OMP_FINISH_CLAUSE, \
-  LANG_HOOKS_OMP_SCALAR_P \
+  LANG_HOOKS_OMP_SCALAR_P, \
+  LANG_HOOKS_OMP_GET_DECL_INIT, \
+  LANG_HOOKS_OMP_FINISH_DECL_INITS \
 }
 
 /* LTO hooks.  */
diff --git a/gcc/langhooks.c b/gcc/langhooks.c
index d82f542..23cbf06 100644
--- a/gcc/langhooks.c
+++ b/gcc/langhooks.c
@@ -632,6 +632,22 @@ lhd_omp_scalar_p (tree decl)
   return false;
 }
 
+/* Return static initializer for DECL.  */
+
+tree *
+lhd_omp_get_decl_init (tree decl)
+{
+  return &DECL_INITIAL (decl);
+}
+
+/* Free any extra memory used to hold initializer information for
+   variable declarations.  */
+
+void
+lhd_omp_finish_decl_inits (void)
+{
+}
+
 /* Register language specific type size variables as potentially OpenMP
    firstprivate variables.  */
 
diff --git a/gcc/langhooks.h b/gcc/langhooks.h
index f12589e..6b90794 100644
--- a/gcc/langhooks.h
+++ b/gcc/langhooks.h
@@ -299,6 +299,16 @@ struct lang_hooks_for_decls
   /* Return true if DECL is a scalar variable (for the purpose of
      implicit firstprivatization).  */
   bool (*omp_scalar_p) (tree decl);
+
+  /* Return a pointer to the tree representing the initializer
+     expression for the non-local variable DECL.  Return NULL if
+     DECL is not initialized.  */
+  tree * (*omp_get_decl_init) (tree decl);
+
+  /* Free any extra memory used to hold initializer information for
+     variable declarations.  omp_get_decl_init must not be called
+     after calling this.  */
+  void (*omp_finish_decl_inits) (void);
 };
 
 /* Language hooks related to LTO serialization.  */
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 9013961..15b735b 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -315,7 +315,7 @@ omp_discover_declare_target_var_r (tree *tp, int *walk_subtrees, void *data)
 	  DECL_ATTRIBUTES (*tp)
 	    = remove_attribute ("omp declare target link", DECL_ATTRIBUTES (*tp));
 	}
-      if (TREE_STATIC (*tp) && DECL_INITIAL (*tp))
+      if (TREE_STATIC (*tp) && lang_hooks.decls.omp_get_decl_init (*tp))
 	((vec<tree> *) data)->safe_push (*tp);
       DECL_ATTRIBUTES (*tp) = tree_cons (id, NULL_TREE, DECL_ATTRIBUTES (*tp));
       symtab_node *node = symtab_node::get (*tp);
@@ -361,14 +361,15 @@ omp_discover_implicit_declare_target (void)
 		   && DECL_STRUCT_FUNCTION (cgn->decl)->has_omp_target)
 	    worklist.safe_push (cgn->decl);
       }
-  FOR_EACH_STATIC_INITIALIZER (vnode)
-    if (omp_declare_target_var_p (vnode->decl))
+  FOR_EACH_VARIABLE (vnode)
+    if (lang_hooks.decls.omp_get_decl_init (vnode->decl)
+	&& omp_declare_target_var_p (vnode->decl))
       worklist.safe_push (vnode->decl);
   while (!worklist.is_empty ())
     {
       tree decl = worklist.pop ();
       if (VAR_P (decl))
-	walk_tree_without_duplicates (&DECL_INITIAL (decl),
+	walk_tree_without_duplicates (lang_hooks.decls.omp_get_decl_init (decl),
 				      omp_discover_declare_target_var_r,
 				      &worklist);
       else if (omp_declare_target_fn_p (decl))
@@ -380,6 +381,8 @@ omp_discover_implicit_declare_target (void)
 				      omp_discover_declare_target_fn_r,
 				      &worklist);
     }
+
+  lang_hooks.decls.omp_finish_decl_inits ();
 }
 
 
diff --git a/gcc/testsuite/g++.dg/gomp/declare-target-3.C b/gcc/testsuite/g++.dg/gomp/declare-target-3.C
new file mode 100644
index 0000000..d2dedaf
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/declare-target-3.C
@@ -0,0 +1,31 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+// Test implicit marking of declare target to.
+
+int foo () { return 1; }
+int bar () { return 2; }	// Implicitly marked (due to b)
+int baz () { return 3; }	// Implicitly marked (due to d via c)
+int qux () { return 4; }	// Implicitly marked (due to g via f and e)
+
+int a = foo ();
+int b = bar ();	// Explicitly marked
+int c = baz ();	// Implicitly marked (due to d)
+int *d = &c;	// Explicitly marked
+int e = qux ();	// Implicitly marked (due to g via f)
+int f = e + 1;	// Implicitly marked (due to g)
+int *g = &f;	// Explicitly marked
+
+#pragma omp declare target to(b, d, g)
+
+// { dg-final { scan-tree-dump-not "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint foo \\\(\\\)" "gimple" } }
+// { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint bar \\\(\\\)" "gimple" } }
+// { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint baz \\\(\\\)" "gimple" } }
+// { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint qux \\\(\\\)" "gimple" } }
+// { dg-final { scan-assembler-not "\\\.offload_var_table:\\n.+\\\.quad\\s+a" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+b" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+c" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+d" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+e" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+f" } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+g" } }
-- 
2.8.1

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Jakub Jelinek]

On Fri, Dec 18, 2020 at 03:10:52PM +0000, Kwok Cheung Yeung wrote:
> --- a/gcc/cp/cp-lang.c
> +++ b/gcc/cp/cp-lang.c
> @@ -34,6 +34,8 @@ static tree cp_eh_personality (void);
>  static tree get_template_innermost_arguments_folded (const_tree);
>  static tree get_template_argument_pack_elems_folded (const_tree);
>  static tree cxx_enum_underlying_base_type (const_tree);
> +static tree * cxx_omp_get_decl_init (tree);

No space between * and cxx_omp_get_decl_init here.

> +extern tree * lhd_omp_get_decl_init (tree);

Ditto.

> +  /* Return a pointer to the tree representing the initializer
> +     expression for the non-local variable DECL.  Return NULL if
> +     DECL is not initialized.  */
> +  tree * (*omp_get_decl_init) (tree decl);

Ditto.

Ok for trunk with those nits fixed.

Thanks!

	Jakub

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Jakub Jelinek]

On Fri, Dec 18, 2020 at 03:10:52PM +0000, Kwok Cheung Yeung wrote:
> 2020-12-17  Kwok Cheung Yeung  <kcy@codesourcery.com>
> 
> 	gcc/testsuite/
> 	* g++.dg/gomp/declare-target-3.C: New.

Note the test fails on the trunk when one doesn't have offloading
configured.  IL scan tests are always problematic, different between
offloading and no offloading...

	Jakub

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Kwok Cheung Yeung]

[-- Attachment #1: Type: text/plain, Size: 600 bytes --]

On 18/12/2020 7:31 pm, Jakub Jelinek wrote:
> On Fri, Dec 18, 2020 at 03:10:52PM +0000, Kwok Cheung Yeung wrote:
>> 2020-12-17  Kwok Cheung Yeung  <kcy@codesourcery.com>
>>
>> 	gcc/testsuite/
>> 	* g++.dg/gomp/declare-target-3.C: New.
> 
> Note the test fails on the trunk when one doesn't have offloading
> configured.  IL scan tests are always problematic, different between
> offloading and no offloading...

Oops. This patch disables the scan for .offload_var_table entries in the 
assembler if offloading is not enabled. The gimple tests appear to be okay though?

Okay for trunk?

Thanks

Kwok

[-- Attachment #2: declare-target-3.C_fix.patch --]
[-- Type: text/plain, Size: 2406 bytes --]

commit f427d4eaddbd1ee4001e057b231c92fdd9fc66f5
Author: Kwok Cheung Yeung <kcy@codesourcery.com>
Date:   Fri Dec 18 12:05:20 2020 -0800

    openmp: Fix g++.dg/gomp/declare-target-3.C testcase when offloading is disabled
    
    2020-12-18  Kwok Cheung Yeung  <kcy@codesourcery.com>
    
    	gcc/testsuite/
    	* g++.dg/gomp/declare-target-3.C: Only check .offload_var_table
    	entries if offloading is enabled.

diff --git a/gcc/testsuite/g++.dg/gomp/declare-target-3.C b/gcc/testsuite/g++.dg/gomp/declare-target-3.C
index d2dedaf..1e23c86 100644
--- a/gcc/testsuite/g++.dg/gomp/declare-target-3.C
+++ b/gcc/testsuite/g++.dg/gomp/declare-target-3.C
@@ -22,10 +22,10 @@ int *g = &f;	// Explicitly marked
 // { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint bar \\\(\\\)" "gimple" } }
 // { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint baz \\\(\\\)" "gimple" } }
 // { dg-final { scan-tree-dump "__attribute__\\\(\\\(omp declare target\\\)\\\)\\\nint qux \\\(\\\)" "gimple" } }
-// { dg-final { scan-assembler-not "\\\.offload_var_table:\\n.+\\\.quad\\s+a" } }
-// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+b" } }
-// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+c" } }
-// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+d" } }
-// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+e" } }
-// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+f" } }
-// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+g" } }
+// { dg-final { scan-assembler-not "\\\.offload_var_table:\\n.+\\\.quad\\s+a" { target { offloading_enabled } } } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+b" { target { offloading_enabled } } } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+c" { target { offloading_enabled } } } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+d" { target { offloading_enabled } } } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+e" { target { offloading_enabled } } } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+f" { target { offloading_enabled } } } }
+// { dg-final { scan-assembler "\\\.offload_var_table:\\n.+\\\.quad\\s+g" { target { offloading_enabled } } } }

Re: [PATCH] openmp: Implicit 'declare target' for C++ static initializers

[Jakub Jelinek]

On Fri, Dec 18, 2020 at 08:15:56PM +0000, Kwok Cheung Yeung wrote:
> On 18/12/2020 7:31 pm, Jakub Jelinek wrote:
> > On Fri, Dec 18, 2020 at 03:10:52PM +0000, Kwok Cheung Yeung wrote:
> > > 2020-12-17  Kwok Cheung Yeung  <kcy@codesourcery.com>
> > > 
> > > 	gcc/testsuite/
> > > 	* g++.dg/gomp/declare-target-3.C: New.
> > 
> > Note the test fails on the trunk when one doesn't have offloading
> > configured.  IL scan tests are always problematic, different between
> > offloading and no offloading...
> 
> Oops. This patch disables the scan for .offload_var_table entries in the
> assembler if offloading is not enabled. The gimple tests appear to be okay
> though?
> 
> Okay for trunk?

Ok, thanks.

> commit f427d4eaddbd1ee4001e057b231c92fdd9fc66f5
> Author: Kwok Cheung Yeung <kcy@codesourcery.com>
> Date:   Fri Dec 18 12:05:20 2020 -0800
> 
>     openmp: Fix g++.dg/gomp/declare-target-3.C testcase when offloading is disabled
>     
>     2020-12-18  Kwok Cheung Yeung  <kcy@codesourcery.com>
>     
>     	gcc/testsuite/
>     	* g++.dg/gomp/declare-target-3.C: Only check .offload_var_table
>     	entries if offloading is enabled.

	Jakub