From patchwork Fri Oct 28 16:33:09 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 79982 Delivered-To: patch@linaro.org Received: by 10.140.97.247 with SMTP id m110csp1266073qge; Fri, 28 Oct 2016 09:33:37 -0700 (PDT) X-Received: by 10.99.37.196 with SMTP id l187mr21634070pgl.91.1477672417743; Fri, 28 Oct 2016 09:33:37 -0700 (PDT) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id e7si11915589pas.335.2016.10.28.09.33.37 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Fri, 28 Oct 2016 09:33:37 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-return-439846-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) client-ip=209.132.180.131; Authentication-Results: mx.google.com; dkim=pass header.i=@gcc.gnu.org; spf=pass (google.com: domain of gcc-patches-return-439846-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-439846-patch=linaro.org@gcc.gnu.org DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=gR412XlQI4/LOZg6rEsPfTKBOL/XvShpVadA6DpMPSwY9OM1fE 5YxHRwGIHpFlOiwEyN//WaWAuNhQf4ux1Q4xOQ6o1DlO4DICHpGqz/r28Rozgfyz w/zeXWsraDTsWsytZ78dBZF2u6jRJPQlEtvvcg/kMG/Lr/9I+v7K18oH8= 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:to :from:subject:message-id:date:mime-version:content-type; s= default; bh=m5s/diTlcGK37hxGx6ryk/wzrZ8=; b=jAXpvXioTqLDIbAX2ORT jAntC9F6zFpd8uVr7ZWTVos99yyeIi9F8dyww4LrEoplP+kmPguYfzmfL7Q2JCBU /ES8EvKiHKUEROupgue0EKWxlbYD/fvLOrTwW+42hEQ2VFdnvXxXbEDfMOiZ0r9n Rp7kS260LPhjL00SIOYGyP0= Received: (qmail 70125 invoked by alias); 28 Oct 2016 16:33:26 -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 70111 invoked by uid 89); 28 Oct 2016 16:33:25 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=gang, Hx-languages-length:3085, gen_rtx_SET, const1_rtx 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, 28 Oct 2016 16:33:15 +0000 Received: from svr-orw-mbx-04.mgc.mentorg.com ([147.34.90.204]) by relay1.mentorg.com with esmtp id 1c0A5g-000466-MH from Cesar_Philippidis@mentor.com for gcc-patches@gcc.gnu.org; Fri, 28 Oct 2016 09:33:12 -0700 Received: from [127.0.0.1] (147.34.91.1) by SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Fri, 28 Oct 2016 09:33:09 -0700 To: "gcc-patches@gcc.gnu.org" From: Cesar Philippidis Subject: [gomp4] propagating conditionals in worker-vector partitioned loops Message-ID: Date: Fri, 28 Oct 2016 09:33:09 -0700 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.3.0 MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To SVR-ORW-MBX-04.mgc.mentorg.com (147.34.90.204) I've applied the patch to gomp-4_0-branch to correct an issue involving the propagation of variables used in conditional expressions to worker and vector partitioned loops. More details regarding this patch can be found here Cesar 2016-10-26 Cesar Philippidis gcc/ * config/nvptx/nvptx.c (nvptx_single): Use a single predicate for loops partitioned across both worker and vector axes. libgomp/ * testsuite/libgomp.oacc-c-c++-common/broadcast-1.c: New test. diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c index 7bf5987..4e6ed60 100644 --- a/gcc/config/nvptx/nvptx.c +++ b/gcc/config/nvptx/nvptx.c @@ -3507,11 +3507,38 @@ nvptx_single (unsigned mask, basic_block from, basic_block to) /* Insert the vector test inside the worker test. */ unsigned mode; rtx_insn *before = tail; + rtx wvpred = NULL_RTX; + bool skip_vector = false; + + /* Create a single predicate for loops containing both worker and + vectors. */ + if (cond_branch + && (GOMP_DIM_MASK (GOMP_DIM_WORKER) & mask) + && (GOMP_DIM_MASK (GOMP_DIM_VECTOR) & mask)) + { + rtx regx = gen_reg_rtx (SImode); + rtx regy = gen_reg_rtx (SImode); + rtx tmp = gen_reg_rtx (SImode); + wvpred = gen_reg_rtx (BImode); + + emit_insn_before (gen_oacc_dim_pos (regx, const1_rtx), head); + emit_insn_before (gen_oacc_dim_pos (regy, const2_rtx), head); + emit_insn_before (gen_rtx_SET (tmp, gen_rtx_IOR (SImode, regx, regy)), + head); + emit_insn_before (gen_rtx_SET (wvpred, gen_rtx_NE (BImode, tmp, + const0_rtx)), + head); + + skip_mask &= ~(GOMP_DIM_MASK (GOMP_DIM_VECTOR)); + skip_vector = true; + } + for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++) if (GOMP_DIM_MASK (mode) & skip_mask) { rtx_code_label *label = gen_label_rtx (); - rtx pred = cfun->machine->axis_predicate[mode - GOMP_DIM_WORKER]; + rtx pred = skip_vector ? wvpred + : cfun->machine->axis_predicate[mode - GOMP_DIM_WORKER]; if (!pred) { diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c new file mode 100644 index 0000000..4dcb60d --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-1.c @@ -0,0 +1,49 @@ +/* Ensure that worker-vector state conditional expressions are + properly handled by the nvptx backend. */ + +#include +#include + + +#define N 1024 + +int A[N][N] ; + +void test(int x) +{ +#pragma acc parallel num_gangs(16) num_workers(4) vector_length(32) copyout(A) + { +#pragma acc loop gang + for(int j=0;j