From patchwork Thu Mar 24 16:59:49 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 601729 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 3qWCNy0TvTz9sBG for ; Fri, 25 Mar 2016 04:00:33 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=DiWmHnTH; 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:cc:message-id:date:mime-version:content-type; q= dns; s=default; b=PSPc23MoTA3s43Ru/oQXkfxpIzHfOQlva+jEW0rNf6tVDm VtKo38HYqTOupWj/ytOkQuPEOqgB58L0y/wGIfdm5k+v1rHvDnqkzfdjEojYcEdZ DxhpUs32IyR2Ed2YkgryEQ/86D5U+QqgR8e65PClaC8D8z/Ak/3Z0cCobBNRI= 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:cc:message-id:date:mime-version:content-type; s= default; bh=uha0fqPMvhpXPTqj2i4rWayIoQQ=; b=DiWmHnTHMA13FIsLiJWO r/VSfdFh3nrvEbDs9efvRSTa7YJUyVrzd59R9Wy2Rboe/FwnX9Wwgidkt4LBH/kB l/IeJ3z6n2BNnFOvdCPAEHI2usmMy20+r2HU3Xza5YyE36Cfdu6zKblxqyBI2d5O nH+KW0QzMFWMN8ammSN9Gug= Received: (qmail 75921 invoked by alias); 24 Mar 2016 17:00:21 -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 75726 invoked by uid 89); 24 Mar 2016 17:00:11 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.0 required=5.0 tests=AWL, BAYES_00, RCVD_IN_DNSWL_NONE, SPF_PASS, UNSUBSCRIBE_BODY autolearn=no version=3.3.2 spammy=stmt 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 (AES256-GCM-SHA384 encrypted) ESMTPS; Thu, 24 Mar 2016 17:00:00 +0000 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 1aj8c2-00042B-Ls from Tom_deVries@mentor.com ; Thu, 24 Mar 2016 09:59:59 -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; Thu, 24 Mar 2016 16:59:56 +0000 To: Jakub Jelinek From: Tom de Vries Subject: [PATCH] Remove incorrect warning for kernels copy clause CC: GCC Patches Message-ID: <56F41D05.7040700@mentor.com> Date: Thu, 24 Mar 2016 17:59:49 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.6.0 MIME-Version: 1.0 Hi, This patch fixes an incorrect warning for the oacc copy clause. Consider this test-case: ... void foo (void) { int i; #pragma acc kernels { i = 1; } } ... When compiling with -fopenacc -Wuninitialized, we get an 'is used uninitialized' warning for variable 'i', which is confusing given that 'i' is not used, but only set in the kernels region. The warning occurs because there's an implicit copy(i) clause on the kernels region, and that copy generates a read of i before the region, and a write to i in region. The patch silences the warning by marking the variable in the copy clause with TREE_NO_WARNING. Build and reg-tested with goacc.exp, gomp.exp and target-libgomp. OK for trunk if bootstrap and reg-test succeeds? Thanks, - Tom Remove incorrect warning for kernels copy clause 2016-03-24 Tom de Vries * omp-low.c (lower_omp_target): Set TREE_NO_WARNING for oacc copy clause. * c-c++-common/goacc/uninit-copy-clause.c: New test. * gfortran.dg/goacc/uninit-copy-clause.f95: New test. --- gcc/omp-low.c | 6 +++- .../c-c++-common/goacc/uninit-copy-clause.c | 38 ++++++++++++++++++++++ .../gfortran.dg/goacc/uninit-copy-clause.f95 | 29 +++++++++++++++++ 3 files changed, 72 insertions(+), 1 deletion(-) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 3fd6eb3..d107961 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -16083,7 +16083,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) || map_kind == GOMP_MAP_POINTER || map_kind == GOMP_MAP_TO_PSET || map_kind == GOMP_MAP_FORCE_DEVICEPTR) - gimplify_assign (avar, var, &ilist); + { + if (is_gimple_omp_oacc (ctx->stmt)) + TREE_NO_WARNING (var) = 1; + gimplify_assign (avar, var, &ilist); + } avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); if ((GOMP_MAP_COPY_FROM_P (map_kind) diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c b/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c new file mode 100644 index 0000000..b3cc445 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/uninit-copy-clause.c @@ -0,0 +1,38 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-Wuninitialized" } */ + +void +foo (void) +{ + int i; + +#pragma acc kernels + { + i = 1; + } + +} + +void +foo2 (void) +{ + int i; + +#pragma acc kernels copy (i) + { + i = 1; + } + +} + +void +foo3 (void) +{ + int i; + +#pragma acc kernels copyin(i) + { + i = 1; + } + +} diff --git a/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 b/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 new file mode 100644 index 0000000..b2aae1d --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/uninit-copy-clause.f95 @@ -0,0 +1,29 @@ +! { dg-do compile } +! { dg-additional-options "-Wuninitialized" } + +subroutine foo + integer :: i + + !$acc kernels + i = 1 + !$acc end kernels + +end subroutine foo + +subroutine foo2 + integer :: i + + !$acc kernels copy (i) + i = 1 + !$acc end kernels + +end subroutine foo2 + +subroutine foo3 + integer :: i + + !$acc kernels copyin (i) + i = 1 + !$acc end kernels + +end subroutine foo3