[gomp4,2/2] Initial support for the OpenACC kernels construct in the C front end.

Message ID	1393579386-11666-2-git-send-email-thomas@codesourcery.com
State	New
Headers	show Return-Path: <gcc-patches-return-362518-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=P4mmpE7GQTiDwWZ0 sFMPfgwK4LYwdp8l8X2G3ycgkfr//jRbrwscpbKvdFMc72dWcQcwX03PCMdhFUK6 tYBpB5X3bEV7RO8WBm09JgMdYGBNVeNsaxWkrHP354uwlDVGh5JyOC5TulcIX00g y709q1ANYNlNZRqzCyGFzHvoPaQ= 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 2/2] Initial support for the OpenACC kernels construct in the C front end. Date: Fri, 28 Feb 2014 10:23:01 +0100 Message-ID: <1393579386-11666-2-git-send-email-thomas@codesourcery.com> In-Reply-To: <1393579386-11666-1-git-send-email-thomas@codesourcery.com> References: <1393579386-11666-1-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 3da377f..3b4a335 100644 --- gcc/c-family/ChangeLog.gomp +++ gcc/c-family/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-02-28 Thomas Schwinge <thomas@codesourcery.com> + + * c-pragma.c (oacc_pragmas): Add "kernels". + * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_KERNELS. + 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> * c-pragma.c (oacc_pragmas): Add "data". diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index 08374aa..ee0ee93 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1170,6 +1170,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 }, + { "kernels", PRAGMA_OACC_KERNELS }, { "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 d092f9f..d55a511 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -28,6 +28,7 @@ typedef enum pragma_kind { PRAGMA_NONE = 0, PRAGMA_OACC_DATA, + PRAGMA_OACC_KERNELS, PRAGMA_OACC_PARALLEL, PRAGMA_OMP_ATOMIC, PRAGMA_OMP_BARRIER, diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 9b95725..0551026 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-02-28 Thomas Schwinge <thomas@codesourcery.com> + + * c-parser.c (OACC_KERNELS_CLAUSE_MASK): New macro definition. + (c_parser_oacc_kernels): New function. + (c_parser_omp_construct): Handle PRAGMA_OACC_KERNELS. + * c-tree.h (c_finish_oacc_kernels): New prototype. + * c-typeck.c (c_finish_oacc_kernels): New function. + 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> * c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition. diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 4643722..c94e442 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -4776,11 +4776,15 @@ c_parser_label (c_parser *parser) openacc-construct: parallel-construct + kernels-construct data-construct parallel-construct: parallel-directive structured-block + kernels-construct: + kernels-directive structured-block + data-construct: data-directive structured-block @@ -11401,6 +11405,41 @@ c_parser_oacc_data (location_t loc, c_parser *parser) } /* OpenACC 2.0: + # pragma acc kernels oacc-kernels-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token. +*/ + +#define OACC_KERNELS_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_kernels (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK, + "#pragma acc kernels"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_kernels (loc, clauses, block); + + return stmt; +} + +/* OpenACC 2.0: # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block @@ -13717,6 +13756,9 @@ c_parser_omp_construct (c_parser *parser) case PRAGMA_OACC_DATA: stmt = c_parser_oacc_data (loc, parser); break; + case PRAGMA_OACC_KERNELS: + stmt = c_parser_oacc_kernels (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 c84d3d7..b6bea31 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_kernels (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); diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 8c4445b..191adfb 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_KERNELS, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_KERNELS. */ + +tree +c_finish_oacc_kernels (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_KERNELS); + TREE_TYPE (stmt) = void_type_node; + OACC_KERNELS_CLAUSES (stmt) = clauses; + OACC_KERNELS_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Generate OACC_DATA, with CLAUSES and BLOCK as its compound statement. LOC is the location of the OACC_DATA. */ diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 41d73b6..1bfb2f3 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,19 @@ +2014-02-28 Thomas Schwinge <thomas@codesourcery.com> + + * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC + kernels construct. + * c-c++-common/goacc/clauses-fail.c: Likewise. + * c-c++-common/goacc/data-clause-duplicate-1.c: Likewise. + * c-c++-common/goacc/deviceptr-1.c: Likewise. + * c-c++-common/goacc/nesting-fail-1.c: Likewise. + * c-c++-common/goacc/kernels-1.c: New file. + * gcc.dg/goacc/parallel-sb-1.c: Rename to... + * gcc.dg/goacc/sb-1.c: ... this new file, and extend for OpenACC + kernels and data constructs. + * gcc.dg/goacc/parallel-sb-2.c: Rename to... + * gcc.dg/goacc/sb-2.c: ... this new file, and extend for OpenACC + kernels and data constructs. + 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC 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 78fb45b..14103a6 100644 --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -9,6 +9,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -18,6 +20,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -30,6 +34,11 @@ f_omp (void) } #pragma omp section { +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; + } +#pragma omp section + { #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -39,6 +48,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -47,6 +58,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -55,6 +68,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -63,6 +78,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -71,6 +88,8 @@ f_omp (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -144,6 +163,71 @@ f_acc_parallel (void) /* TODO: Some of these should either be allowed or fail with a more sensible error message. */ void +f_acc_kernels (void) +{ +#pragma acc kernels + { +#pragma omp parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc kernels + { + int i; +#pragma omp for /* { dg-error "may not be nested" } */ + for (i = 0; i < 3; i++) + ; + } + +#pragma acc kernels + { +#pragma omp sections /* { dg-error "may not be nested" } */ + { + ; + } + } + +#pragma acc kernels + { +#pragma omp single /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc kernels + { +#pragma omp task /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc kernels + { +#pragma omp master /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc kernels + { +#pragma omp critical /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc kernels + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "may not be nested" } */ + } + +#pragma acc kernels + { +#pragma omp ordered /* { dg-error "may not be nested" } */ + ; + } +} + +/* TODO: Some of these should either be allowed or fail with a more sensible + error message. */ +void f_acc_data (void) { #pragma acc data diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c index b0dd042..133bf81 100644 --- gcc/testsuite/c-c++-common/goacc/clauses-fail.c +++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c @@ -4,6 +4,9 @@ f (void) #pragma acc parallel one /* { dg-error "expected clause before 'one'" } */ ; +#pragma acc kernels eins /* { dg-error "expected clause before 'eins'" } */ + ; + #pragma acc data two /* { dg-error "expected clause before 'two'" } */ ; } diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c index 1bcf5be..4cb3cc2 100644 --- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -4,9 +4,9 @@ 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 kernels 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) +#pragma acc data 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 index 0f0cf0c..1ac63bd 100644 --- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -3,10 +3,10 @@ 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 kernels deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ ; -#pragma acc parallel deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */ +#pragma acc data 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 } */ @@ -14,9 +14,9 @@ fun1 (void) ; int i; -#pragma acc parallel deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */ +#pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */ ; -#pragma acc parallel deviceptr(i[0:4]) +#pragma acc data 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 } */ ; @@ -24,13 +24,13 @@ fun1 (void) float fa[10]; #pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */ ; -#pragma acc parallel deviceptr(fa[1:5]) +#pragma acc kernels 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 data deviceptr(fp) ; #pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ ; @@ -41,7 +41,7 @@ fun2 (void) { int i; float *fp; -#pragma acc parallel deviceptr(fp,u,fun2,i,fp) +#pragma acc kernels 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 } */ @@ -53,11 +53,11 @@ void fun3 (void) { float *fp; -#pragma acc parallel deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */ +#pragma acc data 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" } */ +#pragma acc kernels copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ ; } diff --git gcc/testsuite/c-c++-common/goacc/kernels-1.c gcc/testsuite/c-c++-common/goacc/kernels-1.c new file mode 100644 index 0000000..e91b81c --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/kernels-1.c @@ -0,0 +1,6 @@ +void +foo (void) +{ +#pragma acc kernels + ; +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 24a4c11..d88ee8a 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -7,6 +7,24 @@ f_acc_parallel (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { 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_kernels (void) +{ +#pragma acc parallel + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } @@ -21,6 +39,8 @@ f_acc_data (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc kernels /* { dg-error "may not be nested" } */ + ; #pragma acc data /* { dg-error "may not be nested" } */ ; } diff --git gcc/testsuite/gcc.dg/goacc/parallel-sb-1.c gcc/testsuite/gcc.dg/goacc/parallel-sb-1.c deleted file mode 100644 index 3909916..0000000 --- gcc/testsuite/gcc.dg/goacc/parallel-sb-1.c +++ /dev/null @@ -1,22 +0,0 @@ -// { dg-do compile } - -void foo() -{ - bad1: - #pragma acc parallel - goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } - - goto bad2; // { dg-error "invalid entry to OpenACC structured block" } - #pragma acc parallel - { - bad2: ; - } - - #pragma acc parallel - { - int i; - goto ok1; - for (i = 0; i < 10; ++i) - { ok1: break; } - } -} diff --git gcc/testsuite/gcc.dg/goacc/parallel-sb-2.c gcc/testsuite/gcc.dg/goacc/parallel-sb-2.c deleted file mode 100644 index aede042..0000000 --- gcc/testsuite/gcc.dg/goacc/parallel-sb-2.c +++ /dev/null @@ -1,10 +0,0 @@ -// { dg-do compile } - -void foo(int i) -{ - switch (i) // { dg-error "invalid entry to OpenACC structured block" } - { - #pragma acc parallel - { case 0:; } - } -} diff --git gcc/testsuite/gcc.dg/goacc/sb-1.c gcc/testsuite/gcc.dg/goacc/sb-1.c new file mode 100644 index 0000000..24c88fe --- /dev/null +++ gcc/testsuite/gcc.dg/goacc/sb-1.c @@ -0,0 +1,54 @@ +// { dg-do compile } + +void foo() +{ + bad1: + #pragma acc parallel + goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } + #pragma acc kernels + goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } + #pragma acc data + goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } + + goto bad2_parallel; // { dg-error "invalid entry to OpenACC structured block" } + #pragma acc parallel + { + bad2_parallel: ; + } + + goto bad2_kernels; // { dg-error "invalid entry to OpenACC structured block" } + #pragma acc kernels + { + bad2_kernels: ; + } + + goto bad2_data; // { dg-error "invalid entry to OpenACC structured block" } + #pragma acc data + { + bad2_data: ; + } + + #pragma acc parallel + { + int i; + goto ok1_parallel; + for (i = 0; i < 10; ++i) + { ok1_parallel: break; } + } + + #pragma acc kernels + { + int i; + goto ok1_kernels; + for (i = 0; i < 10; ++i) + { ok1_kernels: break; } + } + + #pragma acc data + { + int i; + goto ok1_data; + for (i = 0; i < 10; ++i) + { ok1_data: break; } + } +} diff --git gcc/testsuite/gcc.dg/goacc/sb-2.c gcc/testsuite/gcc.dg/goacc/sb-2.c new file mode 100644 index 0000000..ec3eb95 --- /dev/null +++ gcc/testsuite/gcc.dg/goacc/sb-2.c @@ -0,0 +1,22 @@ +// { dg-do compile } + +void foo(int i) +{ + switch (i) // { dg-error "invalid entry to OpenACC structured block" } + { + #pragma acc parallel + { case 0:; } + } + + switch (i) // { dg-error "invalid entry to OpenACC structured block" } + { + #pragma acc kernels + { case 0:; } + } + + switch (i) // { dg-error "invalid entry to OpenACC structured block" } + { + #pragma acc data + { case 0:; } + } +} diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 3ea5901..7f9ce11 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,9 @@ 2014-02-28 Thomas Schwinge <thomas@codesourcery.com> + * testsuite/libgomp.oacc-c/goacc_kernels.c: New file. + * testsuite/libgomp.oacc-c/kernels-1.c: Likewise. + * testsuite/libgomp.oacc-c/parallel-1.c: Add one missing test. + * libgomp.map (GOACC_2.0): Add GOACC_kernels. * libgomp_g.h (GOACC_kernels): New prototype. * oacc-parallel.c (GOACC_kernels): New function. diff --git libgomp/testsuite/libgomp.oacc-c/goacc_kernels.c libgomp/testsuite/libgomp.oacc-c/goacc_kernels.c new file mode 100644 index 0000000..db1a37d --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/goacc_kernels.c @@ -0,0 +1,25 @@ +/* { dg-do run } */ + +#include "libgomp_g.h" + +extern void abort (); + +volatile int i; + +void +f (void *data) +{ + if (i != -1) + abort (); + i = 42; +} + +int main(void) +{ + i = -1; + GOACC_kernels (0, f, (const void *) 0, 0, (void *) 0, (void *) 0, (void *) 0); + if (i != 42) + abort (); + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c/kernels-1.c libgomp/testsuite/libgomp.oacc-c/kernels-1.c new file mode 100644 index 0000000..8550662 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/kernels-1.c @@ -0,0 +1,170 @@ +/* { dg-do run } */ + +extern void abort (); + +int i; + +int main(void) +{ + int j, v; + +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc kernels /* copyout */ present_or_copyout (v) 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 kernels /* 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 kernels /* 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 kernels /* 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 kernels /* 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 kernels /* 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 kernels /* 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 (); + + i = -1; + j = -2; + v = 0; +#pragma acc kernels /* copyout */ present_or_copyout (v) present_or_create (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != -1 || j != -2) + abort (); + +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc kernels /* 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 kernels /* 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; +} diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c index ff54b9d..68f7de5 100644 --- libgomp/testsuite/libgomp.oacc-c/parallel-1.c +++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c @@ -116,6 +116,20 @@ int main(void) if (v != 1 || i != 2 || j != 1) abort (); + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_create (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != -1 || j != -2) + abort (); + #if 0 i = -1; j = -2;

[gomp4,2/2] Initial support for the OpenACC kernels construct in the C front end.

Commit Message

Patch