From patchwork Tue Sep 29 09:49:57 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 523762 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 60D1B14029C for ; Tue, 29 Sep 2015 19:50:41 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=H6jOX/+u; 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:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=CM9jBSEWrNJ9z3/WOZuSYYgTiS2SveRqYirMQhOQbLqC0hZuyJ 9JmtSxXVPAazp3fNbdzmHWTy3rEI0JiMvLX9YSe/JqRDtHw6VQyS2+Af3DinkfJr fAdDE78KrBOkbbesdEVOkRkAfRaUC2/7cJcxMqraNVw/JsF/6oTmvJC50= 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=BYuXoEE8YXVMnd5YYztdotw42Pg=; b=H6jOX/+uQSMFFzzONdOk tPg2iZAHnBTHW6s/Cy8caJMtqsz1IJ7ZE99F4KMRnvf6kCN7Nf0pYwG6jnp6FHpE ZBIjpZGTuqQBa/SXU8LGXhzz3QKRXbPMHP/QOIXkucBY9xUCtmrRjUzYnNC99bRn BVZUmnfjgBtpXnno7T07KUA= Received: (qmail 110131 invoked by alias); 29 Sep 2015 09:50:34 -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 110119 invoked by uid 89); 29 Sep 2015 09:50:33 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00, SPF_PASS, T_RP_MATCHES_RCVD autolearn=ham version=3.3.2 X-HELO: fencepost.gnu.org Received: from fencepost.gnu.org (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Tue, 29 Sep 2015 09:50:32 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:45362) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1ZgrYL-0002dQ-VY for gcc-patches@gnu.org; Tue, 29 Sep 2015 05:50:30 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1ZgrYI-0007KH-6O for gcc-patches@gnu.org; Tue, 29 Sep 2015 05:50:29 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:45274) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1ZgrYI-0007J8-1n for gcc-patches@gnu.org; Tue, 29 Sep 2015 05:50:26 -0400 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1ZgrYF-0004uN-Ee from Tom_deVries@mentor.com for gcc-patches@gnu.org; Tue, 29 Sep 2015 02:50:23 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Tue, 29 Sep 2015 10:50:21 +0100 To: "gcc-patches@gnu.org" From: Tom de Vries Subject: [gomp4, committed] Ignore reduction clauses in kernels region Message-ID: <560A5EC5.5060009@mentor.com> Date: Tue, 29 Sep 2015 11:49:57 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.2.0 MIME-Version: 1.0 X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 Hi, this patch filters out reduction clauses in an oacc kernels region. This fixes an ICE in the test-case. Committed to gomp-4_0-branch. Thanks, - Tom Ignore reduction clauses in kernels region 2015-09-29 Tom de Vries * omp-low.c (ctx_in_oacc_kernels_region): New function. (scan_omp_for): Filter out reduction clauses in kernels region. * c-c++-common/goacc/kernels-acc-loop-reduction.c: New test. --- gcc/omp-low.c | 18 +++++++++++++++- .../goacc/kernels-acc-loop-reduction.c | 25 ++++++++++++++++++++++ 2 files changed, 42 insertions(+), 1 deletion(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-acc-loop-reduction.c diff --git a/gcc/omp-low.c b/gcc/omp-low.c index a5904eb..597035f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2579,6 +2579,20 @@ oacc_loop_or_target_p (gimple *stmt) && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP)); } +bool +ctx_in_oacc_kernels_region (omp_context *ctx) +{ + for (;ctx != NULL; ctx = ctx->outer) + { + gimple *stmt = ctx->stmt; + if (gimple_code (stmt) == GIMPLE_OMP_TARGET + && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS) + return true; + } + + return false; +} + /* Scan a GIMPLE_OMP_FOR. */ static void @@ -2592,6 +2606,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) bool auto_clause = false; bool seq_clause = false; int gwv_routine = 0; + bool in_oacc_kernels_region = ctx_in_oacc_kernels_region (outer_ctx); if (outer_ctx) outer_type = gimple_code (outer_ctx->stmt); @@ -2665,7 +2680,8 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) /* Filter out any OpenACC clauses which aren't associated with gangs, workers or vectors. Such reductions are no-ops. */ - if (extract_oacc_loop_mask (ctx) == 0) + if (extract_oacc_loop_mask (ctx) == 0 + || in_oacc_kernels_region) { /* First filter out the clauses at the beginning of the chain. */ while (clauses && OMP_CLAUSE_CODE (clauses) == OMP_CLAUSE_REDUCTION) diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-acc-loop-reduction.c b/gcc/testsuite/c-c++-common/goacc/kernels-acc-loop-reduction.c new file mode 100644 index 0000000..f3aa4e7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-loop-reduction.c @@ -0,0 +1,25 @@ +/* { dg-additional-options "-O2" } */ +/* { dg-additional-options "-ftree-parallelize-loops=32" } */ +/* { dg-additional-options "-fdump-tree-parloops_oacc_kernels-all" } */ +/* { dg-additional-options "-fdump-tree-optimized" } */ + +unsigned int +foo (int n, unsigned int *a) +{ + unsigned int sum = 0; + +#pragma acc kernels loop gang reduction(+:sum) + for (int i = 0; i < n; i++) + sum += a[i]; + + return sum; +} + +/* Check that only one loop is analyzed, and that it can be parallelized. */ +/* { dg-final { scan-tree-dump-times "SUCCESS: may be parallelized" 1 "parloops_oacc_kernels" } } */ +/* { dg-final { scan-tree-dump-not "FAILED:" "parloops_oacc_kernels" } } */ + +/* Check that the loop has been split off into a function. */ +/* { dg-final { scan-tree-dump-times "(?n);; Function .*foo.*\\._omp_fn\\.0" 1 "optimized" } } */ + +/* { dg-final { scan-tree-dump-times "(?n)pragma omp target oacc_parallel.*num_gangs\\(32\\)" 1 "parloops_oacc_kernels" } } */ -- 1.9.1