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

Hi Jakub,

Here is a patch for generation of tables containing omp-functions addresses.
It is just a first step, as it lacks generation of similar tables for globals,
but having it in the branch would ease our further development, as we would be
able to base on this.

This patch introduces new function OMP_FINISH_FILE which goes through cgraph and
saves address of every non-external function with 'omp declare target' attribute
to a new symbol.  Then this symbol is stored to a dedicated section
'.omp_fn_table_section'.  When several files are linked together, these sections
are concatenated, so we have an aggregate table with function address.  In a
further patches I'm going to 1) extend this functionality to globals, 2) add a
special symbol at the beginning of the aggregated section (for that I'd add a
dedicated file with this symbol and pass it to the linker as a first input
file).

Is this patch ok for gomp4-branch?

Changelog:

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

	* omp-low.c (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.


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

On Fri, Nov 15, 2013 at 12:35:12PM +0400, Michael V. Zolotukhin wrote:
> 2013-11-15  Michael Zolotukhin  <michael.v.zolotukhin@gmail.com>
> 
> 	* omp-low.c (omp_finish_file): New.
> 	* omp-low.h (omp_finish_file): Declare new function.
> 	* toplev.c: include "omp-low.h"

That would be

	* toplev.c: Include omp-low.h.

> 	(compile_file): Call omp_finish_file.
> 
> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -12101,4 +12101,46 @@ make_pass_omp_simd_clone (gcc::context *ctxt)
>    return new pass_omp_simd_clone (ctxt);
>  }
>  
> +/* Create new symbol containing addresses of functions.  */
> +void
> +omp_finish_file (void)
> +{
> +  struct cgraph_node *node;
> +  const char *section_name = ".omp_fn_table_section";
> +  tree new_decl = build_decl (UNKNOWN_LOCATION,
> +      VAR_DECL, get_identifier (".omp_fn_table"), void_type_node);

Formatting.  Put = on the next line (indented by 4 spaces).
Furthermore, please don't create a decl with a wrong type and then just
change it.

> +  vec<constructor_elt, va_gc> *v;
> +  tree ctor;
> +  int num = 0;
> +
> +  /* TODO: we need to create similar table for global variables.  */
> +  vec_alloc (v, 0);
> +  FOR_EACH_FUNCTION (node)
> +    {
> +      if (DECL_EXTERNAL (node->decl)
> +	  || !lookup_attribute ("omp declare target",
> +				DECL_ATTRIBUTES (node->decl)))
> +	continue;

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.

> +      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.

> +      num++;
> +    }

Please handle here VAR_DECLs here, it is easy.

> +  ctor
> +    = build_constructor (build_array_type_nelts (pointer_sized_int_node, num), v);

Too long line.  Plus calling build_array_type_nelts twice is unnecessary.

> +  TREE_CONSTANT (ctor) = 1;
> +  TREE_STATIC (ctor) = 1;

Why don't you create new_decl only here, with the correct type already?

> +  TREE_STATIC (new_decl) = 1;
> +  DECL_INITIAL (new_decl) = ctor;
> +  TREE_TYPE (new_decl) = build_array_type_nelts (pointer_sized_int_node,
> +						 num);

> +  DECL_SIZE (new_decl) = build_int_cst (integer_type_node,
> +					int_size_in_bytes (pointer_sized_int_node)
> +					* num * BITS_PER_UNIT);
> +  DECL_SIZE_UNIT (new_decl)
> +    = build_int_cst (integer_type_node,
> +		     int_size_in_bytes (pointer_sized_int_node) * num);

Ugh no, layout_decl will do that for you.

> +  DECL_SECTION_NAME (new_decl) = build_string (strlen (section_name), section_name);

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

> +
> +  varpool_assemble_decl ( varpool_node_for_decl (new_decl));

Formatting (extra space after ( ).

> --- 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 ();

Only call it if (flag_openmp).

	Jakub