From patchwork Mon May 23 11:21:57 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 625145 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 3rCx314Kfyz9t0r for ; Mon, 23 May 2016 21:22:21 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=endINsl6; 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:subject:message-id:mime-version:content-type; q=dns; s= default; b=aFz2E0vmJ8UgIk7pkBy/ABNFygh0pqoSMEfDERbger191fTtfIf2t XhzxEckEEJab0pdO9BIE7ocseDyBJ0xOOeF5i4zDgV6rCYMEemetybn1f09SjTzs IzASXmk07j3bf4cuKnPu2nrr+w6PcJblZHB1mCf85T6/Z8GzaHZ5g8= 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:subject:message-id:mime-version:content-type; s= default; bh=qWfBKNOI2+tIwXUnp6uNA88V8fg=; b=endINsl68oTyfEOlNLdy 8EohIeOCJDEOwdSMC63GBZDSdJHv5MfAQ5sEcxu5QXASsCy3Bf5C2XFrOAyE9/0i jCoh8J2oHZF2xzl/PR+BzveTRZFxcg7oE/2tKAVINqIlhPAiCKgkEkAYGRnVk0tB uH9310USV7e5HV5JbCpcZ+4= Received: (qmail 101810 invoked by alias); 23 May 2016 11:22:11 -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 99807 invoked by uid 89); 23 May 2016 11:22:10 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=BAYES_00, SPF_PASS autolearn=ham version=3.3.2 spammy=sbr 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; Mon, 23 May 2016 11:22:00 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 667E2AC61 for ; Mon, 23 May 2016 11:21:57 +0000 (UTC) Date: Mon, 23 May 2016 13:21:57 +0200 From: Martin Jambor To: GCC Patches Subject: [hsa] Avoid segfault in hsa switch expansion Message-ID: <20160523112156.GC2962@virgil.suse.cz> Mail-Followup-To: GCC Patches MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.6.0 (2016-04-01) X-IsSubscribed: yes Hi, when we expand a switch statement to a SBR HSAIL instruction, we must be careful when changing the CFG and avoid adding a new predecessor to the default-value basic block if it has PHI nodes because we do not provide values in the PHIs for the new edge. We can avoid it by splitting the edge to the default-value block and having the new edge lead to this new block, which is what the patch below does. Bootstrapped and tested on x86_64-linux on trunk and the gcc-6 branch. I'll commit it to both momentarily. Thanks, Martin 2016-05-20 Martin Jambor * hsa-gen.c (gen_hsa_insns_for_switch_stmt): Create an empty default block if a PHI node in the original one would be resized. libgomp/ * testsuite/libgomp.hsa.c/switch-sbr-2.c: New test. --- gcc/hsa-gen.c | 6 +++ libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c | 59 ++++++++++++++++++++++++++ 2 files changed, 65 insertions(+) create mode 100644 libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index 697d599..cf7d434 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -3482,6 +3482,12 @@ gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb) basic_block default_label_bb = label_to_block_fn (func, CASE_LABEL (default_label)); + if (!gimple_seq_empty_p (phi_nodes (default_label_bb))) + { + default_label_bb = split_edge (find_edge (e->dest, default_label_bb)); + hsa_init_new_bb (default_label_bb); + } + make_edge (e->src, default_label_bb, EDGE_FALSE_VALUE); hsa_cfun->m_modified_cfg = true; diff --git a/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c b/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c new file mode 100644 index 0000000..06990d1 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/switch-sbr-2.c @@ -0,0 +1,59 @@ +/* { dg-additional-options "-fno-tree-switch-conversion" } */ + +#pragma omp declare target +int +foo (unsigned a) +{ + switch (a) + { + case 1 ... 5: + return 1; + case 9 ... 11: + return a + 3; + case 12 ... 13: + return a + 3; + default: + return 44; + } +} +#pragma omp end declare target + +#define s 100 + +void __attribute__((noinline, noclone)) +verify(int *a) +{ + if (a[0] != 44) + __builtin_abort (); + + for (int i = 1; i <= 5; i++) + if (a[i] != 1) + __builtin_abort (); + + for (int i = 6; i <= 8; i++) + if (a[i] != 44) + __builtin_abort (); + + for (int i = 9; i <= 13; i++) + if (a[i] != i + 3) + __builtin_abort (); + + for (int i = 14; i < s; i++) + if (a[i] != 44) + __builtin_abort (); +} + +int main(int argc) +{ + int array[s]; +#pragma omp target + { + for (int i = 0; i < s; i++) + { + int v = foo (i); + array[i] = v; + } + } + verify (array); + return 0; +}