From patchwork Mon Jan 7 09:01:36 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 1021217 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-493507-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=suse.de Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="rOaw0UQt"; dkim-atps=neutral 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 43Y8Wk4ZrCz9sBn for ; Mon, 7 Jan 2019 20:01:22 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender :subject:from:to:cc:references:message-id:date:mime-version :in-reply-to:content-type; q=dns; s=default; b=WmwoHNCEag4Ta0Dyw DjAcY2jHTSz8/jKTk7irX1u5v4RWAyrJB36Z+Kr7s6VCLqyXma6YAkfm1DPYblsY qDhtPlOklxWserW+pYiXciY9He5Qnn7dXbwtk6Vr5BqK1WVpVCsZ4Avmdie5mqIK A5WOh3l+DPzvOi2BhRPvZcuihc= 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:from:to:cc:references:message-id:date:mime-version :in-reply-to:content-type; s=default; bh=V7qMkMsP1GjIIomPrREY8Bn 5X5I=; b=rOaw0UQtE62oe+sbSBR2yOP5ua9iZck6tnipNL8oF4SB4lZNSDMFyU4 S2qcF6P26AsWxPConlSPnhZJK4rmN69BtoZ8Fvb5nWhsiO3pXgd8tOKSjNTMi1K3 qBy8OQGj7Dlau1glQ7QT8npHog2ziFD6l3rFFzn+niY26y+yirSQ= Received: (qmail 36662 invoked by alias); 7 Jan 2019 09:01:04 -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 36266 invoked by uid 89); 7 Jan 2019 09:01:03 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy=joining X-HELO: mx1.suse.de Received: from mx2.suse.de (HELO mx1.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 07 Jan 2019 09:01:01 +0000 Received: from relay2.suse.de (unknown [195.135.220.254]) by mx1.suse.de (Postfix) with ESMTP id 9642DACBD; Mon, 7 Jan 2019 09:00:59 +0000 (UTC) Subject: [nvptx] Don't emit barriers for empty loops -- fix From: Tom de Vries To: "Schwinge, Thomas" Cc: "gcc-patches@gcc.gnu.org" References: <2ece5d7b-3675-84ab-f255-3c56a2ffd7dc@suse.de> <91b927af-d854-2865-7cbd-9a9a835ab5cc@codesourcery.com> <1394d89c-896e-f6a3-5f9a-78e98b16e85c@suse.de> Message-ID: <6ee3fb54-9204-bf67-e643-7daf6052629f@suse.de> Date: Mon, 7 Jan 2019 10:01:36 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:60.0) Gecko/20100101 Thunderbird/60.3.0 MIME-Version: 1.0 In-Reply-To: <1394d89c-896e-f6a3-5f9a-78e98b16e85c@suse.de> X-IsSubscribed: yes [ was: Re: [nvptx] vector length patch series ] On 14-12-18 20:58, Tom de Vries wrote: > 0022-nvptx-openacc-Don-t-emit-barriers-for-empty-loops.patch Committed without test-case. Thanks, - Tom [nvptx] Don't emit barriers for empty loops -- fix When compiling an empty loop: ... long long v1; #pragma acc parallel num_gangs (640) num_workers(1) vector_length (128) #pragma acc loop for (v1 = 0; v1 < 20; v1 += 2) ; ... the compiler emits two subsequent bar.syncs. This triggers some bug on my quadro m1200 (I'm assuming in the ptxas/JIT compiler) that hangs the testcase. This patch works around the bug by doing an optimization: we detect that this is an empty loop (a forked immediately followed by a joining), and don't emit the barriers. The patch does not include the test-case yet, since vector_length (128) is not yet supported at this point. 2018-12-17 Tom de Vries PR target/85381 * config/nvptx/nvptx.c (nvptx_process_pars): Don't emit barriers for empty loops. --- gcc/config/nvptx/nvptx.c | 15 +++++++++++---- 1 file changed, 11 insertions(+), 4 deletions(-) diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 2166f37b182..26c80716603 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -4636,9 +4636,12 @@ nvptx_process_pars (parallel *par) { nvptx_shared_propagate (false, is_call, par->forked_block, par->forked_insn, !worker); - bool empty = nvptx_shared_propagate (true, is_call, - par->forked_block, par->fork_insn, - !worker); + bool no_prop_p + = nvptx_shared_propagate (true, is_call, par->forked_block, + par->fork_insn, !worker); + bool empty_loop_p + = !is_call && (NEXT_INSN (par->forked_insn) + && NEXT_INSN (par->forked_insn) == par->joining_insn); rtx barrier = GEN_INT (0); int threads = 0; @@ -4648,7 +4651,11 @@ nvptx_process_pars (parallel *par) threads = nvptx_mach_vector_length (); } - if (!empty || !is_call) + if (no_prop_p && empty_loop_p) + ; + else if (no_prop_p && is_call) + ; + else { /* Insert begin and end synchronizations. */ emit_insn_before (nvptx_cta_sync (barrier, threads),