From patchwork Tue Nov 19 08:33:33 2013 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Michael Zolotukhin X-Patchwork-Id: 292321 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (Client did not present a certificate) by ozlabs.org (Postfix) with ESMTPS id 7B0A62C00F8 for ; Tue, 19 Nov 2013 19:34:56 +1100 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:references:mime-version :content-type:in-reply-to; q=dns; s=default; b=et1++4fWuqxMYNM/N FGycgaCWOeFYqF1G27FI+l4uqdQX3iR+mlaDiSmueLH1t4VfjsuH2I7ziu8i0+TB fLZUQAiDaevMJdkLdh5gWs75/8ilAipjT3MpElHij2mFfdoSQ2vn/prc47ICmwlq Y6AEIBgJ9UI5KeX+8OnYyDgsyU= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:references:mime-version :content-type:in-reply-to; s=default; bh=qLqtjwnmfkgs9eoPD0hJHXS gGXs=; b=ZHCoJV5e/JZa5eJCNHjacP1Z9v+3ZKDYX2GQmse8zvJsT22I9+hQnor Tz0ZeH3BxOfeTeEo+B3u2WUu0+rumsVqhn94mMn0Ps2qf6HaRckQ3Z3Y2iuaPiuX EQ3+ZDevEq5i8vtiq3oOn0XZonRaAVJPdi1oV0ngkWCalD9YnKcw= Received: (qmail 13600 invoked by alias); 19 Nov 2013 08:34:36 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 13582 invoked by uid 89); 19 Nov 2013 08:34:36 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.2 required=5.0 tests=AWL, BAYES_60, FREEMAIL_FROM, RDNS_NONE, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-pa0-f47.google.com Received: from Unknown (HELO mail-pa0-f47.google.com) (209.85.220.47) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Tue, 19 Nov 2013 08:34:31 +0000 Received: by mail-pa0-f47.google.com with SMTP id kq14so3450333pab.20 for ; Tue, 19 Nov 2013 00:34:23 -0800 (PST) X-Received: by 10.68.189.197 with SMTP id gk5mr25193181pbc.37.1384850063597; Tue, 19 Nov 2013 00:34:23 -0800 (PST) Received: from msticlxl57.ims.intel.com ([192.55.54.40]) by mx.google.com with ESMTPSA id ry4sm32666871pab.4.2013.11.19.00.34.18 for (version=TLSv1 cipher=RC4-SHA bits=128/128); Tue, 19 Nov 2013 00:34:19 -0800 (PST) Date: Tue, 19 Nov 2013 12:33:33 +0400 From: "Michael V. Zolotukhin" To: Jakub Jelinek Cc: gcc , Kirill Yukhin Subject: Re: [GOMP4] Generation tables with omp-functions addresses for offloading. Message-ID: <20131119083333.GG3461@msticlxl57.ims.intel.com> References: <20131115083512.GF3461@msticlxl57.ims.intel.com> <20131115133655.GD892@tucnak.redhat.com> MIME-Version: 1.0 Content-Disposition: inline In-Reply-To: <20131115133655.GD892@tucnak.redhat.com> User-Agent: Mutt/1.5.21 (2010-09-15) X-IsSubscribed: yes 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 * 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 *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 ();