From patchwork Wed Oct 21 19:00:47 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 534003 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)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 83069140DAE for ; Thu, 22 Oct 2015 06:01:07 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=Xp2al4qL; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=m8owhA+1vJusv73jl jv/TApGaXYoT68Eaawa9FEu6coimbMAR2s7JHZjCS42aOPVCGflZbUvlzdhn5Ioj Q5zXBDIyTx6J0tBbSpSNF60mEhFtyK2NT+DscyRdKGa8y9qFNu7YlS1yIbB6eMIC xqo4zX6zZjyaL1m8SNcuRmHLV8= 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 :subject:to:references:cc:from:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=fi7S1KWBKk2rPRvXVlxYFty M5H4=; b=Xp2al4qLvCWBTNFYxu7EgRJOI/ez6yOiwTZmkXCSK96Ntrj5ec3FdA4 QSAmvtoDB73ZnaKz0m0Q93GDoEVTj9wj87Lq3RgZcwLy5eQxH3dW9A7YAWQhqCgG 3vFUInFeSLnEPhTl2Usoc54RyjXDJas2gsKQj4Rwe0Dn/cwxDSiA= Received: (qmail 3105 invoked by alias); 21 Oct 2015 19:00:54 -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 3070 invoked by uid 89); 21 Oct 2015 19:00:52 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=BAYES_00, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-qk0-f174.google.com Received: from mail-qk0-f174.google.com (HELO mail-qk0-f174.google.com) (209.85.220.174) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Wed, 21 Oct 2015 19:00:51 +0000 Received: by qkbl190 with SMTP id l190so43431665qkb.2 for ; Wed, 21 Oct 2015 12:00:48 -0700 (PDT) X-Received: by 10.55.200.217 with SMTP id t86mr13205406qkl.33.1445454048638; Wed, 21 Oct 2015 12:00:48 -0700 (PDT) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id 68sm3772894qhc.49.2015.10.21.12.00.47 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Wed, 21 Oct 2015 12:00:48 -0700 (PDT) Subject: Re: [OpenACC 1/11] UNIQUE internal function To: GCC Patches References: <5627DD78.9040302@acm.org> Cc: Jakub Jelinek , Bernd Schmidt , Jason Merrill , "Joseph S. Myers" From: Nathan Sidwell Message-ID: <5627E0DF.9050507@acm.org> Date: Wed, 21 Oct 2015 15:00:47 -0400 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <5627DD78.9040302@acm.org> This patch implements a new internal function that has a 'uniqueness' property. Jump-threading cannot clone it and tail-merging cannot combine multiple instances. The uniqueness is implemented by a new gimple fn, gimple_call_internal_unique_p. Routines that check for identical or cloneable calls are augmented to check this property. These are: * tree-ssa-threadedge, which is figuring out if jump threading is a win. Jump threading is inhibited. * gimple_call_same_target_p, used for tail merging and similar transforms. Two calls of IFN_UNIQUE will never be the same target. * tracer.c, which is determining whether to clone a region. Interestingly jump threading avoids cloning volatile asms (which it admits is conservatively safe), but the tracer does not. I wonder if there's a latent problem in tracer? The reason I needed a function with this property is to preserve the looping structure of a function's CFG. As mentioned in the intro, we mark up loops (using this builtin), so the example I gave has the following inserts: #pragma acc parallel ... { // single mode here #pragma acc loop ... IFN_UNIQUE (FORKING ...) for (i = 0; i < N; i++) // loop 1 ... // partitioned mode here IFN_UNIQUE (JOINING ...) if (expr) // single mode here #pragma acc loop ... IFN_UNIQUE (FORKING ...) for (i = 0; i < N; i++) // loop 2 ... // partitioned mode here IFN_UNIQUE (JOINING ...) } The properly nested loop property of the CFG is preserved through the compilation. This is important as (a) it allows later passes to reconstruct this looping structure and (b) hardware constraints require a partioned region end for all partitioned threads at a single instruction. Until I added this unique property, original bring-up of partitioned execution would hit cases of split loops ending in multiple cloned JOINING markers and similar cases. To distinguish different uses of the UNIQUE function, I use the first argument, which is expected to be an INTEGER_CST. I figured this better than using multiple new internal fns, all with the unique property, as the latter would need (at least) a range check in gimple_call_internal_unique_p rather than a simple equality. Jakub, IYR I originally had IFN_FORK and IFN_JOIN as such distinct internal fns. This replaces that scheme. ok? nathan 2015-10-20 Nathan Sidwell Cesar Philippidis * internal-fn.c (expand_UNIQUE): New. * internal-fn.def (IFN_UNIQUE): New. (IFN_UNIQUE_UNSPEC): Define. * gimple.h (gimple_call_internal_unique_p): New. * gimple.c (gimple_call_same_target_p): Check internal fn uniqueness. * tracer.c (ignore_bb_p): Check for IFN_UNIQUE call. * tree-ssa-threadedge.c (record_temporary_equivalences_from_stmts): Likewise. Index: gimple.c =================================================================== --- gimple.c (revision 229096) +++ gimple.c (working copy) @@ -1346,7 +1346,8 @@ gimple_call_same_target_p (const gimple { if (gimple_call_internal_p (c1)) return (gimple_call_internal_p (c2) - && gimple_call_internal_fn (c1) == gimple_call_internal_fn (c2)); + && gimple_call_internal_fn (c1) == gimple_call_internal_fn (c2) + && !gimple_call_internal_unique_p (as_a (c1))); else return (gimple_call_fn (c1) == gimple_call_fn (c2) || (gimple_call_fndecl (c1) Index: gimple.h =================================================================== --- gimple.h (revision 229096) +++ gimple.h (working copy) @@ -2895,6 +2895,14 @@ gimple_call_internal_fn (const gimple *g return gimple_call_internal_fn (gc); } +/* Return true, if this internal gimple call is unique. */ + +static inline bool +gimple_call_internal_unique_p (const gcall *gs) +{ + return gimple_call_internal_fn (gs) == IFN_UNIQUE; +} + /* If CTRL_ALTERING_P is true, mark GIMPLE_CALL S to be a stmt that could alter control flow. */ Index: internal-fn.c =================================================================== --- internal-fn.c (revision 229096) +++ internal-fn.c (working copy) @@ -1958,6 +1958,30 @@ expand_VA_ARG (gcall *stmt ATTRIBUTE_UNU gcc_unreachable (); } +/* Expand the IFN_UNIQUE function according to its first argument. */ + +static void +expand_UNIQUE (gcall *stmt) +{ + rtx pattern = NULL_RTX; + + switch (TREE_INT_CST_LOW (gimple_call_arg (stmt, 0))) + { + default: + gcc_unreachable (); + break; + + case IFN_UNIQUE_UNSPEC: +#ifdef HAVE_unique + pattern = gen_unique (); +#endif + break; + } + + if (pattern) + emit_insn (pattern); +} + /* Routines to expand each internal function, indexed by function number. Each routine has the prototype: Index: internal-fn.def =================================================================== --- internal-fn.def (revision 229096) +++ internal-fn.def (working copy) @@ -65,3 +65,11 @@ DEF_INTERNAL_FN (SUB_OVERFLOW, ECF_CONST DEF_INTERNAL_FN (MUL_OVERFLOW, ECF_CONST | ECF_LEAF | ECF_NOTHROW, NULL) DEF_INTERNAL_FN (TSAN_FUNC_EXIT, ECF_NOVOPS | ECF_LEAF | ECF_NOTHROW, NULL) DEF_INTERNAL_FN (VA_ARG, ECF_NOTHROW | ECF_LEAF, NULL) + +/* An unduplicable, uncombinable function. Generally used to preserve + a CFG property in the face of jump threading, tail merging or + other such optimizations. The first argument distinguishes + between uses. Other arguments are as needed for use. The return + type depends on use too. */ +DEF_INTERNAL_FN (UNIQUE, ECF_NOTHROW | ECF_LEAF, NULL) +#define IFN_UNIQUE_UNSPEC 0 /* Undifferentiated UNIQUE. */ Index: tracer.c =================================================================== --- tracer.c (revision 229096) +++ tracer.c (working copy) @@ -93,6 +93,7 @@ bb_seen_p (basic_block bb) static bool ignore_bb_p (const_basic_block bb) { + gimple_stmt_iterator gsi; gimple *g; if (bb->index < NUM_FIXED_BLOCKS) @@ -106,6 +107,17 @@ ignore_bb_p (const_basic_block bb) if (g && gimple_code (g) == GIMPLE_TRANSACTION) return true; + /* Ignore blocks containing non-clonable function calls. */ + for (gsi = gsi_start_bb (CONST_CAST_BB (bb)); + !gsi_end_p (gsi); gsi_next (&gsi)) + { + g = gsi_stmt (gsi); + + if (is_gimple_call (g) && gimple_call_internal_p (g) + && gimple_call_internal_unique_p (as_a (g))) + return true; + } + return false; } Index: tree-ssa-threadedge.c =================================================================== --- tree-ssa-threadedge.c (revision 229096) +++ tree-ssa-threadedge.c (working copy) @@ -283,6 +283,17 @@ record_temporary_equivalences_from_stmts && gimple_asm_volatile_p (as_a (stmt))) return NULL; + /* If the statement is a unique builtin, we can not thread + through here. */ + if (gimple_code (stmt) == GIMPLE_CALL) + { + gcall *call = as_a (stmt); + + if (gimple_call_internal_p (call) + && gimple_call_internal_unique_p (call)) + return NULL; + } + /* If duplicating this block is going to cause too much code expansion, then do not thread through this block. */ stmt_count++;