From patchwork Fri May 5 08:14:54 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 758898 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 3wK4VM3KSXz9s84 for ; Fri, 5 May 2017 18:16:22 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="tWPJxyr5"; 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:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; q=dns; s=default; b=ndV KAT4WznJrMAWN/l2kacAUiF7gihcL2mstEzrj+MANP3UA4w3qf3T+ajOq3VqJlxt g0UCN4QVZxohesGtrk9n4J8HHlr6AI4hkeLcnhsKZEkolhuBV8ecnNXJFvu+d41k mEDkGtfD/dBTrmD+MShqcuhMmH+yEXCm5TVGTJho= 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 :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; s=default; bh=5uMRm818T G8jWgCHosf9dAhevc0=; b=tWPJxyr5koEJ/gZJk+BH9EQ07K0Iz7/RrvVPw4K75 5xbA0QFT/K+wTURYElcKPqLg8sQLBlEIZnhLlJvSJxPpclQAti0q+3hX3A9k7pBg 2ZUlWP47b1bsH6fxViiJYo8V6hnXXnbSWp2MvVlesAtR/l1duDSa1LscAxiIrgsy ng= Received: (qmail 89850 invoked by alias); 5 May 2017 08:15:24 -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 89601 invoked by uid 89); 5 May 2017 08:15:07 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS, URIBL_RED autolearn=ham version=3.3.2 spammy=1524 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; Fri, 05 May 2017 08:15:01 +0000 Received: from svr-orw-fem-05.mgc.mentorg.com ([147.34.97.43]) by relay1.mentorg.com with esmtp id 1d6YOC-0002gf-5E from Thomas_Schwinge@mentor.com ; Fri, 05 May 2017 01:15:00 -0700 Received: from tftp-cs (147.34.91.1) by svr-orw-fem-05.mgc.mentorg.com (147.34.97.43) with Microsoft SMTP Server id 14.3.224.2; Fri, 5 May 2017 01:14:59 -0700 Received: by tftp-cs (Postfix, from userid 49978) id 3A9FBC231E; Fri, 5 May 2017 01:14:59 -0700 (PDT) From: Thomas Schwinge To: Cesar Philippidis , "gcc-patches@gcc.gnu.org" , Fortran List Subject: Re: [gomp4] Add front end support for the if_present clause with the update directive In-Reply-To: <9f137e97-ea71-b048-cafd-15b58ff095f3@codesourcery.com> References: <9f137e97-ea71-b048-cafd-15b58ff095f3@codesourcery.com> User-Agent: Notmuch/0.9-125-g4686d11 (http://notmuchmail.org) Emacs/24.5.1 (x86_64-pc-linux-gnu) Date: Fri, 5 May 2017 10:14:54 +0200 Message-ID: <87bmr7wwht.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Thu, 4 May 2017 15:11:30 -0700, Cesar Philippidis wrote: > This patch make the c, c++ and fortran FEs aware of the new OpenACC 2.5 > if_present clause for the update directive. Thanks! > The ME and runtime support > will come in a separate followup patch. > Thomas, for some reason I'm seeing a couple of new UNRESOLVED tests for > update-1.C. ;-) You need to instantiate C++ templates for any code to be generated. > The c++ tests running with goacc.exp are built with > -fopenacc, but for some reason the tests in g++.dg/goacc/ are still ran > without -fopenacc for g++.dg/dg.exp. Maybe there's something wrong with > g++.dg/goacc/goacc.exp handling of .C files? I'm not aware of any such a problem, but please do verify. I also added missing handling in gcc/tree-nested.c, and corrsponding test coverage. Committed to gomp-4_0-branch in r247629: commit afb07db9e7f52dac835f6c0a764b8884187bc798 Author: tschwinge Date: Fri May 5 08:09:39 2017 +0000 OpenACC if_present clause fixes gcc/ * tree-nested.c (convert_nonlocal_omp_clauses) (convert_local_omp_clauses): Handle OMP_CLAUSE_IF_PRESENT. gcc/testsuite/ * g++.dg/goacc/update-1.C: Update. * gcc.dg/goacc/nested-function-1.c: Likewise. * gfortran.dg/goacc/nested-function-1.f90: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@247629 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 5 +++++ gcc/testsuite/ChangeLog.gomp | 6 ++++++ gcc/testsuite/g++.dg/goacc/update-1.C | 20 ++++++++++++++++++++ gcc/testsuite/gcc.dg/goacc/nested-function-1.c | 4 ++++ .../gfortran.dg/goacc/nested-function-1.f90 | 4 ++++ gcc/tree-nested.c | 2 ++ 6 files changed, 41 insertions(+) Grüße Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index f24c483..b914cb1 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,8 @@ +2017-05-05 Thomas Schwinge + + * tree-nested.c (convert_nonlocal_omp_clauses) + (convert_local_omp_clauses): Handle OMP_CLAUSE_IF_PRESENT. + 2017-05-04 Cesar Philippidis * gimplify.c (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_IF_PRESENT. diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 12b79e3..d1f7c84 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,9 @@ +2017-05-05 Thomas Schwinge + + * g++.dg/goacc/update-1.C: Update. + * gcc.dg/goacc/nested-function-1.c: Likewise. + * gfortran.dg/goacc/nested-function-1.f90: Likewise. + 2017-05-04 Cesar Philippidis * c-c++-common/goacc/update-if_present-1.c: New test. diff --git gcc/testsuite/g++.dg/goacc/update-1.C gcc/testsuite/g++.dg/goacc/update-1.C index 10c8020..cb4b11e 100644 --- gcc/testsuite/g++.dg/goacc/update-1.C +++ gcc/testsuite/g++.dg/goacc/update-1.C @@ -15,4 +15,24 @@ t () #pragma acc update self(c) device(b) host (a) async(10) if (a == 5) if_present } +class +C +{ +private: + bool a; + float b; + +public: + bool operator== (int x) + { + return a == x; + } +}; + +void +f () +{ + t (); +} + /* { dg-final { scan-tree-dump-times "pragma omp target oacc_update if_present" 4 "omplower" } } */ diff --git gcc/testsuite/gcc.dg/goacc/nested-function-1.c gcc/testsuite/gcc.dg/goacc/nested-function-1.c index e17c0e2..5fc2e46 100644 --- gcc/testsuite/gcc.dg/goacc/nested-function-1.c +++ gcc/testsuite/gcc.dg/goacc/nested-function-1.c @@ -25,6 +25,8 @@ int main () local_a[i] = 5; local_arg = 5; +#pragma acc update device(local_a) if_present + #pragma acc kernels loop \ gang(num:local_arg) worker(local_arg) vector(local_arg) \ wait async(local_arg) @@ -62,6 +64,8 @@ int main () nonlocal_a[i] = 5; nonlocal_arg = 5; +#pragma acc update device(nonlocal_a) if_present + #pragma acc kernels loop \ gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \ wait async(nonlocal_arg) diff --git gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 index 2fcaa40..bbb53c3 100644 --- gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 +++ gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 @@ -25,6 +25,8 @@ contains local_a (:) = 5 local_arg = 5 + !$acc update device(local_a) if_present + !$acc kernels loop & !$acc gang(num:local_arg) worker(local_arg) vector(local_arg) & !$acc wait async(local_arg) @@ -60,6 +62,8 @@ contains nonlocal_a (:) = 5 nonlocal_arg = 5 + !$acc update device(nonlocal_a) if_present + !$acc kernels loop & !$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) & !$acc wait async(nonlocal_arg) diff --git gcc/tree-nested.c gcc/tree-nested.c index b5b73d0..3ddfd65 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1202,6 +1202,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_IF_PRESENT: break; case OMP_CLAUSE_DEVICE_TYPE: @@ -1900,6 +1901,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SEQ: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_IF_PRESENT: break; case OMP_CLAUSE_DEVICE_TYPE: