Ping.
On 19 Nov 12:33, Michael V. Zolotukhin wrote:
> 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.zolotuk...@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
>
