From patchwork Fri Jun 30 15:15:24 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 782893 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 3wzg8X2glMz9s81 for ; Sat, 1 Jul 2017 01:15:56 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="GmV5HK97"; 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:from :subject:to:cc:message-id:date:mime-version:content-type; q=dns; s=default; b=JzllF5AUt5PTvFK1SsySxwKpa2rj3s1rKp0lbZh380DMfU0b2E FnUhhexIxwPh1TEE/bqH5A8Y/tmsIkTDFe3vUwdKXz+5ViTW1uXSZx5ET4EMJDif dQlnoQvO76J29Oh7/3T0cdhivS8nXYAeyTsMH/wQDcKIznfdiRsOplJTM= 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:from :subject:to:cc:message-id:date:mime-version:content-type; s= default; bh=8pC7mXzTG3gyjq73NmrbsJEIg18=; b=GmV5HK97PgwI17xVC6jb mG3whGYjvxtk67oTJoZMcQEPp+2sSkKbNp1FwmtLKcupnLmaAY+Eh4XzoiEVS8YQ IFmdZGGpifdzWaMLViXYwyY8f9mCn7R0EVBLQQiDNFNVzp6HgjqyzcC/h57qexM9 fkeMJLDuc1rCiKjoEYBuCRw= Received: (qmail 41569 invoked by alias); 30 Jun 2017 15:15:45 -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 37196 invoked by uid 89); 30 Jun 2017 15:15:37 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=H*Ad:U*thomas, 10i, ary X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 30 Jun 2017 15:15:35 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1dQxdt-0003Lb-3b from Tom_deVries@mentor.com for gcc-patches@gcc.gnu.org; Fri, 30 Jun 2017 08:15:33 -0700 Received: from [127.0.0.1] (137.202.0.87) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1263.5; Fri, 30 Jun 2017 16:15:28 +0100 From: Tom de Vries Subject: [gomp4, nvptx, committed] Fix assert in nvptx_propagate_unified To: GCC Patches CC: Thomas Schwinge Message-ID: Date: Fri, 30 Jun 2017 17:15:24 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.1.1 MIME-Version: 1.0 X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) Hi, with the openacc test-case in attached patch, I ran into an assert here: ... static void nvptx_propagate_unified (rtx_insn *unified) { rtx_insn *probe = unified; rtx cond_reg = SET_DEST (PATTERN (unified)); rtx pat; /* Find the comparison. (We could skip this and simply scan to he blocks' terminating branch, if we didn't care for self checking.) */ for (;;) { probe = NEXT_INSN (probe); pat = PATTERN (probe); if (GET_CODE (pat) == SET && GET_RTX_CLASS (GET_CODE (SET_SRC (pat))) == RTX_COMPARE && XEXP (SET_SRC (pat), 0) == cond_reg) break; gcc_assert (NONJUMP_INSN_P (probe)); } ... The assert happens when processing insn 56: ... (insn 54 53 56 3 (set (reg:SI 47 [ _71 ]) (unspec:SI [ (reg:SI 36 [ _58 ]) ] UNSPEC_BR_UNIFIED)) 108 {cond_uni} (nil)) (note 56 54 57 3 NOTE_INSN_DELETED) (insn 57 56 58 3 (set (reg:BI 68) (gt:BI (reg:SI 47 [ _71 ]) (const_int 1 [0x1]))) 99 {*cmpsi} (expr_list:REG_DEAD (reg:SI 47 [ _71 ]) (nil))) ... The insn 56 was originally a '(set (reg x) (const_int 1))', but that one has been combined into insn 57 and replaced with a 'NOTE_INSN_DELETED'. So it seems reasonable for the loop to skip over this note. Fixed by making the assert condition less strict. Build on x86_64 with nvptx accelerator. Tested test-case included in the patch. Committed as trivial. Thanks, - Tom Fix assert in nvptx_propagate_unified 2017-06-30 Tom de Vries * config/nvptx/nvptx.c (nvptx_propagate_unified): Fix gcc_assert condition by allowing !INSN_P. * testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c: New test. --- gcc/config/nvptx/nvptx.c | 2 +- .../reduction-cplx-flt-2.c | 32 ++++++++++++++++++++++ 2 files changed, 33 insertions(+), 1 deletion(-) diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 4a93bba..a3abb4b 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -2300,7 +2300,7 @@ nvptx_propagate_unified (rtx_insn *unified) && GET_RTX_CLASS (GET_CODE (SET_SRC (pat))) == RTX_COMPARE && XEXP (SET_SRC (pat), 0) == cond_reg) break; - gcc_assert (NONJUMP_INSN_P (probe)); + gcc_assert (NONJUMP_INSN_P (probe) || !INSN_P (probe)); } rtx pred_reg = SET_DEST (pat); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c new file mode 100644 index 0000000..f80f38c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-cplx-flt-2.c @@ -0,0 +1,32 @@ +#include +#include +#include + +typedef float _Complex Type; + +#define N 32 + +int +main (void) +{ + Type ary[N]; + + for (int ix = 0; ix < N; ix++) + ary[ix] = 1.0 + 1.0i; + + Type tprod = 1.0; + +#pragma acc parallel vector_length(32) + { +#pragma acc loop vector reduction (*:tprod) + for (int ix = 0; ix < N; ix++) + tprod *= ary[ix]; + } + + Type expected = 65536.0; + + if (tprod != expected) + abort (); + + return 0; +}