From patchwork Fri Jan 8 14:59:01 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 564901 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 6A3BC140B0F for ; Sat, 9 Jan 2016 01:59:16 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=bhajznfa; 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:date :from:to:cc:subject:message-id:mime-version:content-type; q=dns; s=default; b=SzREsCnGBpWzu/78/2/VSEpTD/ReSj8mGWLolDJ7bo5ElFWwUA erH+U21bsqSh5dcfbC8o3ggmEutgbneT2kwvwQ3LZ3kU/IdolhMfSJ9sHT7mAqM+ B230PuQKByD4lmcc1CMvJ/M92yRXJaqPn9tmm2RdLfPDbIjOXRgUDkSZ4= 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:mime-version:content-type; s= default; bh=5RMbBNRqxnq5Uqk0QfAWDZiH0yc=; b=bhajznfavWAkWL9GbLzT WNNMJEjMQoGNwQMRtL+lqv06baSQ21He5zCS1maQfFpOpQsBfQa4l1mFHgElJcAQ FIUfzvPQNwtxflL0VJHAkeGe2wkuR7Fj898T2HCGnl7IUK7dU3k/tcFEjg3YmXdP I7rs+8yhU5OXA1IPO3lHhpA= Received: (qmail 13380 invoked by alias); 8 Jan 2016 14:59:09 -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 13359 invoked by uid 89); 8 Jan 2016 14:59:08 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-3.1 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_LOW, SPF_PASS, T_FILL_THIS_FORM_SHORT autolearn=ham version=3.3.2 spammy=forbidden, labels, Walk, unwanted X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (CAMELLIA256-SHA encrypted) ESMTPS; Fri, 08 Jan 2016 14:59:07 +0000 Received: from relay2.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 5A454AD53; Fri, 8 Jan 2016 14:59:01 +0000 (UTC) Date: Fri, 8 Jan 2016 15:59:01 +0100 From: Martin Jambor To: GCC Patches Cc: Jakub Jelinek , Jan Hubicka Subject: [patch] Avoid an unwanted decl re-map in copy_gimple_seq_and_replace_locals Message-ID: <20160108145901.GD3982@virgil.suse.cz> Mail-Followup-To: GCC Patches , Jakub Jelinek , Jan Hubicka MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.24 (2015-08-30) X-IsSubscribed: yes Hi, I ran into an ICE when compiling the following function on the HSA branch: foo (int n, int m, int o, int (*a)[m][o]) { int i, j, k; #pragma omp target teams distribute parallel for shared(a) firstprivate(n, m, o) private(i,j,k) for (i = 0; i < n; i++) for (j = 0; j < m; j++) for (k = 0; k < o; k++) a[i][j][k] = i + j + k; } The problem is that when I duplicate the loop with copy_gimple_seq_and_replace_locals, I get one extra re-mapping. Specifically, I feed the function this: { int i.2; #pragma omp teams shared(a) firstprivate(n) firstprivate(m) firstprivate(o) shared(m.1) shared(D.3275) shared(o.0) { #pragma omp distribute private(i.2) for (i.2 = 0; i.2 < n; i.2 = i.2 + 1) { #pragma omp parallel shared(a) firstprivate(n) firstprivate(m) firstprivate(o) private(i) private(j) private(k) shared(m.1) shared(D.3275) shared(o.0) { sizetype D.3286; long unsigned int D.3287; sizetype D.3288; sizetype D.3289; sizetype D.3290; long unsigned int D.3291; long unsigned int D.3292; int[0:D.3279][0:D.3271] * D.3293; int D.3294; int D.3295; #pragma omp for nowait for (i = 0; i < n; i = i + 1) { j = 0; goto ; : k = 0; goto ; : D.3286 = D.3275 /[ex] 4; <--- here I get wrog decl D.3287 = (long unsigned int) i; D.3288 = (sizetype) o.0; D.3289 = (sizetype) m.1; D.3290 = D.3288 * D.3289; D.3291 = D.3287 * D.3290; D.3292 = D.3291 * 4; D.3293 = a + D.3292; D.3294 = i + j; D.3295 = D.3294 + k; *D.3293[j]{lb: 0 sz: D.3286 * 4}[k] = D.3295; k = k + 1; : if (k < o) goto ; else goto ; : j = j + 1; : if (j < m) goto ; else goto ; : } } } } } and it replaces D.3275 with its new copy with undefined value. The mapping is created when an array type where the size is defined in terms of that variable declaration is copied. The comment in type-remapping code says that we "use the already remaped data" but that is not true. My solution was to prevent declaration duplication in this case with yet another state variable in struct copy_body_data that holds a special value when we are running copy_gimple_seq_and_replace_locals and another when we are within type-remapping. I'll be happy for any suggestion how to deal with this without cluttering copy_body_date even more but so far I have not found any. If nobody has a better idea, is the following good for trunk? (I am about to commit it to the hsa branch.) It has passed bootstrap and testing on x86_64-linux. Thanks, Martin 2016-01-06 Martin Jambor * tree-inline.h (copy_body_data): New field decl_creation_prevention_level. Moved remap_var_for_cilk to minimize padding. * tree-inline.c (remap_decl): Return original decls if decl_creation_prevention_level is two or bigger. (remap_type_1): Increment and decrement decl_creation_prevention_level if appropriate. (copy_gimple_seq_and_replace_locals): Set decl_creation_prevention_level to 1. diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c index 88a6753..2df11a2 100644 --- a/gcc/tree-inline.c +++ b/gcc/tree-inline.c @@ -340,8 +340,20 @@ remap_decl (tree decl, copy_body_data *id) return decl; } - /* If we didn't already have an equivalent for this declaration, - create one now. */ + /* If decl copying is forbidden (which happens when copying a type with size + defined outside of the copied sequence) work with the original decl. */ + if (!n + && id->decl_creation_prevention_level > 1 + && (VAR_P (decl) || TREE_CODE (decl) == PARM_DECL)) + { + if (id->do_not_unshare) + return decl; + else + return unshare_expr (decl); + } + + /* If we didn't already have an equivalent for this declaration, create one + now. */ if (!n) { /* Make a copy of the variable or label. */ @@ -526,7 +538,10 @@ remap_type_1 (tree type, copy_body_data *id) gcc_unreachable (); } - /* All variants of type share the same size, so use the already remaped data. */ + /* All variants of type share the same size, so use the already remaped + data. */ + if (id->decl_creation_prevention_level > 0) + id->decl_creation_prevention_level++; if (TYPE_MAIN_VARIANT (new_tree) != new_tree) { gcc_checking_assert (TYPE_SIZE (type) == TYPE_SIZE (TYPE_MAIN_VARIANT (type))); @@ -540,6 +555,8 @@ remap_type_1 (tree type, copy_body_data *id) walk_tree (&TYPE_SIZE (new_tree), copy_tree_body_r, id, NULL); walk_tree (&TYPE_SIZE_UNIT (new_tree), copy_tree_body_r, id, NULL); } + if (id->decl_creation_prevention_level > 1) + id->decl_creation_prevention_level--; return new_tree; } @@ -5276,6 +5293,7 @@ copy_gimple_seq_and_replace_locals (gimple_seq seq) id.transform_return_to_modify = false; id.transform_parameter = false; id.transform_lang_insert_block = NULL; + id.decl_creation_prevention_level = 1; /* Walk the tree once to find local labels. */ memset (&wi, 0, sizeof (wi)); diff --git a/gcc/tree-inline.h b/gcc/tree-inline.h index b8fb2a2..f737daf 100644 --- a/gcc/tree-inline.h +++ b/gcc/tree-inline.h @@ -140,14 +140,21 @@ struct copy_body_data the originals have been mapped to a value rather than to a variable. */ hash_map *debug_map; - - /* Cilk keywords currently need to replace some variables that - ordinary nested functions do not. */ - bool remap_var_for_cilk; /* A map from the inlined functions dependence info cliques to equivalents in the function into which it is being inlined. */ hash_map *dependence_map; + + /* Cilk keywords currently need to replace some variables that + ordinary nested functions do not. */ + bool remap_var_for_cilk; + + /* When zero, do nothing. When one or higher, increment during type + remapping, When two or higher, do not create new variables when remapping + decls. Used when remapping types with variable size, but when the size is + defined outside the sequence copied by + copy_gimple_seq_and_replace_locals. */ + unsigned decl_creation_prevention_level; }; /* Weights of constructions for estimate_num_insns. */