From patchwork Tue Jan 14 15:10:08 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 310761 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 7F0192C0091 for ; Wed, 15 Jan 2014 02:11: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:date:message-id:in-reply-to:references :mime-version:content-type; q=dns; s=default; b=NYS7xYaW5Bt0XtRv dGmaLWkFE8SPZOUohzM63rj24GspI6nJM0S20XCHkWyuFtfp5RsF7T82ESp4T2ym gzBuXpTZadYOLprM4jbw2Ev6WlQyAaSdzpvniD35cEXYz709/EldSG157q5CROiy RPMkjGVORO+aBgVjxTbRdujf9ho= 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:date:message-id:in-reply-to:references :mime-version:content-type; s=default; bh=DjcEf3lDgSPy4Y+rFGBFI2 z3UpA=; b=W58WROmZdjLaX1m8SvPFju4/cckPOe4tt2nhZxQ5BWJGlxkUZ9XkUY ZcxFpCRC8ta2P3CyfDLyScNNfKEvGolQjXpSlTr0Tipvd+lz5x0bHttByVXNPkOP CjmqT4h8ODgdp6UfuxRwlI7q8rhcCNev1bL45969xaVDYubeez49M= Received: (qmail 12227 invoked by alias); 14 Jan 2014 15:10:36 -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 12158 invoked by uid 89); 14 Jan 2014 15:10:35 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.3 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (208.118.235.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-SHA encrypted) ESMTPS; Tue, 14 Jan 2014 15:10:34 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1W35dI-0006Pa-FM for gcc-patches@gcc.gnu.org; Tue, 14 Jan 2014 10:10:31 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:62397) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1W35dI-0006NA-4t for gcc-patches@gcc.gnu.org; Tue, 14 Jan 2014 10:10:24 -0500 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1W35dC-0002Fv-LE from Thomas_Schwinge@mentor.com ; Tue, 14 Jan 2014 07:10:18 -0800 Received: from SVR-ORW-FEM-02.mgc.mentorg.com ([147.34.96.206]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Tue, 14 Jan 2014 07:10:18 -0800 Received: from build5-lucid-cs (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.2.247.3; Tue, 14 Jan 2014 07:10:17 -0800 Received: by build5-lucid-cs (Postfix, from userid 49978) id CB7F1321CA1; Tue, 14 Jan 2014 07:10:16 -0800 (PST) From: To: , CC: Thomas Schwinge Subject: [gomp4 6/6] Enable initial support in the C front end for OpenACC data clauses. Date: Tue, 14 Jan 2014 16:10:08 +0100 Message-ID: <1389712208-416-6-git-send-email-thomas@codesourcery.com> 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> MIME-Version: 1.0 X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 From: Thomas Schwinge gcc/c/ * c-parser.c (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, 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/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: New file. * c-c++-common/goacc/deviceptr-1.c: New file. libgomp/ * testsuite/libgomp.oacc-c/parallel-1.c: Extend. --- gcc/c/c-parser.c | 14 +- .../c-c++-common/goacc/data-clause-duplicate-1.c | 13 ++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 64 +++++++++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c | 150 +++++++++++++++++++-- 4 files changed, 228 insertions(+), 13 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-1.c diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 48c55e6..d6a2af0 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -11225,8 +11225,17 @@ c_parser_omp_structured_block (c_parser *parser) LOC is the location of the #pragma token. */ -#define OACC_PARALLEL_CLAUSE_MASK \ - (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE) +#define OACC_PARALLEL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) ) static tree c_parser_oacc_parallel (location_t loc, c_parser *parser) @@ -11235,7 +11244,6 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser) clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, "#pragma acc parallel"); - gcc_assert (clauses == NULL); block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser)); diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c new file mode 100644 index 0000000..1bcf5be --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -0,0 +1,13 @@ +void +fun (void) +{ + float *fp; +#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel create(fp[:10]) deviceptr(fp) + /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */ + /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */ + ; +} diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c new file mode 100644 index 0000000..0f0cf0c --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -0,0 +1,64 @@ +void +fun1 (void) +{ +#pragma acc parallel deviceptr(u) /* { dg-error "'u' undeclared" } */ + ; +#pragma acc parallel deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + ; + +#pragma acc parallel deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */ + ; +#pragma acc parallel deviceptr(fun1[2:5]) + /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 11 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 11 } */ + ; + + int i; +#pragma acc parallel deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */ + ; +#pragma acc parallel deviceptr(i[0:4]) + /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 19 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 19 } */ + ; + + float fa[10]; +#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */ + ; +#pragma acc parallel deviceptr(fa[1:5]) + /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 27 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 27 } */ + ; + + float *fp; +#pragma acc parallel deviceptr(fp) + ; +#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + ; +} + +void +fun2 (void) +{ + int i; + float *fp; +#pragma acc parallel deviceptr(fp,u,fun2,i,fp) + /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 44 } */ + /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 44 } */ + /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 44 } */ + /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 44 } */ + ; +} + +void +fun3 (void) +{ + float *fp; +#pragma acc parallel deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +} + +/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */ diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c index b40545d..ff54b9d 100644 --- libgomp/testsuite/libgomp.oacc-c/parallel-1.c +++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c @@ -2,25 +2,155 @@ extern void abort (); -volatile int i; +int i; int main(void) { - volatile int j; + int j, v; - i = -0x42; - j = -42; -#pragma acc parallel +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) copyin (i, j) { - if (i != -0x42 || j != -42) + if (i != -1 || j != -2) abort (); - i = 42; - j = 0x42; - if (i != 42 || j != 0x42) + i = 2; + j = 1; + if (i != 2 || j != 1) abort (); + v = 1; } - if (i != 42 || j != 0x42) + if (v != 1 || i != -1 || j != -2) abort (); + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) copyout (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) copy (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) create (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != -1 || j != -2) + abort (); +#endif + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyin (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != -1 || j != -2) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyout (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copy (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); +#endif + +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); +#endif + return 0; }