From patchwork Wed Feb 12 11:16:44 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 319616 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 DDF9F2C00B3 for ; Wed, 12 Feb 2014 22:17:16 +1100 (EST) 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:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; q=dns; s=default; b=dkrSSBsLG0s/uX9s Awk479KavIgGGBIc0ad0uVxF8njRjOnY6nWgXOUlloaU5qdvqzvWWynBQpnHq3cc q0R6TAwX7vJGD3jTBYO/J1AdSxINwBzQHP4ttyQpCybSTI8jE7b0Z5swrSaOd9tU POkG2bi+SdyWPRIFHKL/KZSZiN8= 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:cc:subject:in-reply-to:references:date:message-id :mime-version:content-type; s=default; bh=A426a6iyM666DASsd/FioY vo3lQ=; b=m7WhKzbM1WAk9ET+7dNrrlETmpxB1SsiiV1w2W1Uj9uwp8TIz/81Y5 8lJDTAocFoloLjtNCe9cJWJcNvheSHXks3tP43Hd/EWs1++IvNxuRu2ypTktnfan YRq8MyTlRV1lj7Q6hZZIhtkh3JQqb90c5Irjs2jO/88Pc8Np/9vNU= Received: (qmail 32229 invoked by alias); 12 Feb 2014 11:17:09 -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 32214 invoked by uid 89); 12 Feb 2014 11:17:08 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.7 required=5.0 tests=AWL, BAYES_00, UNWANTED_LANGUAGE_BODY autolearn=ham version=3.3.2 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; Wed, 12 Feb 2014 11:17:07 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1WDXoM-0006lg-Ez from Thomas_Schwinge@mentor.com ; Wed, 12 Feb 2014 03:17:02 -0800 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Wed, 12 Feb 2014 03:17:02 -0800 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.2.247.3; Wed, 12 Feb 2014 11:17:00 +0000 From: Thomas Schwinge To: , Ilmir Usmanov CC: Subject: Re: [gomp4 5/6] Initial support in the C front end for OpenACC data clauses. In-Reply-To: <1389712208-416-5-git-send-email-thomas@codesourcery.com> References: <87ppnuvbv6.fsf@schwinge.name> <1389712208-416-1-git-send-email-thomas@codesourcery.com> <1389712208-416-2-git-send-email-thomas@codesourcery.com> <1389712208-416-3-git-send-email-thomas@codesourcery.com> <1389712208-416-4-git-send-email-thomas@codesourcery.com> <1389712208-416-5-git-send-email-thomas@codesourcery.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (i486-pc-linux-gnu) Date: Wed, 12 Feb 2014 12:16:44 +0100 Message-ID: <87a9dwy40z.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! On Tue, 14 Jan 2014 16:10:07 +0100, I wrote: > gcc/c-family/ > * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY, > PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, > PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEVICEPTR, > PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and > PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. > gcc/c/ > * c-parser.c (c_parser_omp_clause_name): Handle these. > (c_parser_oacc_data_clause, c_parser_oacc_data_clause_deviceptr): > New functions. > (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_COPY, > PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT, > PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DELETE, > PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and > PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. > gcc/ > * tree-core.h (omp_clause_code): Update description for > OMP_CLAUSE_MAP. This I committed to gomp-4_0-branch as r207177. In , Ilmir mentioned that I'm missing to handle the »short names: pcopy, pcopyin, pcopyout and pcreate (see 2.6.5.9 - 2.6.5.12 of OpenACC 2.0«. Unless there are any comments, I'll soon commit the following to gomp-4_0-branch: commit 9a1f6c075f6198c9ae3281387b875e6012e4387e Author: Thomas Schwinge Date: Wed Feb 12 11:59:51 2014 +0100 OpenACC: pcopy, pcopyin, pcopyout, pcreate clauses. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin, pcopyout, pcreate clauses. (c_parser_oacc_data_clause): Update comment. gcc/ * tree-core.h (omp_clause_code) : Mention pcopy, pcopyin, pcopyout, pcreate OpenACC clauses. gcc/testsuite/ * c-c++-common/goacc/pcopy.c: New file. * c-c++-common/goacc/pcopyin.c: Likewise. * c-c++-common/goacc/pcopyout.c: Likewise. * c-c++-common/goacc/pcreate.c: Likewise. Grüße, Thomas diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 6e89471..f401cef 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -9671,13 +9671,17 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_PARALLEL; else if (!strcmp ("present", p)) result = PRAGMA_OMP_CLAUSE_PRESENT; - else if (!strcmp ("present_or_copy", p)) + else if (!strcmp ("present_or_copy", p) + || !strcmp ("pcopy", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY; - else if (!strcmp ("present_or_copyin", p)) + else if (!strcmp ("present_or_copyin", p) + || !strcmp ("pcopyin", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN; - else if (!strcmp ("present_or_copyout", p)) + else if (!strcmp ("present_or_copyout", p) + || !strcmp ("pcopyout", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT; - else if (!strcmp ("present_or_create", p)) + else if (!strcmp ("present_or_create", p) + || !strcmp ("pcreate", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE; else if (!strcmp ("private", p)) result = PRAGMA_OMP_CLAUSE_PRIVATE; @@ -9870,9 +9874,13 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, delete ( variable-list ) present ( variable-list ) present_or_copy ( variable-list ) + pcopy ( variable-list ) present_or_copyin ( variable-list ) + pcopyin ( variable-list ) present_or_copyout ( variable-list ) - present_or_create ( variable-list ) */ + pcopyout ( variable-list ) + present_or_create ( variable-list ) + pcreate ( variable-list ) */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, diff --git gcc/testsuite/c-c++-common/goacc/pcopy.c gcc/testsuite/c-c++-common/goacc/pcopy.c new file mode 100644 index 0000000..fd16525 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcopy.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcopy(cp[3:5]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/testsuite/c-c++-common/goacc/pcopyin.c gcc/testsuite/c-c++-common/goacc/pcopyin.c new file mode 100644 index 0000000..c009d24 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcopyin.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcopyin(cp[4:6]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/testsuite/c-c++-common/goacc/pcopyout.c gcc/testsuite/c-c++-common/goacc/pcopyout.c new file mode 100644 index 0000000..6099eff --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcopyout.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcopyout(cp[5:7]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/testsuite/c-c++-common/goacc/pcreate.c gcc/testsuite/c-c++-common/goacc/pcreate.c new file mode 100644 index 0000000..2f6e836 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcreate.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcreate(cp[6:8]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/tree-core.h gcc/tree-core.h index a5a95cd..2d9bf0c 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -259,8 +259,9 @@ enum omp_clause_code { OMP_CLAUSE_TO, /* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr, - present, present_or_copy, present_or_copyin, present_or_copyout, - present_or_create} (variable-list). + present, present_or_copy (pcopy), present_or_copyin (pcopyin), + present_or_copyout (pcopyout), present_or_create (pcreate)} + (variable-list). OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP,