summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--gcc/ChangeLog6
-rw-r--r--gcc/c-family/ChangeLog9
-rw-r--r--gcc/c-family/c-common.h1
-rw-r--r--gcc/c-family/c-omp.c40
-rw-r--r--gcc/c/ChangeLog19
-rw-r--r--gcc/c/c-parser.c148
-rw-r--r--gcc/c/c-tree.h3
-rw-r--r--gcc/c/c-typeck.c36
-rw-r--r--gcc/cp/ChangeLog18
-rw-r--r--gcc/cp/cp-tree.h3
-rw-r--r--gcc/cp/parser.c147
-rw-r--r--gcc/cp/semantics.c54
-rw-r--r--gcc/fortran/trans-openmp.c4
-rw-r--r--gcc/gimplify.c16
-rw-r--r--gcc/testsuite/ChangeLog13
-rw-r--r--gcc/testsuite/c-c++-common/goacc/loop-1.c10
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/coarray_2.f901
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/combined_loop.f909
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/cray.f952
-rw-r--r--gcc/testsuite/gfortran.dg/goacc/parameter.f952
-rw-r--r--gcc/tree-pretty-print.c11
-rw-r--r--gcc/tree.def8
-rw-r--r--gcc/tree.h20
-rw-r--r--libgomp/ChangeLog5
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c52
-rw-r--r--libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f9037
26 files changed, 406 insertions, 268 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index bad27b54f91..20fd5ececa7 100644
--- a/gcc/ChangeLog
+++ b/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 a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index f2beefdd5c4..4d225726c56 100644
--- a/gcc/c-family/ChangeLog
+++ b/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 a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index 4b5cac85446..6a381f23435 100644
--- a/gcc/c-family/c-common.h
+++ b/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 a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 36f9b6627e5..93bff2e2ec5 100644
--- a/gcc/c-family/c-omp.c
+++ b/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 a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index aed7a627932..da4c8016834 100644
--- a/gcc/c/ChangeLog
+++ b/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 a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index c8c6a2dff90..ec918c14ca8 100644
--- a/gcc/c/c-parser.c
+++ b/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 a/gcc/c/c-tree.h b/gcc/c/c-tree.h
index bee03d3dc45..04991f7261c 100644
--- a/gcc/c/c-tree.h
+++ b/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 a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 61c5313c25b..64ea1c272f8 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -11481,39 +11481,19 @@ c_expr_to_decl (tree expr, bool *tc ATTRIBUTE_UNUSED, bool *se)
return expr;
}
-/* 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;
-
- 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);
+ body = c_end_compound_stmt (loc, body, 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 a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index e900e297494..979ebb96171 100644
--- a/gcc/cp/ChangeLog
+++ b/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 a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 16db41f79e9..af2ba64409c 100644
--- a/gcc/cp/cp-tree.h
+++ b/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 a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7555bf3c87a..0354029a060 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);
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 11315d90668..f9e86d09b43 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7106,69 +7106,49 @@ finish_omp_structured_block (tree block)
return do_poplevel (block);
}
-/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_DATA. */
+/* Similarly, except force the retention of the BLOCK. */
tree
-finish_oacc_data (tree clauses, tree block)
+begin_omp_parallel (void)
{
- tree stmt;
-
- block = finish_omp_structured_block (block);
-
- stmt = make_node (OACC_DATA);
- TREE_TYPE (stmt) = void_type_node;
- OACC_DATA_CLAUSES (stmt) = clauses;
- OACC_DATA_BODY (stmt) = block;
-
- return add_stmt (stmt);
+ keep_next_level (true);
+ return begin_omp_structured_block ();
}
-/* Generate OACC_KERNELS, with CLAUSES and BLOCK as its compound
- statement. LOC is the location of the OACC_KERNELS. */
+/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
+ statement. */
tree
-finish_oacc_kernels (tree clauses, tree block)
+finish_oacc_data (tree clauses, tree block)
{
tree stmt;
block = finish_omp_structured_block (block);
- stmt = make_node (OACC_KERNELS);
+ stmt = make_node (OACC_DATA);
TREE_TYPE (stmt) = void_type_node;
- OACC_KERNELS_CLAUSES (stmt) = clauses;
- OACC_KERNELS_BODY (stmt) = block;
+ OACC_DATA_CLAUSES (stmt) = clauses;
+ OACC_DATA_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. */
+/* Generate OMP construct CODE, with BODY and CLAUSES as its compound
+ statement. */
tree
-finish_oacc_parallel (tree clauses, tree block)
+finish_omp_construct (enum tree_code code, tree body, tree clauses)
{
- tree stmt;
-
- block = finish_omp_structured_block (block);
+ body = finish_omp_structured_block (body);
- 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 a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index b6284c9c957..def8afb4d6d 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/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 a/gcc/gimplify.c b/gcc/gimplify.c
index ab9e540e76d..a2d71d2f958 100644
--- a/gcc/gimplify.c
+++ b/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 a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 7f4f8a09941..e1cb93ca871 100644
--- a/gcc/testsuite/ChangeLog
+++ b/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 a/gcc/testsuite/c-c++-common/goacc/loop-1.c b/gcc/testsuite/c-c++-common/goacc/loop-1.c
index 2d8f4d5a8d3..4cf5429c9bb 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-1.c
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90 b/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
index f35d4b9b18e..f9cf9ac6ca0 100644
--- a/gcc/testsuite/gfortran.dg/goacc/coarray_2.f90
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90 b/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
index b8be6499937..abb10f93067 100644
--- a/gcc/testsuite/gfortran.dg/goacc/combined_loop.f90
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/cray.f95 b/gcc/testsuite/gfortran.dg/goacc/cray.f95
index 8f2c077beee..52789fe2998 100644
--- a/gcc/testsuite/gfortran.dg/goacc/cray.f95
+++ b/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 a/gcc/testsuite/gfortran.dg/goacc/parameter.f95 b/gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 1364181b33c..454924aed3d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parameter.f95
+++ b/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 a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 11f90051107..b1f76d52c33 100644
--- a/gcc/tree-pretty-print.c
+++ b/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 a/gcc/tree.def b/gcc/tree.def
index d0a3bd68398..fc7490a48d8 100644
--- a/gcc/tree.def
+++ b/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 a/gcc/tree.h b/gcc/tree.h
index ece083bed3e..65c3117b5dd 100644
--- a/gcc/tree.h
+++ b/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 a/libgomp/ChangeLog b/libgomp/ChangeLog
index afc49ae164c..ca34af8203c 100644
--- a/libgomp/ChangeLog
+++ b/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 a/libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
new file mode 100644
index 00000000000..a7def920b94
--- /dev/null
+++ b/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 a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
new file mode 100644
index 00000000000..0cd8a670abb
--- /dev/null
+++ b/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
View Lorry Controller Status