[gomp4,3/3] OpenACC data construct support in the C front end.

Message ID	1393014736-19719-3-git-send-email-thomas@codesourcery.com
State	New
Headers	show Return-Path: <gcc-patches-return-362219-incoming=patchwork.ozlabs.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 :to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-type; q=dns; s=default; b=Cr1JN7bj3IwgtmxN Md9wHU+oIvEwChLV1S0IJ4dSWSmEinEhbdoTX0t4uRfkcOkggTRugQZO+hTXDYZK RO1D1jxUGJ49WLGy03kDk/3gYkkHi0ic9k1wmUTa3gLRYJj3mbt+6swHeFPz1DO8 c3gAgezddK/j0RiNf2yxU7nBU5Q= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org From: Thomas Schwinge <thomas@codesourcery.com> To: <gcc-patches@gcc.gnu.org> CC: <jakub@redhat.com> Subject: [gomp4 3/3] OpenACC data construct support in the C front end. Date: Fri, 21 Feb 2014 21:32:15 +0100 Message-ID: <1393014736-19719-3-git-send-email-thomas@codesourcery.com> In-Reply-To: <1393014736-19719-2-git-send-email-thomas@codesourcery.com> References: <877g8os0vx.fsf@kepler.schwinge.homeip.net> <1393014736-19719-1-git-send-email-thomas@codesourcery.com> <1393014736-19719-2-git-send-email-thomas@codesourcery.com> MIME-Version: 1.0 Content-Type: text/plain

diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp index e092d53..3da377f 100644 --- gcc/c-family/ChangeLog.gomp +++ gcc/c-family/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * c-pragma.c (oacc_pragmas): Add "data". + * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA. + 2014-01-28 Thomas Schwinge <thomas@codesourcery.com> * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY, diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index f69486a..08374aa 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1169,6 +1169,7 @@ static vec<pragma_ns_name> registered_pp_pragmas; struct omp_pragma_def { const char *name; unsigned int id; }; static const struct omp_pragma_def oacc_pragmas[] = { + { "data", PRAGMA_OACC_DATA }, { "parallel", PRAGMA_OACC_PARALLEL }, }; static const struct omp_pragma_def omp_pragmas[] = { diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index 1ea5b1d..d092f9f 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see typedef enum pragma_kind { PRAGMA_NONE = 0, + PRAGMA_OACC_DATA, PRAGMA_OACC_PARALLEL, PRAGMA_OMP_ATOMIC, PRAGMA_OMP_BARRIER, diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index b199957..9b95725 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition. + (c_parser_oacc_data): New function. + (c_parser_omp_construct): Handle PRAGMA_OACC_DATA. + * c-tree.h (c_finish_oacc_data): New prototype. + * c-typeck.c (c_finish_oacc_data): New function. + 2014-02-17 Thomas Schwinge <thomas@codesourcery.com> * c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin, diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 7850eab..4643722 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -4776,10 +4776,14 @@ c_parser_label (c_parser *parser) openacc-construct: parallel-construct + data-construct parallel-construct: parallel-directive structured-block + data-construct: + data-directive structured-block + OpenMP: statement: @@ -11362,6 +11366,41 @@ c_parser_omp_structured_block (c_parser *parser) } /* OpenACC 2.0: + # pragma acc data oacc-data-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token. +*/ + +#define OACC_DATA_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_data (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK, + "#pragma acc data"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_data (loc, clauses, block); + + return stmt; +} + +/* OpenACC 2.0: # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block @@ -13675,6 +13714,9 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { + case PRAGMA_OACC_DATA: + stmt = c_parser_oacc_data (loc, parser); + break; case PRAGMA_OACC_PARALLEL: stmt = c_parser_oacc_parallel (loc, parser); break; diff --git gcc/c/c-tree.h gcc/c/c-tree.h index c174c7a..c84d3d7 100644 --- gcc/c/c-tree.h +++ gcc/c/c-tree.h @@ -634,6 +634,7 @@ extern tree c_finish_goto_label (location_t, tree); extern tree c_finish_goto_ptr (location_t, tree); extern tree c_expr_to_decl (tree, bool *, bool *); extern tree c_finish_oacc_parallel (location_t, tree, tree); +extern tree c_finish_oacc_data (location_t, tree, tree); extern tree c_begin_omp_parallel (void); extern tree c_finish_omp_parallel (location_t, tree, tree); extern tree c_begin_omp_task (void); diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 76d655b..8c4445b 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -11122,6 +11122,25 @@ c_finish_oacc_parallel (location_t loc, tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_DATA. */ + +tree +c_finish_oacc_data (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_DATA_CLAUSES (stmt) = clauses; + OACC_DATA_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Like c_begin_compound_stmt, except force the retention of the BLOCK. */ tree diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index fbccfa3..41d73b6 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,13 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC + data construct. + * c-c++-common/goacc/nesting-fail-1.c: Likewise. + * c-c++-common/goacc/parallel-fail-1.c: Rename to... + * c-c++-common/goacc/clauses-fail.c: ... this new file. Extend + for OpenACC data construct. + * c-c++-common/goacc/data-1.c: New file. + 2014-02-18 Thomas Schwinge <thomas@codesourcery.com> * gcc.dg/goacc/parallel-sb-1.c: New file. diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 875ec66..78fb45b 100644 --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -1,7 +1,7 @@ /* TODO: Some of these should either be allowed or fail with a more sensible error message. */ void -f1 (void) +f_omp (void) { int i; @@ -9,6 +9,8 @@ f1 (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp for @@ -16,49 +18,68 @@ f1 (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp sections { + { #pragma acc parallel /* { dg-error "may not be nested" } */ - ; + ; + } +#pragma omp section + { +#pragma acc data /* { dg-error "may not be nested" } */ + ; + } } #pragma omp single { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp task { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp master { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp critical { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp ordered { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } } /* TODO: Some of these should either be allowed or fail with a more sensible error message. */ void -f2 (void) +f_acc_parallel (void) { #pragma acc parallel { @@ -119,3 +140,68 @@ f2 (void) ; } } + +/* TODO: Some of these should either be allowed or fail with a more sensible + error message. */ +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma omp parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { + int i; +#pragma omp for /* { dg-error "may not be nested" } */ + for (i = 0; i < 3; i++) + ; + } + +#pragma acc data + { +#pragma omp sections /* { dg-error "may not be nested" } */ + { + ; + } + } + +#pragma acc data + { +#pragma omp single /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { +#pragma omp task /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { +#pragma omp master /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { +#pragma omp critical /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "may not be nested" } */ + } + +#pragma acc data + { +#pragma omp ordered /* { dg-error "may not be nested" } */ + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c new file mode 100644 index 0000000..b0dd042 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c @@ -0,0 +1,9 @@ +void +f (void) +{ +#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */ + ; + +#pragma acc data two /* { dg-error "expected clause before 'two'" } */ + ; +} diff --git gcc/testsuite/c-c++-common/goacc/data-1.c gcc/testsuite/c-c++-common/goacc/data-1.c new file mode 100644 index 0000000..8094575 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/data-1.c @@ -0,0 +1,6 @@ +void +foo (void) +{ +#pragma acc data + ; +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 6501397..24a4c11 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -1,11 +1,27 @@ /* TODO: While the OpenACC specification does allow for certain kinds of nesting, we don't support that yet. */ void -f1 (void) +f_acc_parallel (void) { #pragma acc parallel { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; + } +} + +/* TODO: While the OpenACC specification does allow for certain kinds of + nesting, we don't support that yet. */ +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } } diff --git gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c deleted file mode 100644 index efc6f14..0000000 --- gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c +++ /dev/null @@ -1,6 +0,0 @@ -void -foo (void) -{ -#pragma acc parallel foo /* { dg-error "expected clause before 'foo'" } */ - foo (); -} diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 5c15656..b90b09b 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,7 @@ 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + * testsuite/libgomp.oacc-c/data-1.c: New file. + * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New diff --git libgomp/testsuite/libgomp.oacc-c/data-1.c libgomp/testsuite/libgomp.oacc-c/data-1.c new file mode 100644 index 0000000..8f9a17a --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/data-1.c @@ -0,0 +1,170 @@ +/* { dg-do run } */ + +extern void abort (); + +int i; + +int main(void) +{ + int j; + +#if 0 + i = -1; + j = -2; +#pragma acc data copyin (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != 2 || j != 1) + abort (); + + i = -1; + j = -2; +#pragma acc data copyout (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); + + i = -1; + j = -2; +#pragma acc data copy (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); + + i = -1; + j = -2; +#pragma acc data create (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + + i = -1; + j = -2; +#pragma acc data present_or_copyin (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != 2 || j != 1) + abort (); + +#if 0 + i = -1; + j = -2; +#pragma acc data present_or_copyout (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + + i = -1; + j = -2; +#pragma acc data present_or_copy (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); + +#if 0 + i = -1; + j = -2; +#pragma acc data present_or_create (i, j) + { + // TODO: check that variables have been mapped. + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + +#if 0 + i = -1; + j = -2; +#pragma acc data present (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + +#if 0 + i = -1; + j = -2; +#pragma acc data + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + + return 0; +}

[gomp4,3/3] OpenACC data construct support in the C front end.

Commit Message

Patch