From patchwork Thu Dec 8 23:05:10 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Cesar Philippidis X-Patchwork-Id: 87353 Delivered-To: patch@linaro.org Received: by 10.140.20.101 with SMTP id 92csp24834qgi; Thu, 8 Dec 2016 15:05:41 -0800 (PST) X-Received: by 10.98.152.212 with SMTP id d81mr78300356pfk.12.1481238340928; Thu, 08 Dec 2016 15:05:40 -0800 (PST) Return-Path: Received: from sourceware.org (server1.sourceware.org. [209.132.180.131]) by mx.google.com with ESMTPS id z102si30611337plh.253.2016.12.08.15.05.40 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Thu, 08 Dec 2016 15:05:40 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-return-443841-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-443841-patch=linaro.org@gcc.gnu.org designates 209.132.180.131 as permitted sender) smtp.mailfrom=gcc-patches-return-443841-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:from :subject:to:message-id:date:mime-version:content-type; q=dns; s= default; b=PK1Pl+dBZR9zIGwFMeiS+NodLOf58fpnmuIHGd16MUFqdTZWr0WIw Dxn2y95oDOYfKNoaYvD2Rl9SVAfruwlAEJJ3je3lwdQgSw1/5C3x8voAl900VHnZ 1m7py9THlV+0jSpHz1GCJ4HRxtZAzbGyORaob3b1zHCntqLgdTYSns= 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:message-id:date:mime-version:content-type; s= default; bh=BHJJhoQnNAIpsboMLZhWe8hXC/o=; b=GTo+RwKXiUxcnQIHCARt eyYgh1NynPmyR9FsHWMA3qeT6jrY/lRTs1CN/KD/hhRsoZNA8WV8t9Hruw9XxiJN DqxsOGTn6q5d/lYmTV6AMiV8UoHmPIYBdpED8nLU4RI/Ujx2DvEXy4vCsx6K9Cjr XR6raeLrtmO5M9VdjLGAwKA= Received: (qmail 1793 invoked by alias); 8 Dec 2016 23:05:28 -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 1771 invoked by uid 89); 8 Dec 2016 23:05:27 -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=classified, gang, acc, Cesar X-Spam-User: qpsmtpd, 2 recipients 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; Thu, 08 Dec 2016 23:05:17 +0000 Received: from svr-orw-mbx-01.mgc.mentorg.com ([147.34.90.201]) by relay1.mentorg.com with esmtp id 1cF7kX-00015C-Fc from Cesar_Philippidis@mentor.com ; Thu, 08 Dec 2016 15:05:13 -0800 Received: from [127.0.0.1] (147.34.91.1) by svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Thu, 8 Dec 2016 15:05:10 -0800 From: Cesar Philippidis Subject: [PATCH] Fix PR78027 To: "gcc-patches@gcc.gnu.org" , , , Fortran List Message-ID: Date: Thu, 8 Dec 2016 15:05:10 -0800 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:45.0) Gecko/20100101 Thunderbird/45.5.1 MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-04.mgc.mentorg.com (147.34.90.204) To svr-orw-mbx-01.mgc.mentorg.com (147.34.90.201) PR78027 was classified as a fortran bug, but the underlying problem turned out to be more generic. Basically, the IPA-ICF pass usually ignores functions with omp decl attributes as candidates for aliasing. Usually that works for OpenACC, because offloaded functions generated by OpenACC PARALLEL or KERNELS regions usually have an 'omp target entrypoint' attribute. However that is not the case in two situations: 1. ACC ROUTINES do not have an 'omp target entrypoint' attribute. 2. Nested offloaded regions inside 'omp declare target' functions do not have 'omp target entrypoint' attributes either. The solution I chose for this problem is to teach the IPA-ICF pass to ignore function with 'oacc *' decl attributes. Arguably 2) should be a parser error when an OpenACC offloaded region is embedded within an OpenMP offloaded function. The LTO linker does report an error for nvptx targets, but not the host. With that in mind, this patch is still necessary for case 1. Is this OK for trunk? Cesar 2016-12-08 Cesar Philippidis PR fortran/78027 gcc/ * ipa-icf.c (sem_function::parse): Don't process functions with oacc decl attributes, as they may be OpenACC routines. gcc/testsuite/ * c-c++-common/goacc/acc-icf.c: New test. * gfortran.dg/goacc/pr78027.f90: New test. diff --git a/gcc/ipa-icf.c b/gcc/ipa-icf.c index 553b81e..31061e9 100644 --- a/gcc/ipa-icf.c +++ b/gcc/ipa-icf.c @@ -1689,6 +1689,10 @@ sem_function::parse (cgraph_node *node, bitmap_obstack *stack) if (lookup_attribute_by_prefix ("omp ", DECL_ATTRIBUTES (node->decl)) != NULL) return NULL; + if (lookup_attribute_by_prefix ("oacc ", + DECL_ATTRIBUTES (node->decl)) != NULL) + return NULL; + /* PR ipa/70306. */ if (DECL_STATIC_CONSTRUCTOR (node->decl) || DECL_STATIC_DESTRUCTOR (node->decl)) diff --git a/gcc/testsuite/c-c++-common/goacc/acc-icf.c b/gcc/testsuite/c-c++-common/goacc/acc-icf.c new file mode 100644 index 0000000..ecfe3f2 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/acc-icf.c @@ -0,0 +1,49 @@ +/* Ensure that IPA-ICF is disabled on OpenACC routines. */ + +/* { dg-additional-options "-fopenacc -O2 -fdump-ipa-icf" } */ + +#pragma acc routine gang +int +routine1 (int n) +{ + int i; + + #pragma acc loop + for (i = 0; i < n; i++) + ; + + return n + 1; +} + +#pragma acc routine gang +int +routine2 (int n) +{ + int i; + + #pragma acc loop + for (i = 0; i < n; i++) + ; + + return n + 1; +} + +int +main () +{ + int i; + + #pragma acc parallel loop + for (i = 0; i < 8; i++) + ; + + #pragma acc parallel loop + for (i = 0; i < 8; i++) + ; + + return 0; +} + +/* { dg-final { scan-ipa-dump-times "Not parsed function:" 4 "icf" } } */ +/* { dg-final { scan-ipa-dump "Parsed function:main" "icf" } } */ + diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78027.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90 new file mode 100644 index 0000000..db65063 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr78027.f90 @@ -0,0 +1,20 @@ +! { dg-additional-options "-fopenmp -O2 -fdump-ipa-icf" } + +real function f() + !$omp declare target(f) + f = 1. + !$acc parallel + !$acc loop + do i = 1, 8 + end do + !$acc end parallel + !$acc parallel + !$acc loop + do i = 1, 8 + end do + !$acc end parallel + end + +! { dg-final { scan-ipa-dump "Not parsed function:f_._omp_fn.1" "icf" } } +! { dg-final { scan-ipa-dump "Not parsed function:f_._omp_fn.0" "icf" } } +! { dg-final { scan-ipa-dump "Not parsed function:f_" "icf" } }