* [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
@ 2015-10-08 16:39 Thomas Schwinge
2015-10-08 16:47 ` Joseph Myers
` (2 more replies)
0 siblings, 3 replies; 9+ messages in thread
From: Thomas Schwinge @ 2015-10-08 16:39 UTC (permalink / raw)
To: gcc-patches, Joseph Myers, Jason Merrill, Nathan Sidwell
Cc: Jakub Jelinek, James Norris, Cesar Philippidis
[-- Attachment #1: Type: text/plain, Size: 37079 bytes --]
Hi!
Some bits extracted out of gomp-4_0-branch, and some other bits
rewritten; here is a patch to support OpenACC Combined Directives in C,
C++. (The Fortran front end already does support these.)
As far as I know, Jakub is not available at this time, so maybe the C
(Joseph) and C++ (Jason, Nathan) front end maintainers could please
review this, instead of him? (The front end changes as well as the few
other cleanup changes should all be straight forward.) OK for trunk once
bootstrap tested?
commit 9626356d641129381306f2ad5d884d5b7f7a5fc7
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu Oct 8 15:59:54 2015 +0200
[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.
---
gcc/c-family/c-common.h | 1 +
gcc/c-family/c-omp.c | 39 +++++-
gcc/c/c-parser.c | 148 +++++++++-----------
gcc/c/c-tree.h | 3 +-
gcc/c/c-typeck.c | 36 ++---
gcc/cp/cp-tree.h | 3 +-
gcc/cp/parser.c | 149 ++++++++++++---------
gcc/cp/semantics.c | 54 +++-----
gcc/fortran/trans-openmp.c | 4 -
gcc/gimplify.c | 16 +--
gcc/testsuite/c-c++-common/goacc/loop-1.c | 10 +-
gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 | 1 +
gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 | 9 +-
gcc/testsuite/gfortran.dg/goacc/cray.f95 | 2 +-
gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 2 +-
gcc/tree-pretty-print.c | 11 +-
gcc/tree.def | 8 +-
gcc/tree.h | 20 ---
.../libgomp.oacc-c-c++-common/combdir-1.c | 52 +++++++
.../testsuite/libgomp.oacc-fortran/combdir-1.f90 | 37 +++++
20 files changed, 336 insertions(+), 269 deletions(-)
diff --git gcc/c-family/c-common.h gcc/c-family/c-common.h
index d5fb499..94c68b9 100644
--- gcc/c-family/c-common.h
+++ gcc/c-family/c-common.h
@@ -1268,6 +1268,7 @@ extern void c_finish_omp_taskyield (location_t);
extern tree c_finish_omp_for (location_t, enum tree_code, tree, tree, tree,
tree, tree, tree);
extern tree c_finish_oacc_wait (location_t, tree, tree);
+extern tree c_oacc_split_loop_clauses (tree, tree *);
extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask,
tree, tree *);
extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree);
diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c
index cdd2ee8..1a64884 100644
--- gcc/c-family/c-omp.c
+++ gcc/c-family/c-omp.c
@@ -679,13 +679,46 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv,
}
}
-/* Right now we have 14 different combined constructs, this
- function attempts to split or duplicate clauses for combined
+/* This function splits clauses for OpenACC combined loop
+ constructs. OpenACC combined loop constructs are:
+ #pragma acc kernels loop
+ #pragma acc parallel loop
+*/
+
+tree
+c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
+{
+ tree next, loop_clauses;
+
+ loop_clauses = *not_loop_clauses = NULL_TREE;
+ for (; clauses ; clauses = next)
+ {
+ next = OMP_CLAUSE_CHAIN (clauses);
+
+ switch (OMP_CLAUSE_CODE (clauses))
+ {
+ case OMP_CLAUSE_COLLAPSE:
+ case OMP_CLAUSE_REDUCTION:
+ OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
+ loop_clauses = clauses;
+ break;
+
+ default:
+ OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses;
+ *not_loop_clauses = clauses;
+ break;
+ }
+ }
+
+ return loop_clauses;
+}
+
+/* This function attempts to split or duplicate clauses for OpenMP combined
constructs. CODE is the innermost construct in the combined construct,
and MASK allows to determine which constructs are combined together,
as every construct has at least one clause that no other construct
has (except for OMP_SECTIONS, but that can be only combined with parallel).
- Combined constructs are:
+ OpenMP combined constructs are:
#pragma omp parallel for
#pragma omp parallel sections
#pragma omp parallel for simd
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 2d24c21..ca36a13 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1237,7 +1237,6 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
unsigned int * = NULL);
static void c_parser_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
-static tree c_parser_oacc_loop (location_t, c_parser *, char *);
static void c_parser_omp_construct (c_parser *);
static void c_parser_omp_threadprivate (c_parser *);
static void c_parser_omp_barrier (c_parser *);
@@ -12118,60 +12117,6 @@ 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_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
-c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
-{
- tree stmt, clauses = NULL_TREE, block;
-
- strcat (p_name, " kernels");
-
- if (c_parser_next_token_is (parser, CPP_NAME))
- {
- const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
- if (strcmp (p, "loop") == 0)
- {
- c_parser_consume_token (parser);
- block = c_begin_omp_parallel ();
- c_parser_oacc_loop (loc, parser, p_name);
- stmt = c_finish_oacc_kernels (loc, clauses, block);
- OACC_KERNELS_COMBINED (stmt) = 1;
- return stmt;
- }
- }
-
- clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
- p_name);
-
- 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 enter data oacc-enter-data-clause[optseq] new-line
or
@@ -12261,16 +12206,25 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
static tree
-c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
+c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
+ omp_clause_mask mask, tree *cclauses)
{
- tree stmt, clauses, block;
-
strcat (p_name, " loop");
+ mask |= OACC_LOOP_CLAUSE_MASK;
- clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
+ tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
+ cclauses == NULL);
+ if (cclauses)
+ {
+ clauses = c_oacc_split_loop_clauses (clauses, cclauses);
+ if (*cclauses)
+ c_finish_omp_clauses (*cclauses);
+ if (clauses)
+ c_finish_omp_clauses (clauses);
+ }
- block = c_begin_compound_stmt (true);
- stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
+ tree block = c_begin_compound_stmt (true);
+ tree stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
block = c_end_compound_stmt (loc, block, true);
add_stmt (block);
@@ -12278,12 +12232,32 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
}
/* 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
LOC is the location of the #pragma token.
*/
+#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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
@@ -12304,11 +12278,26 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static tree
-c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
+c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser,
+ enum pragma_kind p_kind, char *p_name)
{
- tree stmt, clauses = NULL_TREE, block;
-
- strcat (p_name, " parallel");
+ omp_clause_mask mask;
+ enum tree_code code;
+ switch (p_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 ();
+ }
if (c_parser_next_token_is (parser, CPP_NAME))
{
@@ -12316,23 +12305,21 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
if (strcmp (p, "loop") == 0)
{
c_parser_consume_token (parser);
- block = c_begin_omp_parallel ();
- c_parser_oacc_loop (loc, parser, p_name);
- stmt = c_finish_oacc_parallel (loc, clauses, block);
- OACC_PARALLEL_COMBINED (stmt) = 1;
- return stmt;
+ mask |= OACC_LOOP_CLAUSE_MASK;
+
+ tree block = c_begin_omp_parallel ();
+ tree clauses;
+ c_parser_oacc_loop (loc, parser, p_name, mask, &clauses);
+ return c_finish_omp_construct (loc, code, block, clauses);
}
}
- clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
- p_name);
+ tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
- block = c_begin_omp_parallel ();
+ tree block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
- stmt = c_finish_oacc_parallel (loc, clauses, block);
-
- return stmt;
+ return c_finish_omp_construct (loc, code, block, clauses);
}
/* OpenACC 2.0:
@@ -14766,16 +14753,13 @@ c_parser_omp_construct (c_parser *parser)
stmt = c_parser_oacc_data (loc, parser);
break;
case PRAGMA_OACC_KERNELS:
+ case PRAGMA_OACC_PARALLEL:
strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_kernels (loc, parser, p_name);
+ stmt = c_parser_oacc_kernels_parallel (loc, parser, p_kind, p_name);
break;
case PRAGMA_OACC_LOOP:
strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_loop (loc, parser, p_name);
- break;
- case PRAGMA_OACC_PARALLEL:
- strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_parallel (loc, parser, p_name);
+ stmt = c_parser_oacc_loop (loc, parser, p_name, mask, NULL);
break;
case PRAGMA_OACC_WAIT:
strcpy (p_name, "#pragma wait");
diff --git gcc/c/c-tree.h gcc/c/c-tree.h
index 667529a..6c5c07b 100644
--- gcc/c/c-tree.h
+++ gcc/c/c-tree.h
@@ -640,8 +640,7 @@ extern tree c_finish_bc_stmt (location_t, tree *, bool);
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_omp_construct (location_t, enum tree_code, 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 df3245a4..9c64c9a 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11453,39 +11453,19 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se)
return expr;
}
\f
-/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_PARALLEL. */
+/* Generate OMP construct CODE, with BODY and CLAUSES as its compound
+ statement. LOC is the location of the construct. */
tree
-c_finish_oacc_parallel (location_t loc, tree clauses, tree block)
+c_finish_omp_construct (location_t loc, enum tree_code code, tree body,
+ tree clauses)
{
- tree stmt;
+ body = c_end_compound_stmt (loc, body, true);
- block = c_end_compound_stmt (loc, block, true);
-
- stmt = make_node (OACC_PARALLEL);
- TREE_TYPE (stmt) = void_type_node;
- OACC_PARALLEL_CLAUSES (stmt) = clauses;
- OACC_PARALLEL_BODY (stmt) = block;
- SET_EXPR_LOCATION (stmt, loc);
-
- 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 stmt = make_node (code);
TREE_TYPE (stmt) = void_type_node;
- OACC_KERNELS_CLAUSES (stmt) = clauses;
- OACC_KERNELS_BODY (stmt) = block;
+ OMP_BODY (stmt) = body;
+ OMP_CLAUSES (stmt) = clauses;
SET_EXPR_LOCATION (stmt, loc);
return add_stmt (stmt);
diff --git gcc/cp/cp-tree.h gcc/cp/cp-tree.h
index 5acb065..1d1bcc4 100644
--- gcc/cp/cp-tree.h
+++ gcc/cp/cp-tree.h
@@ -6300,8 +6300,7 @@ extern void finish_omp_threadprivate (tree);
extern tree begin_omp_structured_block (void);
extern tree finish_omp_structured_block (tree);
extern tree finish_oacc_data (tree, tree);
-extern tree finish_oacc_kernels (tree, tree);
-extern tree finish_oacc_parallel (tree, tree);
+extern tree finish_omp_construct (enum tree_code, tree, tree);
extern tree begin_omp_parallel (void);
extern tree finish_omp_parallel (tree, tree);
extern tree begin_omp_task (void);
diff --git gcc/cp/parser.c gcc/cp/parser.c
index d4ef7f9..f769539 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -33132,69 +33132,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);
+ if (clauses)
+ finish_omp_clauses (clauses);
+ }
- 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) \
@@ -33213,23 +33208,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:
@@ -33981,13 +34006,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);
- 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_kernels_parallel (parser, pragma_tok, p_name);
+ break;
+ case PRAGMA_OACC_LOOP:
+ 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);
diff --git gcc/cp/semantics.c gcc/cp/semantics.c
index c1f4330..3b9b714 100644
--- gcc/cp/semantics.c
+++ gcc/cp/semantics.c
@@ -6124,8 +6124,17 @@ finish_omp_structured_block (tree block)
return do_poplevel (block);
}
+/* Similarly, except force the retention of the BLOCK. */
+
+tree
+begin_omp_parallel (void)
+{
+ keep_next_level (true);
+ return begin_omp_structured_block ();
+}
+
/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_DATA. */
+ statement. */
tree
finish_oacc_data (tree clauses, tree block)
@@ -6142,51 +6151,22 @@ finish_oacc_data (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. */
+/* Generate OMP construct CODE, with BODY and CLAUSES as its compound
+ statement. */
tree
-finish_oacc_kernels (tree clauses, tree block)
+finish_omp_construct (enum tree_code code, tree body, tree clauses)
{
- tree stmt;
+ body = finish_omp_structured_block (body);
- block = finish_omp_structured_block (block);
-
- stmt = make_node (OACC_KERNELS);
- TREE_TYPE (stmt) = void_type_node;
- OACC_KERNELS_CLAUSES (stmt) = clauses;
- OACC_KERNELS_BODY (stmt) = block;
-
- return add_stmt (stmt);
-}
-
-/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_PARALLEL. */
-
-tree
-finish_oacc_parallel (tree clauses, tree block)
-{
- tree stmt;
-
- block = finish_omp_structured_block (block);
-
- stmt = make_node (OACC_PARALLEL);
+ tree stmt = make_node (code);
TREE_TYPE (stmt) = void_type_node;
- OACC_PARALLEL_CLAUSES (stmt) = clauses;
- OACC_PARALLEL_BODY (stmt) = block;
+ OMP_BODY (stmt) = body;
+ OMP_CLAUSES (stmt) = clauses;
return add_stmt (stmt);
}
-/* Similarly, except force the retention of the BLOCK. */
-
-tree
-begin_omp_parallel (void)
-{
- keep_next_level (true);
- return begin_omp_structured_block ();
-}
-
tree
finish_omp_parallel (tree clauses, tree body)
{
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 294b6ef..bf04b7f 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -3459,10 +3459,6 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
poplevel (0, 0);
stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
oacc_clauses);
- if (code->op == EXEC_OACC_KERNELS_LOOP)
- OACC_KERNELS_COMBINED (stmt) = 1;
- else
- OACC_PARALLEL_COMBINED (stmt) = 1;
gfc_add_expr_to_block (&block, stmt);
return gfc_finish_block (&block);
}
diff --git gcc/gimplify.c gcc/gimplify.c
index 25a81f6..c541847 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -8589,23 +8589,9 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_DATA:
case OACC_KERNELS:
- if (OACC_KERNELS_COMBINED (*expr_p))
- sorry ("directive not yet implemented");
- else
- gimplify_omp_workshare (expr_p, pre_p);
- ret = GS_ALL_DONE;
- break;
-
case OACC_PARALLEL:
- if (OACC_PARALLEL_COMBINED (*expr_p))
- sorry ("directive not yet implemented");
- else
- gimplify_omp_workshare (expr_p, pre_p);
- ret = GS_ALL_DONE;
- break;
-
- case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
case OMP_TARGET:
diff --git gcc/testsuite/c-c++-common/goacc/loop-1.c gcc/testsuite/c-c++-common/goacc/loop-1.c
index fea40e0..fc2e1a8 100644
--- gcc/testsuite/c-c++-common/goacc/loop-1.c
+++ gcc/testsuite/c-c++-common/goacc/loop-1.c
@@ -1,5 +1,3 @@
-/* { dg-skip-if "not yet" { c++ } } */
-
int test1()
{
int i, j, k, b[10];
@@ -69,4 +67,10 @@ int test1()
}
return 0;
}
-/* { dg-prune-output "sorry, unimplemented: directive not yet implemented" } */
+
+// PR64765
+void PR64765(float *f, double *r) {
+ int i;
+ #pragma acc kernels loop create(f) copy(r)
+ for(i = 64; i < 76; i += 5) {}
+}
diff --git gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
index f35d4b9..f9cf9ac 100644
--- gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
+++ gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
@@ -2,6 +2,7 @@
! { dg-additional-options "-fcoarray=lib" }
!
! PR fortran/63861
+! { dg-xfail-if "<http://gcc.gnu.org/PR63861>" { *-*-* } }
module test
contains
diff --git gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
index b8be649..abb10f9 100644
--- gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
+++ gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
@@ -6,7 +6,14 @@ subroutine oacc1()
implicit none
integer :: i
integer :: a
- !$acc parallel loop reduction(+:a) ! { dg-excess-errors "sorry, unimplemented: directive not yet implemented" }
+ !$acc parallel loop reduction(+:a)
do i = 1,5
enddo
+ !$acc end parallel loop
+ !$acc kernels loop collapse(2)
+ do i = 2,6
+ do a = 3,5
+ enddo
+ enddo
+ !$acc end kernels loop
end subroutine oacc1
diff --git gcc/testsuite/gfortran.dg/goacc/cray.f95 gcc/testsuite/gfortran.dg/goacc/cray.f95
index 8f2c077..52789fe 100644
--- gcc/testsuite/gfortran.dg/goacc/cray.f95
+++ gcc/testsuite/gfortran.dg/goacc/cray.f95
@@ -53,4 +53,4 @@ contains
!$acc update self (ptr)
end subroutine oacc1
end module test
-! { dg-prune-output "unimplemented" }
+! { dg-prune-output "ACC cache unimplemented" }
diff --git gcc/testsuite/gfortran.dg/goacc/parameter.f95 gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 1364181..454924a 100644
--- gcc/testsuite/gfortran.dg/goacc/parameter.f95
+++ gcc/testsuite/gfortran.dg/goacc/parameter.f95
@@ -29,4 +29,4 @@ contains
!$acc update self (a) ! { dg-error "not a variable" }
end subroutine oacc1
end module test
-! { dg-prune-output "unimplemented" }
+! { dg-prune-output "ACC cache unimplemented" }
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index ce3f6a8..48c5c98 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -2537,13 +2537,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, int flags,
case OACC_PARALLEL:
pp_string (pp, "#pragma acc parallel");
- dump_omp_clauses (pp, OACC_PARALLEL_CLAUSES (node), spc, flags);
- goto dump_omp_body;
+ goto dump_omp_clauses_body;
case OACC_KERNELS:
pp_string (pp, "#pragma acc kernels");
- dump_omp_clauses (pp, OACC_KERNELS_CLAUSES (node), spc, flags);
- goto dump_omp_body;
+ goto dump_omp_clauses_body;
case OACC_DATA:
pp_string (pp, "#pragma acc data");
@@ -2583,6 +2581,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, int flags,
case OMP_PARALLEL:
pp_string (pp, "#pragma omp parallel");
dump_omp_clauses (pp, OMP_PARALLEL_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ dump_omp_clauses_body:
+ dump_omp_clauses (pp, OMP_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
dump_omp_body:
if (!(flags & TDF_SLIM) && OMP_BODY (node))
diff --git gcc/tree.def gcc/tree.def
index 56580af..3e16e82 100644
--- gcc/tree.def
+++ gcc/tree.def
@@ -1039,14 +1039,14 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
not change the ordering of these codes. */
/* OpenACC - #pragma acc parallel [clause1 ... clauseN]
- Operand 0: OACC_PARALLEL_BODY: Code to be executed in parallel.
- Operand 1: OACC_PARALLEL_CLAUSES: List of clauses. */
+ Operand 0: OMP_BODY: Code to be executed in parallel.
+ Operand 1: OMP_CLAUSES: List of clauses. */
DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
/* OpenACC - #pragma acc kernels [clause1 ... clauseN]
- Operand 0: OACC_KERNELS_BODY: Sequence of kernels.
- Operand 1: OACC_KERNELS_CLAUSES: List of clauses. */
+ Operand 0: OMP_BODY: Sequence of kernels.
+ Operand 1: OMP_CLAUSES: List of clauses. */
DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
diff --git gcc/tree.h gcc/tree.h
index 4c803f4..fb34e4b 100644
--- gcc/tree.h
+++ gcc/tree.h
@@ -1,4 +1,3 @@
-
/* Definitions for the ubiquitous 'tree' type for GNU compilers.
Copyright (C) 1989-2015 Free Software Foundation, Inc.
@@ -1212,16 +1211,6 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_STANDALONE_CLAUSES(NODE) \
TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_CACHE, OMP_TARGET_UPDATE), 0)
-#define OACC_PARALLEL_BODY(NODE) \
- TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 0)
-#define OACC_PARALLEL_CLAUSES(NODE) \
- TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1)
-
-#define OACC_KERNELS_BODY(NODE) \
- TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 0)
-#define OACC_KERNELS_CLAUSES(NODE) \
- TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 1)
-
#define OACC_DATA_BODY(NODE) \
TREE_OPERAND (OACC_DATA_CHECK (NODE), 0)
#define OACC_DATA_CLAUSES(NODE) \
@@ -1316,15 +1305,6 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_SECTION_LAST(NODE) \
(OMP_SECTION_CHECK (NODE)->base.private_flag)
-/* True on an OACC_KERNELS statement if is represents combined kernels loop
- directive. */
-#define OACC_KERNELS_COMBINED(NODE) \
- (OACC_KERNELS_CHECK (NODE)->base.private_flag)
-
-/* Like OACC_KERNELS_COMBINED, but for parallel loop directive. */
-#define OACC_PARALLEL_COMBINED(NODE) \
- (OACC_PARALLEL_CHECK (NODE)->base.private_flag)
-
/* True on an OMP_PARALLEL statement if it represents an explicit
combined parallel work-sharing constructs. */
#define OMP_PARALLEL_COMBINED(NODE) \
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
new file mode 100644
index 0000000..a7def92
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+ const int N = 32;
+ float a[N], b[N];
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 1.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc parallel loop copy (a[0:N]) copy (b[0:N])
+ for (i = 0; i < N; i++)
+ {
+ b[i] = 2.0;
+ a[i] = a[i] + b[i];
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 2.0)
+ abort ();
+ }
+
+#pragma acc kernels loop copy (a[0:N]) copy (b[0:N])
+ for (i = 0; i < N; i++)
+ {
+ b[i] = 3.0;
+ a[i] = a[i] + b[i];
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ return 0;
+
+}
diff --git libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90 libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
new file mode 100644
index 0000000..0cd8a67
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
@@ -0,0 +1,37 @@
+! { dg-do run }
+
+program main
+ integer, parameter :: n = 32
+ real :: a(n), b(n);
+ integer :: i
+
+ do i = 1, n
+ a(i) = 1.0
+ b(i) = 0.0
+ end do
+
+ !$acc parallel loop copy (a(1:n)) copy (b(1:n))
+ do i = 1, n
+ b(i) = 2.0
+ a(i) = a(i) + b(i)
+ end do
+
+ do i = 1, n
+ if (a(i) .ne. 3.0) call abort
+
+ if (b(i) .ne. 2.0) call abort
+ end do
+
+ !$acc kernels loop copy (a(1:n)) copy (b(1:n))
+ do i = 1, n
+ b(i) = 3.0;
+ a(i) = a(i) + b(i)
+ end do
+
+ do i = 1, n
+ if (a(i) .ne. 6.0) call abort
+
+ if (b(i) .ne. 3.0) call abort
+ end do
+
+end program main
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-08 16:39 [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++ Thomas Schwinge
@ 2015-10-08 16:47 ` Joseph Myers
2015-10-09 12:26 ` Nathan Sidwell
2015-10-27 8:52 ` Thomas Schwinge
2 siblings, 0 replies; 9+ messages in thread
From: Joseph Myers @ 2015-10-08 16:47 UTC (permalink / raw)
To: Thomas Schwinge
Cc: gcc-patches, Jason Merrill, Nathan Sidwell, Jakub Jelinek,
James Norris, Cesar Philippidis
The C front-end changes are OK, but please follow up with fixes for any
issues Jakub identifies.
--
Joseph S. Myers
joseph@codesourcery.com
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-08 16:39 [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++ Thomas Schwinge
2015-10-08 16:47 ` Joseph Myers
@ 2015-10-09 12:26 ` Nathan Sidwell
2015-10-09 13:26 ` Thomas Schwinge
2015-10-27 8:52 ` Thomas Schwinge
2 siblings, 1 reply; 9+ messages in thread
From: Nathan Sidwell @ 2015-10-09 12:26 UTC (permalink / raw)
To: Thomas Schwinge, gcc-patches, Joseph Myers, Jason Merrill,
Nathan Sidwell
Cc: Jakub Jelinek, James Norris, Cesar Philippidis
On 10/08/15 12:39, Thomas Schwinge wrote:
> Hi!
>
> Some bits extracted out of gomp-4_0-branch, and some other bits
> rewritten; here is a patch to support OpenACC Combined Directives in C,
> C++. (The Fortran front end already does support these.)
>
> As far as I know, Jakub is not available at this time, so maybe the C
> (Joseph) and C++ (Jason, Nathan) front end maintainers could please
> review this, instead of him? (The front end changes as well as the few
> other cleanup changes should all be straight forward.) OK for trunk once
> bootstrap tested?
> --- gcc/cp/parser.c
> +++ gcc/cp/parser.c
> @@ -33132,69 +33132,64 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
> -/* 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)
>
Needs documentation.
> 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)
Likewise,
> diff --git gcc/cp/semantics.c gcc/cp/semantics.c
> index c1f4330..3b9b714 100644
> --- gcc/cp/semantics.c
> +++ gcc/cp/semantics.c
> @@ -6124,8 +6124,17 @@ finish_omp_structured_block (tree block)
> return do_poplevel (block);
> }
>
> +/* Similarly, except force the retention of the BLOCK. */
similar to what?
> +
> +tree
> +begin_omp_parallel (void)
nathan
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-09 12:26 ` Nathan Sidwell
@ 2015-10-09 13:26 ` Thomas Schwinge
2015-10-09 13:34 ` Nathan Sidwell
0 siblings, 1 reply; 9+ messages in thread
From: Thomas Schwinge @ 2015-10-09 13:26 UTC (permalink / raw)
To: Nathan Sidwell, Joseph Myers
Cc: Jakub Jelinek, James Norris, Cesar Philippidis, gcc-patches,
Jason Merrill
[-- Attachment #1: Type: text/plain, Size: 2770 bytes --]
Hi!
Thanks to Joseph and you for your review!
On Fri, 9 Oct 2015 08:26:37 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> On 10/08/15 12:39, Thomas Schwinge wrote:
> > Some bits extracted out of gomp-4_0-branch, and some other bits
> > rewritten; here is a patch to support OpenACC Combined Directives in C,
> > C++. (The Fortran front end already does support these.)
> >
> > As far as I know, Jakub is not available at this time, so maybe the C
> > (Joseph) and C++ (Jason, Nathan) front end maintainers could please
> > review this, instead of him? (The front end changes as well as the few
> > other cleanup changes should all be straight forward.) OK for trunk once
> > bootstrap tested?
Assuming the following review comments resolved, does this constitute
approval for the whole patch? (Will of course address any later review
comments, as usual.)
> > --- gcc/cp/parser.c
> > +++ gcc/cp/parser.c
> > @@ -33132,69 +33132,64 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
>
> > -/* 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)
> >
>
> Needs documentation.
>
>
> > 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)
>
> Likewise,
You mean the cp_parser_oacc_loop and cp_parser_oacc_kernels_parallel
functions need documentation? I agree it's a bit terse, but documenting
these by just listing the accepted parsing tokens "# pragma acc loop"
etc., followed by the *_CLAUSE_MASKs is what's done for the other
OpenACC/OpenMP directives in the C/C++ front ends. So, I don't see a
reason to be different for these two?
> > --- gcc/cp/semantics.c
> > +++ gcc/cp/semantics.c
> > @@ -6124,8 +6124,17 @@ finish_omp_structured_block (tree block)
> > return do_poplevel (block);
> > }
> >
> > +/* Similarly, except force the retention of the BLOCK. */
>
> similar to what?
>
> > +
> > +tree
> > +begin_omp_parallel (void)
I just moved begin_omp_parallel up a little bit in the file (where it
originally resided); the "similarly" applies to the function defined just
before.
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-09 13:26 ` Thomas Schwinge
@ 2015-10-09 13:34 ` Nathan Sidwell
2015-10-09 13:57 ` Jakub Jelinek
2015-10-09 14:00 ` Thomas Schwinge
0 siblings, 2 replies; 9+ messages in thread
From: Nathan Sidwell @ 2015-10-09 13:34 UTC (permalink / raw)
To: Thomas Schwinge, Joseph Myers
Cc: Jakub Jelinek, James Norris, Cesar Philippidis, gcc-patches,
Jason Merrill
On 10/09/15 09:26, Thomas Schwinge wrote:
> Hi!
> You mean the cp_parser_oacc_loop and cp_parser_oacc_kernels_parallel
> functions need documentation? I agree it's a bit terse, but documenting
> these by just listing the accepted parsing tokens "# pragma acc loop"
> etc., followed by the *_CLAUSE_MASKs is what's done for the other
> OpenACC/OpenMP directives in the C/C++ front ends. So, I don't see a
> reason to be different for these two?
What's the p_name argument for?
nathan
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-09 13:34 ` Nathan Sidwell
@ 2015-10-09 13:57 ` Jakub Jelinek
2015-10-09 14:00 ` Thomas Schwinge
1 sibling, 0 replies; 9+ messages in thread
From: Jakub Jelinek @ 2015-10-09 13:57 UTC (permalink / raw)
To: Nathan Sidwell
Cc: Thomas Schwinge, Joseph Myers, James Norris, Cesar Philippidis,
gcc-patches, Jason Merrill
On Fri, Oct 09, 2015 at 09:34:51AM -0400, Nathan Sidwell wrote:
> On 10/09/15 09:26, Thomas Schwinge wrote:
> >Hi!
>
> >You mean the cp_parser_oacc_loop and cp_parser_oacc_kernels_parallel
> >functions need documentation? I agree it's a bit terse, but documenting
> >these by just listing the accepted parsing tokens "# pragma acc loop"
> >etc., followed by the *_CLAUSE_MASKs is what's done for the other
> >OpenACC/OpenMP directives in the C/C++ front ends. So, I don't see a
> >reason to be different for these two?
>
>
> What's the p_name argument for?
That is a pointer to the name of the combined/composite construct.
This stuff works basically that if you have a combined construct,
you call the outermost parsing routine with a buffer long enough to
contain the longest possible combined construct name, and
strcpy (p_name, "#pragma omp {omp,acc}");
first, and then each parsing routine strcats the name of itself
after that (and merges in its own clause mask), and finally on
the innermost construct you have the full name of your combined construct in
p_name and the right clause mask in mask, so you call the clause
parsing with those two parameters, which might use the name
for diagnostics.
Jakub
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-09 13:34 ` Nathan Sidwell
2015-10-09 13:57 ` Jakub Jelinek
@ 2015-10-09 14:00 ` Thomas Schwinge
2015-10-11 22:19 ` Nathan Sidwell
1 sibling, 1 reply; 9+ messages in thread
From: Thomas Schwinge @ 2015-10-09 14:00 UTC (permalink / raw)
To: Nathan Sidwell
Cc: Jakub Jelinek, James Norris, Cesar Philippidis, gcc-patches,
Jason Merrill, Joseph Myers
[-- Attachment #1: Type: text/plain, Size: 827 bytes --]
Hi Nathan!
On Fri, 9 Oct 2015 09:34:51 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> On 10/09/15 09:26, Thomas Schwinge wrote:
> > You mean the cp_parser_oacc_loop and cp_parser_oacc_kernels_parallel
> > functions need documentation? I agree it's a bit terse, but documenting
> > these by just listing the accepted parsing tokens "# pragma acc loop"
> > etc., followed by the *_CLAUSE_MASKs is what's done for the other
> > OpenACC/OpenMP directives in the C/C++ front ends. So, I don't see a
> > reason to be different for these two?
>
>
> What's the p_name argument for?
It's s string describing the pragma as parsed thus far. Again, not
documenting that as well as our usage of it is totally "standard", see
OpenMP's cp_parser_omp_parallel, cp_parser_omp_for, and many more.
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-09 14:00 ` Thomas Schwinge
@ 2015-10-11 22:19 ` Nathan Sidwell
0 siblings, 0 replies; 9+ messages in thread
From: Nathan Sidwell @ 2015-10-11 22:19 UTC (permalink / raw)
To: Thomas Schwinge
Cc: Jakub Jelinek, James Norris, Cesar Philippidis, gcc-patches,
Jason Merrill, Joseph Myers
On 10/09/15 09:59, Thomas Schwinge wrote:
> It's s string describing the pragma as parsed thus far. Again, not
> documenting that as well as our usage of it is totally "standard", see
> OpenMP's cp_parser_omp_parallel, cp_parser_omp_for, and many more.
Ok, I'm not going to hold this to a higher than normal standard then. C++ bits
ok, thanks.
nathan
^ permalink raw reply [flat|nested] 9+ messages in thread
* Re: [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++
2015-10-08 16:39 [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++ Thomas Schwinge
2015-10-08 16:47 ` Joseph Myers
2015-10-09 12:26 ` Nathan Sidwell
@ 2015-10-27 8:52 ` Thomas Schwinge
2 siblings, 0 replies; 9+ messages in thread
From: Thomas Schwinge @ 2015-10-27 8:52 UTC (permalink / raw)
To: gcc-patches
Cc: Jakub Jelinek, James Norris, Cesar Philippidis, Joseph Myers,
Jason Merrill, Nathan Sidwell
[-- Attachment #1: Type: text/plain, Size: 42009 bytes --]
Hi!
On Thu, 8 Oct 2015 18:39:25 +0200, I wrote:
> Some bits extracted out of gomp-4_0-branch, and some other bits
> rewritten; here is a patch to support OpenACC Combined Directives in C,
> C++. (The Fortran front end already does support these.)
>
> As far as I know, Jakub is not available at this time, so maybe the C
> (Joseph) and C++ (Jason, Nathan) front end maintainers could please
> review this, instead of him? (The front end changes as well as the few
> other cleanup changes should all be straight forward.) OK for trunk once
> bootstrap tested?
Given the approval by Joseph and Nathan (thanks!), I have now committed
this, in r229404:
commit 2c4c872529d6e1051fb9cdd641ad5cc9fbc0e2cb
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue Oct 27 08:39:15 2015 +0000
[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.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@229404 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 6 +
gcc/c-family/ChangeLog | 9 ++
gcc/c-family/c-common.h | 1 +
gcc/c-family/c-omp.c | 40 +++++-
gcc/c/ChangeLog | 19 +++
gcc/c/c-parser.c | 148 +++++++++-----------
gcc/c/c-tree.h | 3 +-
gcc/c/c-typeck.c | 36 ++---
gcc/cp/ChangeLog | 18 +++
gcc/cp/cp-tree.h | 3 +-
gcc/cp/parser.c | 149 ++++++++++++---------
gcc/cp/semantics.c | 54 +++-----
gcc/fortran/trans-openmp.c | 4 -
gcc/gimplify.c | 16 +--
gcc/testsuite/ChangeLog | 13 ++
gcc/testsuite/c-c++-common/goacc/loop-1.c | 10 +-
gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 | 1 +
gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 | 9 +-
gcc/testsuite/gfortran.dg/goacc/cray.f95 | 2 +-
gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 2 +-
gcc/tree-pretty-print.c | 11 +-
gcc/tree.def | 8 +-
gcc/tree.h | 20 ---
libgomp/ChangeLog | 5 +
.../libgomp.oacc-c-c++-common/combdir-1.c | 52 +++++++
.../testsuite/libgomp.oacc-fortran/combdir-1.f90 | 37 +++++
26 files changed, 407 insertions(+), 269 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index bad27b5..20fd5ec 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,9 @@
+2015-10-27 Thomas Schwinge <thomas@codesourcery.com>
+
+ * 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.
+
2015-10-27 Tom de Vries <tom@codesourcery.com>
* tree-ssa-structalias.c (push_fields_onto_fieldstack): Add and use var
diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog
index f2beefd..4d22572 100644
--- gcc/c-family/ChangeLog
+++ gcc/c-family/ChangeLog
@@ -1,3 +1,12 @@
+2015-10-27 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ 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.
+
2015-10-21 Martin Sebor <msebor@redhat.com>
PR driver/68043
diff --git gcc/c-family/c-common.h gcc/c-family/c-common.h
index 4b5cac8..6a381f2 100644
--- gcc/c-family/c-common.h
+++ gcc/c-family/c-common.h
@@ -1269,6 +1269,7 @@ extern void c_finish_omp_taskyield (location_t);
extern tree c_finish_omp_for (location_t, enum tree_code, tree, tree, tree,
tree, tree, tree, tree);
extern tree c_finish_oacc_wait (location_t, tree, tree);
+extern tree c_oacc_split_loop_clauses (tree, tree *);
extern void c_omp_split_clauses (location_t, enum tree_code, omp_clause_mask,
tree, tree *);
extern tree c_omp_declare_simd_clauses_to_numbers (tree, tree);
diff --git gcc/c-family/c-omp.c gcc/c-family/c-omp.c
index 36f9b66..93bff2e 100644
--- gcc/c-family/c-omp.c
+++ gcc/c-family/c-omp.c
@@ -691,13 +691,47 @@ c_finish_omp_for (location_t locus, enum tree_code code, tree declv,
}
}
-/* Right now we have 21 different combined/composite constructs, this
- function attempts to split or duplicate clauses for combined
+/* This function splits clauses for OpenACC combined loop
+ constructs. OpenACC combined loop constructs are:
+ #pragma acc kernels loop
+ #pragma acc parallel loop
+*/
+
+tree
+c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses)
+{
+ tree next, loop_clauses;
+
+ loop_clauses = *not_loop_clauses = NULL_TREE;
+ for (; clauses ; clauses = next)
+ {
+ next = OMP_CLAUSE_CHAIN (clauses);
+
+ switch (OMP_CLAUSE_CODE (clauses))
+ {
+ case OMP_CLAUSE_COLLAPSE:
+ case OMP_CLAUSE_REDUCTION:
+ OMP_CLAUSE_CHAIN (clauses) = loop_clauses;
+ loop_clauses = clauses;
+ break;
+
+ default:
+ OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses;
+ *not_loop_clauses = clauses;
+ break;
+ }
+ }
+
+ return loop_clauses;
+}
+
+/* This function attempts to split or duplicate clauses for OpenMP
+ combined/composite constructs. Right now there are 21 different
constructs. CODE is the innermost construct in the combined construct,
and MASK allows to determine which constructs are combined together,
as every construct has at least one clause that no other construct
has (except for OMP_SECTIONS, but that can be only combined with parallel).
- Combined/composite constructs are:
+ OpenMP combined/composite constructs are:
#pragma omp distribute parallel for
#pragma omp distribute parallel for simd
#pragma omp distribute simd
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index aed7a62..da4c801 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,22 @@
+2015-10-27 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ 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.
+
2015-10-22 Richard Biener <rguenther@suse.de>
* c-typeck.c (c_finish_omp_clauses): Properly convert operands
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index c8c6a2d..ec918c1 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1237,7 +1237,6 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
unsigned int * = NULL);
static void c_parser_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
-static tree c_parser_oacc_loop (location_t, c_parser *, char *);
static void c_parser_omp_construct (c_parser *);
static void c_parser_omp_threadprivate (c_parser *);
static void c_parser_omp_barrier (c_parser *);
@@ -12875,60 +12874,6 @@ 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_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
-c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
-{
- tree stmt, clauses = NULL_TREE, block;
-
- strcat (p_name, " kernels");
-
- if (c_parser_next_token_is (parser, CPP_NAME))
- {
- const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
- if (strcmp (p, "loop") == 0)
- {
- c_parser_consume_token (parser);
- block = c_begin_omp_parallel ();
- c_parser_oacc_loop (loc, parser, p_name);
- stmt = c_finish_oacc_kernels (loc, clauses, block);
- OACC_KERNELS_COMBINED (stmt) = 1;
- return stmt;
- }
- }
-
- clauses = c_parser_oacc_all_clauses (parser, OACC_KERNELS_CLAUSE_MASK,
- p_name);
-
- 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 enter data oacc-enter-data-clause[optseq] new-line
or
@@ -13018,16 +12963,25 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) )
static tree
-c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
+c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
+ omp_clause_mask mask, tree *cclauses)
{
- tree stmt, clauses, block;
-
strcat (p_name, " loop");
+ mask |= OACC_LOOP_CLAUSE_MASK;
- clauses = c_parser_oacc_all_clauses (parser, OACC_LOOP_CLAUSE_MASK, p_name);
+ tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name,
+ cclauses == NULL);
+ if (cclauses)
+ {
+ clauses = c_oacc_split_loop_clauses (clauses, cclauses);
+ if (*cclauses)
+ c_finish_omp_clauses (*cclauses, false);
+ if (clauses)
+ c_finish_omp_clauses (clauses, false);
+ }
- block = c_begin_compound_stmt (true);
- stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
+ tree block = c_begin_compound_stmt (true);
+ tree stmt = c_parser_omp_for_loop (loc, parser, OACC_LOOP, clauses, NULL);
block = c_end_compound_stmt (loc, block, true);
add_stmt (block);
@@ -13035,12 +12989,32 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
}
/* 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
LOC is the location of the #pragma token.
*/
+#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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
@@ -13061,11 +13035,26 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static tree
-c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
+c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser,
+ enum pragma_kind p_kind, char *p_name)
{
- tree stmt, clauses = NULL_TREE, block;
-
- strcat (p_name, " parallel");
+ omp_clause_mask mask;
+ enum tree_code code;
+ switch (p_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 ();
+ }
if (c_parser_next_token_is (parser, CPP_NAME))
{
@@ -13073,23 +13062,21 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
if (strcmp (p, "loop") == 0)
{
c_parser_consume_token (parser);
- block = c_begin_omp_parallel ();
- c_parser_oacc_loop (loc, parser, p_name);
- stmt = c_finish_oacc_parallel (loc, clauses, block);
- OACC_PARALLEL_COMBINED (stmt) = 1;
- return stmt;
+ mask |= OACC_LOOP_CLAUSE_MASK;
+
+ tree block = c_begin_omp_parallel ();
+ tree clauses;
+ c_parser_oacc_loop (loc, parser, p_name, mask, &clauses);
+ return c_finish_omp_construct (loc, code, block, clauses);
}
}
- clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
- p_name);
+ tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name);
- block = c_begin_omp_parallel ();
+ tree block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
- stmt = c_finish_oacc_parallel (loc, clauses, block);
-
- return stmt;
+ return c_finish_omp_construct (loc, code, block, clauses);
}
/* OpenACC 2.0:
@@ -16079,16 +16066,13 @@ c_parser_omp_construct (c_parser *parser)
stmt = c_parser_oacc_data (loc, parser);
break;
case PRAGMA_OACC_KERNELS:
+ case PRAGMA_OACC_PARALLEL:
strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_kernels (loc, parser, p_name);
+ stmt = c_parser_oacc_kernels_parallel (loc, parser, p_kind, p_name);
break;
case PRAGMA_OACC_LOOP:
strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_loop (loc, parser, p_name);
- break;
- case PRAGMA_OACC_PARALLEL:
- strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_parallel (loc, parser, p_name);
+ stmt = c_parser_oacc_loop (loc, parser, p_name, mask, NULL);
break;
case PRAGMA_OACC_WAIT:
strcpy (p_name, "#pragma wait");
diff --git gcc/c/c-tree.h gcc/c/c-tree.h
index bee03d3..04991f7 100644
--- gcc/c/c-tree.h
+++ gcc/c/c-tree.h
@@ -640,8 +640,7 @@ extern tree c_finish_bc_stmt (location_t, tree *, bool);
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_omp_construct (location_t, enum tree_code, 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 61c5313..64ea1c2 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11481,39 +11481,19 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se)
return expr;
}
\f
-/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_PARALLEL. */
+/* Generate OMP construct CODE, with BODY and CLAUSES as its compound
+ statement. LOC is the location of the construct. */
tree
-c_finish_oacc_parallel (location_t loc, tree clauses, tree block)
+c_finish_omp_construct (location_t loc, enum tree_code code, tree body,
+ tree clauses)
{
- tree stmt;
+ body = c_end_compound_stmt (loc, body, true);
- block = c_end_compound_stmt (loc, block, true);
-
- stmt = make_node (OACC_PARALLEL);
- TREE_TYPE (stmt) = void_type_node;
- OACC_PARALLEL_CLAUSES (stmt) = clauses;
- OACC_PARALLEL_BODY (stmt) = block;
- SET_EXPR_LOCATION (stmt, loc);
-
- 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 stmt = make_node (code);
TREE_TYPE (stmt) = void_type_node;
- OACC_KERNELS_CLAUSES (stmt) = clauses;
- OACC_KERNELS_BODY (stmt) = block;
+ OMP_BODY (stmt) = body;
+ OMP_CLAUSES (stmt) = clauses;
SET_EXPR_LOCATION (stmt, loc);
return add_stmt (stmt);
diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog
index e900e29..979ebb9 100644
--- gcc/cp/ChangeLog
+++ gcc/cp/ChangeLog
@@ -1,3 +1,21 @@
+2015-10-27 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ 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.
+
2015-10-25 Jason Merrill <jason@redhat.com>
DR 2179
diff --git gcc/cp/cp-tree.h gcc/cp/cp-tree.h
index 16db41f..af2ba64 100644
--- gcc/cp/cp-tree.h
+++ gcc/cp/cp-tree.h
@@ -6318,8 +6318,7 @@ extern void finish_omp_threadprivate (tree);
extern tree begin_omp_structured_block (void);
extern tree finish_omp_structured_block (tree);
extern tree finish_oacc_data (tree, tree);
-extern tree finish_oacc_kernels (tree, tree);
-extern tree finish_oacc_parallel (tree, tree);
+extern tree finish_omp_construct (enum tree_code, tree, tree);
extern tree begin_omp_parallel (void);
extern tree finish_omp_parallel (tree, tree);
extern tree begin_omp_task (void);
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 7555bf3..0354029 100644
--- gcc/cp/parser.c
+++ 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);
- 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_kernels_parallel (parser, pragma_tok, p_name);
+ break;
+ case PRAGMA_OACC_LOOP:
+ 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);
diff --git gcc/cp/semantics.c gcc/cp/semantics.c
index 11315d9..f9e86d0 100644
--- gcc/cp/semantics.c
+++ gcc/cp/semantics.c
@@ -7106,8 +7106,17 @@ finish_omp_structured_block (tree block)
return do_poplevel (block);
}
+/* Similarly, except force the retention of the BLOCK. */
+
+tree
+begin_omp_parallel (void)
+{
+ keep_next_level (true);
+ return begin_omp_structured_block ();
+}
+
/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_DATA. */
+ statement. */
tree
finish_oacc_data (tree clauses, tree block)
@@ -7124,51 +7133,22 @@ finish_oacc_data (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. */
+/* Generate OMP construct CODE, with BODY and CLAUSES as its compound
+ statement. */
tree
-finish_oacc_kernels (tree clauses, tree block)
+finish_omp_construct (enum tree_code code, tree body, tree clauses)
{
- tree stmt;
+ body = finish_omp_structured_block (body);
- block = finish_omp_structured_block (block);
-
- stmt = make_node (OACC_KERNELS);
- TREE_TYPE (stmt) = void_type_node;
- OACC_KERNELS_CLAUSES (stmt) = clauses;
- OACC_KERNELS_BODY (stmt) = block;
-
- return add_stmt (stmt);
-}
-
-/* Generate OACC_PARALLEL, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_PARALLEL. */
-
-tree
-finish_oacc_parallel (tree clauses, tree block)
-{
- tree stmt;
-
- block = finish_omp_structured_block (block);
-
- stmt = make_node (OACC_PARALLEL);
+ tree stmt = make_node (code);
TREE_TYPE (stmt) = void_type_node;
- OACC_PARALLEL_CLAUSES (stmt) = clauses;
- OACC_PARALLEL_BODY (stmt) = block;
+ OMP_BODY (stmt) = body;
+ OMP_CLAUSES (stmt) = clauses;
return add_stmt (stmt);
}
-/* Similarly, except force the retention of the BLOCK. */
-
-tree
-begin_omp_parallel (void)
-{
- keep_next_level (true);
- return begin_omp_structured_block ();
-}
-
tree
finish_omp_parallel (tree clauses, tree body)
{
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index b6284c9..def8afb 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -3463,10 +3463,6 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
poplevel (0, 0);
stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
oacc_clauses);
- if (code->op == EXEC_OACC_KERNELS_LOOP)
- OACC_KERNELS_COMBINED (stmt) = 1;
- else
- OACC_PARALLEL_COMBINED (stmt) = 1;
gfc_add_expr_to_block (&block, stmt);
return gfc_finish_block (&block);
}
diff --git gcc/gimplify.c gcc/gimplify.c
index ab9e540..a2d71d2 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -9600,23 +9600,9 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_DATA:
case OACC_KERNELS:
- if (OACC_KERNELS_COMBINED (*expr_p))
- sorry ("directive not yet implemented");
- else
- gimplify_omp_workshare (expr_p, pre_p);
- ret = GS_ALL_DONE;
- break;
-
case OACC_PARALLEL:
- if (OACC_PARALLEL_COMBINED (*expr_p))
- sorry ("directive not yet implemented");
- else
- gimplify_omp_workshare (expr_p, pre_p);
- ret = GS_ALL_DONE;
- break;
-
- case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
case OMP_TARGET:
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 7f4f8a0..e1cb93c 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,16 @@
+2015-10-27 Thomas Schwinge <thomas@codesourcery.com>
+
+ 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.
+
2015-10-26 Louis Krupp <louis.krupp@zoho.com>
PR fortran/66056
diff --git gcc/testsuite/c-c++-common/goacc/loop-1.c gcc/testsuite/c-c++-common/goacc/loop-1.c
index 2d8f4d5..4cf5429 100644
--- gcc/testsuite/c-c++-common/goacc/loop-1.c
+++ gcc/testsuite/c-c++-common/goacc/loop-1.c
@@ -1,5 +1,3 @@
-/* { dg-skip-if "not yet" { c++ } } */
-
int test1()
{
int i, j, k, b[10];
@@ -69,4 +67,10 @@ int test1()
}
return 0;
}
-/* { dg-prune-output "sorry, unimplemented: directive not yet implemented" } */
+
+// PR64765
+void PR64765(float *f, double *r) {
+ int i;
+ #pragma acc kernels loop create(f) copy(r)
+ for(i = 64; i < 76; i += 5) {}
+}
diff --git gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
index f35d4b9..f9cf9ac 100644
--- gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
+++ gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
@@ -2,6 +2,7 @@
! { dg-additional-options "-fcoarray=lib" }
!
! PR fortran/63861
+! { dg-xfail-if "<http://gcc.gnu.org/PR63861>" { *-*-* } }
module test
contains
diff --git gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
index b8be649..abb10f9 100644
--- gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
+++ gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
@@ -6,7 +6,14 @@ subroutine oacc1()
implicit none
integer :: i
integer :: a
- !$acc parallel loop reduction(+:a) ! { dg-excess-errors "sorry, unimplemented: directive not yet implemented" }
+ !$acc parallel loop reduction(+:a)
do i = 1,5
enddo
+ !$acc end parallel loop
+ !$acc kernels loop collapse(2)
+ do i = 2,6
+ do a = 3,5
+ enddo
+ enddo
+ !$acc end kernels loop
end subroutine oacc1
diff --git gcc/testsuite/gfortran.dg/goacc/cray.f95 gcc/testsuite/gfortran.dg/goacc/cray.f95
index 8f2c077..52789fe 100644
--- gcc/testsuite/gfortran.dg/goacc/cray.f95
+++ gcc/testsuite/gfortran.dg/goacc/cray.f95
@@ -53,4 +53,4 @@ contains
!$acc update self (ptr)
end subroutine oacc1
end module test
-! { dg-prune-output "unimplemented" }
+! { dg-prune-output "ACC cache unimplemented" }
diff --git gcc/testsuite/gfortran.dg/goacc/parameter.f95 gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 1364181..454924a 100644
--- gcc/testsuite/gfortran.dg/goacc/parameter.f95
+++ gcc/testsuite/gfortran.dg/goacc/parameter.f95
@@ -29,4 +29,4 @@ contains
!$acc update self (a) ! { dg-error "not a variable" }
end subroutine oacc1
end module test
-! { dg-prune-output "unimplemented" }
+! { dg-prune-output "ACC cache unimplemented" }
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index 11f90051..b1f76d5 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -2676,13 +2676,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, int flags,
case OACC_PARALLEL:
pp_string (pp, "#pragma acc parallel");
- dump_omp_clauses (pp, OACC_PARALLEL_CLAUSES (node), spc, flags);
- goto dump_omp_body;
+ goto dump_omp_clauses_body;
case OACC_KERNELS:
pp_string (pp, "#pragma acc kernels");
- dump_omp_clauses (pp, OACC_KERNELS_CLAUSES (node), spc, flags);
- goto dump_omp_body;
+ goto dump_omp_clauses_body;
case OACC_DATA:
pp_string (pp, "#pragma acc data");
@@ -2722,6 +2720,11 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, int flags,
case OMP_PARALLEL:
pp_string (pp, "#pragma omp parallel");
dump_omp_clauses (pp, OMP_PARALLEL_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ dump_omp_clauses_body:
+ dump_omp_clauses (pp, OMP_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
dump_omp_body:
if (!(flags & TDF_SLIM) && OMP_BODY (node))
diff --git gcc/tree.def gcc/tree.def
index d0a3bd6..fc7490a 100644
--- gcc/tree.def
+++ gcc/tree.def
@@ -1047,14 +1047,14 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
not change the ordering of these codes. */
/* OpenACC - #pragma acc parallel [clause1 ... clauseN]
- Operand 0: OACC_PARALLEL_BODY: Code to be executed in parallel.
- Operand 1: OACC_PARALLEL_CLAUSES: List of clauses. */
+ Operand 0: OMP_BODY: Code to be executed in parallel.
+ Operand 1: OMP_CLAUSES: List of clauses. */
DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
/* OpenACC - #pragma acc kernels [clause1 ... clauseN]
- Operand 0: OACC_KERNELS_BODY: Sequence of kernels.
- Operand 1: OACC_KERNELS_CLAUSES: List of clauses. */
+ Operand 0: OMP_BODY: Sequence of kernels.
+ Operand 1: OMP_CLAUSES: List of clauses. */
DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
diff --git gcc/tree.h gcc/tree.h
index ece083b..65c3117 100644
--- gcc/tree.h
+++ gcc/tree.h
@@ -1,4 +1,3 @@
-
/* Definitions for the ubiquitous 'tree' type for GNU compilers.
Copyright (C) 1989-2015 Free Software Foundation, Inc.
@@ -1218,16 +1217,6 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_STANDALONE_CLAUSES(NODE) \
TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_CACHE, OMP_TARGET_EXIT_DATA), 0)
-#define OACC_PARALLEL_BODY(NODE) \
- TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 0)
-#define OACC_PARALLEL_CLAUSES(NODE) \
- TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1)
-
-#define OACC_KERNELS_BODY(NODE) \
- TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 0)
-#define OACC_KERNELS_CLAUSES(NODE) \
- TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 1)
-
#define OACC_DATA_BODY(NODE) \
TREE_OPERAND (OACC_DATA_CHECK (NODE), 0)
#define OACC_DATA_CLAUSES(NODE) \
@@ -1332,15 +1321,6 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_SECTION_LAST(NODE) \
(OMP_SECTION_CHECK (NODE)->base.private_flag)
-/* True on an OACC_KERNELS statement if is represents combined kernels loop
- directive. */
-#define OACC_KERNELS_COMBINED(NODE) \
- (OACC_KERNELS_CHECK (NODE)->base.private_flag)
-
-/* Like OACC_KERNELS_COMBINED, but for parallel loop directive. */
-#define OACC_PARALLEL_COMBINED(NODE) \
- (OACC_PARALLEL_CHECK (NODE)->base.private_flag)
-
/* True on an OMP_PARALLEL statement if it represents an explicit
combined parallel work-sharing constructs. */
#define OMP_PARALLEL_COMBINED(NODE) \
diff --git libgomp/ChangeLog libgomp/ChangeLog
index afc49ae..ca34af8 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2015-10-27 James Norris <jnorris@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/combdir-1.c: New file.
+ * testsuite/libgomp.oacc-fortran/combdir-1.f90: Likewise.
+
2015-10-26 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/abort-1.c: Print to stderr.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
new file mode 100644
index 0000000..a7def92
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
@@ -0,0 +1,52 @@
+/* { dg-do run } */
+
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+ const int N = 32;
+ float a[N], b[N];
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 1.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc parallel loop copy (a[0:N]) copy (b[0:N])
+ for (i = 0; i < N; i++)
+ {
+ b[i] = 2.0;
+ a[i] = a[i] + b[i];
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 2.0)
+ abort ();
+ }
+
+#pragma acc kernels loop copy (a[0:N]) copy (b[0:N])
+ for (i = 0; i < N; i++)
+ {
+ b[i] = 3.0;
+ a[i] = a[i] + b[i];
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ return 0;
+
+}
diff --git libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90 libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
new file mode 100644
index 0000000..0cd8a67
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
@@ -0,0 +1,37 @@
+! { dg-do run }
+
+program main
+ integer, parameter :: n = 32
+ real :: a(n), b(n);
+ integer :: i
+
+ do i = 1, n
+ a(i) = 1.0
+ b(i) = 0.0
+ end do
+
+ !$acc parallel loop copy (a(1:n)) copy (b(1:n))
+ do i = 1, n
+ b(i) = 2.0
+ a(i) = a(i) + b(i)
+ end do
+
+ do i = 1, n
+ if (a(i) .ne. 3.0) call abort
+
+ if (b(i) .ne. 2.0) call abort
+ end do
+
+ !$acc kernels loop copy (a(1:n)) copy (b(1:n))
+ do i = 1, n
+ b(i) = 3.0;
+ a(i) = a(i) + b(i)
+ end do
+
+ do i = 1, n
+ if (a(i) .ne. 6.0) call abort
+
+ if (b(i) .ne. 3.0) call abort
+ end do
+
+end program main
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 9+ messages in thread
end of thread, other threads:[~2015-10-27 8:40 UTC | newest]
Thread overview: 9+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-10-08 16:39 [PR c/64765, c/64880] Support OpenACC Combined Directives in C, C++ Thomas Schwinge
2015-10-08 16:47 ` Joseph Myers
2015-10-09 12:26 ` Nathan Sidwell
2015-10-09 13:26 ` Thomas Schwinge
2015-10-09 13:34 ` Nathan Sidwell
2015-10-09 13:57 ` Jakub Jelinek
2015-10-09 14:00 ` Thomas Schwinge
2015-10-11 22:19 ` Nathan Sidwell
2015-10-27 8:52 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).