Re: [GOMP4] Generation tables with omp-functions addresses for offloading.

["Michael V. Zolotukhin" <michael dot v dot zolotukhin at gmail dot com>]

Hi Jakub,

Thanks for the remarks.  Updated patch is attached, and my answers are below.

> This will add into the table all "omp declare target" functions, but you
> actually want there only the outlined #pragma omp target bodies.
> The question is how to find them here reliably.  At least ignoring
> !DECL_ARTIFICIAL (node->decl) functions would help, but still would add
> e.g. cloned functions, or say #pragma omp parallel/task outlined bodies
> in omp declare target functions, etc.
> So, perhaps either add some extra attribute, or some flag in cgraph node,
> or change create_omp_child_function_name + create_omp_child_function
> + callers that it doesn't create "_omp_fn" clone suffix for GOMP_target,
> but say "_omp_tgtfn".  I think the last thing would be nice in any case.
> Then you can just check if it is DECL_ARTIFICIAL with strstr (DECL_NAME
> (node->decl), "_omp_tgtfn") != NULL.
That's right.  I added check for DECL_ARTIFICIAL for now.  Renaming to
'_omp_tgtfn' could go in a separate patch I guess.  Variables are handled now as
well.

> 
> > +      CONSTRUCTOR_APPEND_ELT (v, NULL_TREE, build_fold_addr_expr (node->decl));
> 
> As explained, the table shouldn't contain just pointers, but pairs of
> pointer, size.  For FUNCTION_DECLs just use size 1, we want to look up
> only the first byte in them, but for var decls you want to know also
> the size to register in the mapping tree.
> So you need to create the structure with the two pairs, or create the
> table as pointers but push two elements for each function (and VAR_DECL),
> first one address, second one 1 (or size in bytes) fold_converted into
> pointer type.
Yep, I fixed that using the option with table of just pointers, storing
(address, size) pairs.

> You need to check if target actually supports named sections, if not, don't
> create anything.
Fixed.

> > +      omp_finish_file ();
> 
> Only call it if (flag_openmp).
Currently that would break work with '-flto' as we don't have '-fopenmp' in the
options list when calling lto1-compiler.  I think that would be fixed soon, when
we finish with all command-line options stuff.  Also, when no symbols have 'omp
declare target' attribute, this call won't do anything except traversal through
all symbols.  What do you think, is it still worth to be guarded with
if (flag_openmp)?

Changelog:

2013-11-19  Michael Zolotukhin  <michael.v.zolotukhin@gmail.com>

	* omp-low.c: Include common/common-target.h.
	(omp_finish_file): New.
	* omp-low.h (omp_finish_file): Declare new function.
	* toplev.c: Include omp-low.h.
	(compile_file): Call omp_finish_file.


Thanks,
Michael
> 	Jakub

---
 gcc/omp-low.c | 70 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
 gcc/omp-low.h |  1 +
 gcc/toplev.c  |  3 +++
 3 files changed, 74 insertions(+)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 797a492..be458eb 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -51,6 +51,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "optabs.h"
 #include "cfgloop.h"
 #include "target.h"
+#include "common/common-target.h"
 #include "omp-low.h"
 #include "gimple-low.h"
 #include "tree-cfgcleanup.h"
@@ -12101,4 +12102,73 @@ make_pass_omp_simd_clone (gcc::context *ctxt)
   return new pass_omp_simd_clone (ctxt);
 }
 
+/* Create new symbol containing (address, size) pairs for omp-marked
+   functions and global variables.  */
+void
+omp_finish_file (void)
+{
+  struct cgraph_node *node;
+  struct varpool_node *vnode;
+  const char *section_name = ".omp_table_section";
+  tree new_decl, new_decl_type;
+  vec<constructor_elt, va_gc> *v;
+  tree ctor;
+  int num = 0;
+
+  if (!targetm_common.have_named_sections)
+    return;
+
+  vec_alloc (v, 0);
+
+  /* Collect all omp-target functions.  */
+  FOR_EACH_DEFINED_FUNCTION (node)
+    {
+      /* TODO: This check could fail on functions, created by omp
+	 parallel/task pragmas.  It's better to name outlined for offloading
+	 functions in some different way and to check here the function name.
+	 It could be something like "*_omp_tgtfn" in contrast with "*_omp_fn"
+	 for functions from omp parallel/task pragmas.  */
+      if (!lookup_attribute ("omp declare target",
+			     DECL_ATTRIBUTES (node->decl))
+	  || !DECL_ARTIFICIAL (node->decl))
+	continue;
+      CONSTRUCTOR_APPEND_ELT (v, NULL_TREE, build_fold_addr_expr (node->decl));
+      CONSTRUCTOR_APPEND_ELT (v, NULL_TREE,
+			      fold_convert (const_ptr_type_node,
+					    integer_one_node));
+      num += 2;
+    }
+
+  /* Collect all omp-target global variables.  */
+  FOR_EACH_DEFINED_VARIABLE (vnode)
+    {
+      if (!lookup_attribute ("omp declare target",
+			     DECL_ATTRIBUTES (vnode->decl))
+	  || TREE_CODE (vnode->decl) != VAR_DECL
+	  || DECL_SIZE (vnode->decl) == 0)
+	continue;
+
+      CONSTRUCTOR_APPEND_ELT (v, NULL_TREE, build_fold_addr_expr (vnode->decl));
+      CONSTRUCTOR_APPEND_ELT (v, NULL_TREE,
+			      fold_convert (const_ptr_type_node,
+					    DECL_SIZE (vnode->decl)));
+      num += 2;
+    }
+  if (num == 0)
+    return;
+
+  new_decl_type = build_array_type_nelts (pointer_sized_int_node, num);
+  ctor = build_constructor (new_decl_type, v);
+  TREE_CONSTANT (ctor) = 1;
+  TREE_STATIC (ctor) = 1;
+  new_decl = build_decl (UNKNOWN_LOCATION, VAR_DECL,
+			 get_identifier (".omp_table"), new_decl_type);
+  TREE_STATIC (new_decl) = 1;
+  DECL_INITIAL (new_decl) = ctor;
+  DECL_SECTION_NAME (new_decl) = build_string (strlen (section_name),
+					       section_name);
+
+  varpool_assemble_decl (varpool_node_for_decl (new_decl));
+}
+
 #include "gt-omp-low.h"
diff --git a/gcc/omp-low.h b/gcc/omp-low.h
index 6b5a2ff..813189d 100644
--- a/gcc/omp-low.h
+++ b/gcc/omp-low.h
@@ -27,5 +27,6 @@ extern void omp_expand_local (basic_block);
 extern void free_omp_regions (void);
 extern tree omp_reduction_init (tree, tree);
 extern bool make_gimple_omp_edges (basic_block, struct omp_region **);
+extern void omp_finish_file (void);
 
 #endif /* GCC_OMP_LOW_H */
diff --git a/gcc/toplev.c b/gcc/toplev.c
index f78912e..595705a 100644
--- a/gcc/toplev.c
+++ b/gcc/toplev.c
@@ -76,6 +76,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "diagnostic-color.h"
 #include "context.h"
 #include "pass_manager.h"
+#include "omp-low.h"
 
 #if defined(DBX_DEBUGGING_INFO) || defined(XCOFF_DEBUGGING_INFO)
 #include "dbxout.h"
@@ -574,6 +575,8 @@ compile_file (void)
       if (flag_sanitize & SANITIZE_THREAD)
 	tsan_finish_file ();
 
+      omp_finish_file ();
+
       output_shared_constant_pool ();
       output_object_blocks ();
       finish_tm_clone_pairs ();
-- 
1.8.3.1