From 88bae6f494dc867edd8e6257658974d629bdc53b Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 27 Oct 2015 09:39:15 +0100 Subject: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++ gcc/c-family/ PR c/64765 PR c/64880 * c-common.h (c_oacc_split_loop_clauses): Declare function. * c-omp.c (c_oacc_split_loop_clauses): New function. gcc/c/ PR c/64765 PR c/64880 * c-parser.c (c_parser_oacc_loop): Add mask, cclauses formal parameters, and handle these. Adjust all users. (c_parser_oacc_kernels, c_parser_oacc_parallel): Merge functions into... (c_parser_oacc_kernels_parallel): ... this new function. Adjust all users. * c-tree.h (c_finish_oacc_parallel, c_finish_oacc_kernels): Don't declare functions. (c_finish_omp_construct): Declare function. * c-typeck.c (c_finish_oacc_parallel, c_finish_oacc_kernels): Merge functions into... (c_finish_omp_construct): ... this new function. gcc/cp/ PR c/64765 PR c/64880 * cp-tree.h (finish_oacc_kernels, finish_oacc_parallel): Don't declare functions. (finish_omp_construct): Declare function. * parser.c (cp_parser_oacc_loop): Add p_name, mask, cclauses formal parameters, and handle these. Adjust all users. (cp_parser_oacc_kernels, cp_parser_oacc_parallel): Merge functions into... (cp_parser_oacc_kernels_parallel): ... this new function. Adjust all users. * semantics.c (finish_oacc_kernels, finish_oacc_parallel): Merge functions into... (finish_omp_construct): ... this new function. gcc/ * tree.h (OACC_PARALLEL_BODY, OACC_PARALLEL_CLAUSES) (OACC_KERNELS_BODY, OACC_KERNELS_CLAUSES, OACC_KERNELS_COMBINED) (OACC_PARALLEL_COMBINED): Don't define macros. Adjust all users. gcc/testsuite/ PR c/64765 PR c/64880 * c-c++-common/goacc/loop-1.c: Don't skip for C++. Don't prune sorry message. (PR64765): New function. * gfortran.dg/goacc/coarray_2.f90: XFAIL. * gfortran.dg/goacc/combined_loop.f90: Extend. Don't prune sorry message. * gfortran.dg/goacc/cray.f95: Refine prune directive. * gfortran.dg/goacc/parameter.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/combdir-1.c: New file. * testsuite/libgomp.oacc-fortran/combdir-1.f90: Likewise. From-SVN: r229404 --- gcc/cp/parser.c | 147 +++++++++++++++++++++++++++++++++----------------------- 1 file changed, 86 insertions(+), 61 deletions(-) (limited to 'gcc/cp/parser.c') diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7555bf3..0354029 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -34298,69 +34298,64 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, } /* OpenACC 2.0: - # pragma acc kernels oacc-kernels-clause[optseq] new-line - structured-block */ - -#define OACC_KERNELS_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) - -static tree -cp_parser_oacc_kernels (cp_parser *parser, cp_token *pragma_tok) -{ - tree stmt, clauses, block; - unsigned int save; - - clauses = cp_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK, - "#pragma acc kernels", pragma_tok); - - block = begin_omp_parallel (); - save = cp_parser_begin_omp_structured_block (parser); - cp_parser_statement (parser, NULL_TREE, false, NULL); - cp_parser_end_omp_structured_block (parser, save); - stmt = finish_oacc_kernels (clauses, block); - return stmt; -} - -/* OpenACC 2.0: # pragma acc loop oacc-loop-clause[optseq] new-line structured-block */ #define OACC_LOOP_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) ) static tree -cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok) +cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, + omp_clause_mask mask, tree *cclauses) { - tree stmt, clauses, block; - int save; + strcat (p_name, " loop"); + mask |= OACC_LOOP_CLAUSE_MASK; - clauses = cp_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, - "#pragma acc loop", pragma_tok); + tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, + cclauses == NULL); + if (cclauses) + { + clauses = c_oacc_split_loop_clauses (clauses, cclauses); + if (*cclauses) + finish_omp_clauses (*cclauses, false); + if (clauses) + finish_omp_clauses (clauses, false); + } - block = begin_omp_structured_block (); - save = cp_parser_begin_omp_structured_block (parser); - stmt = cp_parser_omp_for_loop (parser, OACC_LOOP, clauses, NULL); + tree block = begin_omp_structured_block (); + int save = cp_parser_begin_omp_structured_block (parser); + tree stmt = cp_parser_omp_for_loop (parser, OACC_LOOP, clauses, NULL); cp_parser_end_omp_structured_block (parser, save); add_stmt (finish_omp_structured_block (block)); + return stmt; } /* OpenACC 2.0: + # pragma acc kernels oacc-kernels-clause[optseq] new-line + structured-block + + or + # pragma acc parallel oacc-parallel-clause[optseq] new-line - structured-block */ + structured-block +*/ + +#define OACC_KERNELS_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) #define OACC_PARALLEL_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ @@ -34379,23 +34374,53 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static tree -cp_parser_oacc_parallel (cp_parser *parser, cp_token *pragma_tok) +cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, + char *p_name) { - tree stmt, clauses, block; - unsigned int save; + omp_clause_mask mask; + enum tree_code code; + switch (pragma_tok->pragma_kind) + { + case PRAGMA_OACC_KERNELS: + strcat (p_name, " kernels"); + mask = OACC_KERNELS_CLAUSE_MASK; + code = OACC_KERNELS; + break; + case PRAGMA_OACC_PARALLEL: + strcat (p_name, " parallel"); + mask = OACC_PARALLEL_CLAUSE_MASK; + code = OACC_PARALLEL; + break; + default: + gcc_unreachable (); + } - clauses = cp_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, - "#pragma acc parallel", pragma_tok); + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) + { + const char *p + = IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value); + if (strcmp (p, "loop") == 0) + { + cp_lexer_consume_token (parser->lexer); + mask |= OACC_LOOP_CLAUSE_MASK; - block = begin_omp_parallel (); - save = cp_parser_begin_omp_structured_block (parser); + tree block = begin_omp_parallel (); + tree clauses; + cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, &clauses); + return finish_omp_construct (code, block, clauses); + } + } + + tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok); + + tree block = begin_omp_parallel (); + unsigned int save = cp_parser_begin_omp_structured_block (parser); cp_parser_statement (parser, NULL_TREE, false, NULL); cp_parser_end_omp_structured_block (parser, save); - stmt = finish_oacc_parallel (clauses, block); - return stmt; + return finish_omp_construct (code, block, clauses); } /* OpenACC 2.0: @@ -35290,13 +35315,13 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok) stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, false); break; case PRAGMA_OACC_KERNELS: - stmt = cp_parser_oacc_kernels (parser, pragma_tok); + case PRAGMA_OACC_PARALLEL: + strcpy (p_name, "#pragma acc"); + stmt = cp_parser_oacc_kernels_parallel (parser, pragma_tok, p_name); break; case PRAGMA_OACC_LOOP: - stmt = cp_parser_oacc_loop (parser, pragma_tok); - break; - case PRAGMA_OACC_PARALLEL: - stmt = cp_parser_oacc_parallel (parser, pragma_tok); + strcpy (p_name, "#pragma acc"); + stmt = cp_parser_oacc_loop (parser, pragma_tok, p_name, mask, NULL); break; case PRAGMA_OACC_UPDATE: stmt = cp_parser_oacc_update (parser, pragma_tok); -- cgit v1.1