This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4] acc enter/exit data
- From: Cesar Philippidis <cesar at codesourcery dot com>
- To: "gcc-patches at gcc dot gnu dot org" <gcc-patches at gcc dot gnu dot org>, <fortran at gcc dot gnu dot org>
- Date: Thu, 30 Oct 2014 17:11:04 -0700
- Subject: [gomp4] acc enter/exit data
- Authentication-results: sourceware.org; auth=none
This patch add support for OpenACC's enter/exit data directive. Note
that there is a problem in the 2.0a spec regarding the live ranges of
variables in data clauses. Section 2.6.5.7 states that exit data delete
should deallocate memory without writing it back. However, that may
conflict with an acc data variable as the following example demonstrates.
#pragma acc data copy (A)
{
...
#pragma acc exit data delete (A)
...
} // end of acc data block
The OpenACC technical committee has informed me that this issue has been
corrected in a future revision of OpenACC. For now though, acc exit data
delete will decrement A's refcount and the GC will delete it when it's
no longer necessary. To be clear, this example will result in a runtime
failure at when the acc data block terminates.
One note regarding the mystery 3 refcount in gomp_acc_remove_pointer.
When gomp_acc_insert_pointer creates a mapping for a pset, the array
data itself has three references: (1) the data itself, (2) the pointer,
and (3) the pset. However, when it comes time to deleting the pset,
gomp_acc_remove_pointer is really removing the data itself. The mystery
2 argument comes from acc_unmap_vars. I suspect a similar argument can
be used for 2 which is only used for pointers (really subarrays): (1)
data, and (2) pointer.
Thomas has already approved this patch internally, so I'll commit it to
gomp-4_0-branch in the next few days unless someone complains.
Thanks,
Cesar
2014-10-30 Cesar Philippidis <cesar@codesourcery.com>
gcc/c-family/
* c-pragma.c (oacc_pragmas): Add entries for PRAGMA_OACC_ENTER_DATA
and PRAGMA_OACC_EXIT_DATA.
* c-pragma.h (pragma_kind): Likewise.
gcc/c/
* c-parser.c (c_parser_oacc_enter_exit_data): New function.
(c_parser_pragma): Handle PRAGMA_OACC_ENTER_DATA and
PRAGMA_OACC_EXIT_DATA.
(OACC_ENTER_DATA_CLAUSE_MASK): New macro.
(OACC_EXIT_DATA_CLAUSE_MASK): New macro.
(c_parser_oacc_update): Don't create a new stmt if the pragma
is bogus.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Also consider CPP_KEYWORD
typed tokens as clauses for delete.
(OACC_ENTER_DATA_CLAUSE_MASK): New macro.
(OACC_EXIT_DATA_CLAUSE_MASK): New macro.
(cp_parser_oacc_enter_exit_data): New function.
(cp_parser_omp_construct): Handle PRAGMA_OACC_ENTER_DATA and
PRAGMA_OACC_EXIT_DATA.
(cp_parser_pragma): Likewise.
gcc/fortran/
* gfortran.h (enum OMP_LIST_HOST): Remove.
(enum OMP_LIST_DEVICE, OMP_LIST_DEVICE): Remove.
* dump-parse-tree.c (show_omp_clauses): Remove OMP_LIST_HOST and
OMP_LIST_DEVICE from here also.
* openmp.c (OMP_CLAUSE_SELF): New define.
(gfc_match_omp_clauses): Update handling of OMP_CLAUSE_HOST and
OMP_CLAUSE_DEVICE. Add support for OMP_CLAUSE_SELF.
* trans-openmp.c (gfc_trans_omp_clauses): Remove support for
OMP_LIST_HOST and OMP_LIST_DEVICE since they are treated as memory
maps now.
(gfc_trans_oacc_executable_directive): Remove stale EXEC_OACC_WAIT.
gcc/
* gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
* gimple-pretty-print.c (dump_gimple_omp_target): Handle it.
* gimplify.c (gimplify_scan_omp_clauses): Remove switch stmt which
declared OMP_CLAUSE_MAP_FORCE_DEALLOC as unimplemented.
(gimplify_omp_target_update): Handle OACC_ENTER_DATA and
OACC_EXIT_DATA.
(gimplify_expr): Shuffle around OACC_ENTER_DATA, OACC_EXIT_DATA and
OACC_WAIT.
* oacc-builtins.def (BUILD_INT_GOACC_ENTER_EXIT_DATA): New built-in
function.
* omp-low.c (expand_omp_target): Handle
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA. Don't use quick_push when
there is an unknown number of wait args.
(lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
gcc/testsuite/
* c-c++-common/goacc/data-1.c: Exercise enter/exit data pragma.
* c-c++-common/goacc/update-1.c: Ensure that fortran subarrays err.
* gcc/testsuite/c-c++-common/goacc/data-2.c: New test.
* gcc/testsuite/c-c++-common/goacc/update-1.c: Check for malformed
subarrays.
libgomp/
* libgomp.map (GOACC_enter_exit_data): Declare as global.
* libgomp_g.h (GOACC_enter_exit_data): Declare.
(GOACC_update): Declare.
(gomp_acc_insert_pointer): Declare.
(gomp_acc_remove_pointer): Declare.
* oacc-mem.c (gomp_acc_insert_pointer): New function.
(gomp_acc_remove_pointer): New function.
* oacc-parallel.c (find_pset): New function.
(GOACC_enter_exit_data): New function.
(GOACC_update): Handle GOMP_MAP_TO_PSET.
* testsuite/libgomp.oacc-c++/c++.exp (check_efective_target_oacc_c):
New proc.
* testsuite/libgomp.oacc-c-c++-common/data-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/data-3.c: New test.
* testsuite/libgomp.oacc-c/c.exp (check_efective_target_oacc_c):
New proc.
* testsuite/libgomp.oacc-fortran/data-1.f90: New test.
* testsuite/libgomp.oacc-fortran/data-2.f90: New test.
* testsuite/libgomp.oacc-fortran/data-3.f90: New test.
* testsuite/libgomp.oacc-fortran/data-4.f90: New test.
diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index 39634ea..e98b555 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1183,6 +1183,8 @@ struct omp_pragma_def { const char *name; unsigned int id; };
static const struct omp_pragma_def oacc_pragmas[] = {
{ "cache", PRAGMA_OACC_CACHE },
{ "data", PRAGMA_OACC_DATA },
+ { "enter", PRAGMA_OACC_ENTER_DATA },
+ { "exit", PRAGMA_OACC_EXIT_DATA },
{ "kernels", PRAGMA_OACC_KERNELS },
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 4722d51..d495849 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -29,6 +29,8 @@ typedef enum pragma_kind {
PRAGMA_OACC_CACHE,
PRAGMA_OACC_DATA,
+ PRAGMA_OACC_ENTER_DATA,
+ PRAGMA_OACC_EXIT_DATA,
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index cb2fc63..3df8d28 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -1242,6 +1242,7 @@ static vec<tree, va_gc> *c_parser_expr_list (c_parser *, bool, bool,
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_oacc_enter_exit_data (c_parser *, bool);
static void c_parser_oacc_update (c_parser *);
static void c_parser_omp_barrier (c_parser *);
static void c_parser_omp_flush (c_parser *);
@@ -9544,6 +9545,14 @@ c_parser_pragma (c_parser *parser, enum pragma_context context)
switch (id)
{
+ case PRAGMA_OACC_ENTER_DATA:
+ c_parser_oacc_enter_exit_data (parser, true);
+ return false;
+
+ case PRAGMA_OACC_EXIT_DATA:
+ c_parser_oacc_enter_exit_data (parser, false);
+ return false;
+
case PRAGMA_OACC_UPDATE:
if (context != pragma_compound)
{
@@ -11937,6 +11946,87 @@ c_parser_oacc_data (location_t loc, c_parser *parser)
}
/* OpenACC 2.0:
+ # pragma acc enter data oacc-enter-data-clause[optseq] new-line
+
+ or
+
+ # pragma acc exit data oacc-exit-data-clause[optseq] new-line
+
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_ENTER_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
+
+#define OACC_EXIT_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
+
+static void
+c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
+{
+ location_t loc = c_parser_peek_token (parser)->location;
+ tree clauses, stmt;
+
+ c_parser_consume_pragma (parser);
+
+ if (!c_parser_next_token_is (parser, CPP_NAME))
+ {
+ c_parser_error (parser, enter
+ ? "expected %<data%> in %<#pragma acc enter data%>"
+ : "expected %<data%> in %<#pragma acc exit data%>");
+ c_parser_skip_to_pragma_eol (parser);
+ return;
+ }
+
+ const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ if (strcmp (p, "data") != 0)
+ {
+ c_parser_error (parser, "invalid pragma");
+ c_parser_skip_to_pragma_eol (parser);
+ return;
+ }
+
+ c_parser_consume_token (parser);
+
+ if (enter)
+ clauses = c_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK,
+ "#pragma acc enter data");
+ else
+ clauses = c_parser_oacc_all_clauses (parser, OACC_EXIT_DATA_CLAUSE_MASK,
+ "#pragma acc exit data");
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (loc, enter
+ ? "%<#pragma acc enter data%> has no data movement clause"
+ : "%<#pragma acc exit data%> has no data movement clause");
+ return;
+ }
+
+ stmt = enter ? make_node (OACC_ENTER_DATA) : make_node (OACC_EXIT_DATA);;
+ TREE_TYPE (stmt) = void_type_node;
+ if (enter)
+ OACC_ENTER_DATA_CLAUSES (stmt) = clauses;
+ else
+ OACC_EXIT_DATA_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+ add_stmt (stmt);
+}
+
+
+/* OpenACC 2.0:
+
# pragma acc loop oacc-loop-clause[optseq] new-line
structured-block
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 9a9ace1..3987081 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -31456,6 +31456,84 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
}
/* OpenACC 2.0:
+ # pragma acc enter data oacc-enter-data-clause[optseq] new-line
+
+ or
+
+ # pragma acc exit data oacc-exit-data-clause[optseq] new-line
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_ENTER_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
+
+#define OACC_EXIT_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_WAIT) )
+
+static tree
+cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok,
+ bool enter)
+{
+ tree stmt, clauses;
+
+ if (cp_lexer_next_token_is (parser->lexer, CPP_PRAGMA_EOL)
+ || cp_lexer_next_token_is_not (parser->lexer, CPP_NAME))
+ {
+ cp_parser_error (parser, enter
+ ? "expected %<data%> in %<#pragma acc enter data%>"
+ : "expected %<data%> in %<#pragma acc exit data%>");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return NULL_TREE;
+ }
+
+ const char *p =
+ IDENTIFIER_POINTER (cp_lexer_peek_token (parser->lexer)->u.value);
+ if (strcmp (p, "data") != 0)
+ {
+ cp_parser_error (parser, "invalid pragma");
+ cp_parser_skip_to_pragma_eol (parser, pragma_tok);
+ return NULL_TREE;
+ }
+
+ cp_lexer_consume_token (parser->lexer);
+
+ if (enter)
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_ENTER_DATA_CLAUSE_MASK,
+ "#pragma acc enter data", pragma_tok);
+ else
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_EXIT_DATA_CLAUSE_MASK,
+ "#pragma acc exit data", pragma_tok);
+
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (pragma_tok->location,
+ "%<#pragma acc enter data%> has no data movement clause");
+ return NULL_TREE;
+ }
+
+ stmt = enter ? make_node (OACC_ENTER_DATA) : make_node (OACC_EXIT_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ if (enter)
+ OACC_ENTER_DATA_CLAUSES (stmt) = clauses;
+ else
+ OACC_EXIT_DATA_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, pragma_tok->location);
+ add_stmt (stmt);
+ return stmt;
+}
+
+/* OpenACC 2.0:
# pragma acc kernels oacc-kernels-clause[optseq] new-line
structured-block */
@@ -32298,6 +32376,12 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok)
case PRAGMA_OACC_DATA:
stmt = cp_parser_oacc_data (parser, pragma_tok);
break;
+ case PRAGMA_OACC_ENTER_DATA:
+ stmt = cp_parser_oacc_enter_exit_data (parser, pragma_tok, true);
+ break;
+ case PRAGMA_OACC_EXIT_DATA:
+ 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;
@@ -32857,6 +32941,8 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context)
case PRAGMA_OACC_CACHE:
case PRAGMA_OACC_DATA:
+ case PRAGMA_OACC_ENTER_DATA:
+ case PRAGMA_OACC_EXIT_DATA:
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_LOOP:
diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c
index f85f6b6..57af730 100644
--- a/gcc/fortran/dump-parse-tree.c
+++ b/gcc/fortran/dump-parse-tree.c
@@ -1255,8 +1255,6 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
case OMP_LIST_DEVICEPTR: type = "DEVICEPTR"; break;
case OMP_LIST_USE_DEVICE: type = "USE_DEVICE"; break;
case OMP_LIST_DEVICE_RESIDENT: type = "USE_DEVICE"; break;
- case OMP_LIST_HOST: type = "HOST"; break;
- case OMP_LIST_DEVICE: type = "DEVICE"; break;
case OMP_LIST_CACHE: type = ""; break;
case OMP_LIST_PRIVATE: type = "PRIVATE"; break;
case OMP_LIST_FIRSTPRIVATE: type = "FIRSTPRIVATE"; break;
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e685b67..6bd131c 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1188,8 +1188,6 @@ enum
OMP_LIST_DATA_CLAUSE_LAST = OMP_LIST_DEVICEPTR,
OMP_LIST_DEVICE_RESIDENT,
OMP_LIST_USE_DEVICE,
- OMP_LIST_HOST,
- OMP_LIST_DEVICE,
OMP_LIST_CACHE,
OMP_LIST_NUM,
OMP_LIST_LAST = OMP_LIST_NUM
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 1970730..c7af004 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -451,6 +451,7 @@ match_oacc_clause_gang (gfc_omp_clauses *cp)
#define OMP_CLAUSE_DELETE (1ULL << 55)
#define OMP_CLAUSE_AUTO (1ULL << 56)
#define OMP_CLAUSE_TILE (1ULL << 57)
+#define OMP_CLAUSE_SELF (1ULL << 58)
/* Helper function for OpenACC and OpenMP clauses involving memory
mapping. */
@@ -682,18 +683,23 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask,
== MATCH_YES)
continue;
if ((mask & OMP_CLAUSE_HOST)
- && gfc_match_omp_variable_list ("host (",
- &c->lists[OMP_LIST_HOST], true)
- == MATCH_YES)
+ && gfc_match ("host ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_FORCE_FROM))
continue;
if ((mask & OMP_CLAUSE_OACC_DEVICE)
- && gfc_match_omp_variable_list ("device (",
- &c->lists[OMP_LIST_DEVICE], true)
- == MATCH_YES)
+ && gfc_match ("device ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_FORCE_TO))
continue;
if ((mask & OMP_CLAUSE_TILE)
&& match_oacc_expr_list ("tile (", &c->tile_list, true) == MATCH_YES)
continue;
+ if ((mask & OMP_CLAUSE_SELF)
+ && gfc_match ("self ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_FORCE_FROM))
+ continue;
if ((mask & OMP_CLAUSE_SEQ) && !c->seq
&& gfc_match ("seq") == MATCH_YES)
{
@@ -1164,7 +1170,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask,
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
| OMP_CLAUSE_PRESENT_OR_CREATE)
#define OACC_UPDATE_CLAUSES \
- (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST | OMP_CLAUSE_OACC_DEVICE)
+ (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST | OMP_CLAUSE_SELF \
+ | OMP_CLAUSE_OACC_DEVICE | OMP_CLAUSE_WAIT)
#define OACC_ENTER_DATA_CLAUSES \
(OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT | OMP_CLAUSE_COPYIN \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT_OR_COPYIN \
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 4d7f3ea..2de7127 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1806,12 +1806,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_LIST_DEVICE_RESIDENT:
clause_code = OMP_CLAUSE_DEVICE_RESIDENT;
goto add_clause;
- case OMP_LIST_HOST:
- clause_code = OMP_CLAUSE_HOST;
- goto add_clause;
- case OMP_LIST_DEVICE:
- clause_code = OMP_CLAUSE_OACC_DEVICE;
- goto add_clause;
case OMP_LIST_CACHE:
clause_code = OMP_NO_CLAUSE_CACHE;
goto add_clause;
@@ -2558,17 +2552,14 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (clauses->wait_list)
{
gfc_expr_list *el;
- tree list = NULL;
for (el = clauses->wait_list; el; el = el->next)
{
c = build_omp_clause (where.lb->location, OMP_CLAUSE_WAIT);
OMP_CLAUSE_DECL (c) = gfc_convert_expr_to_tree (block, el->expr);
- OMP_CLAUSE_CHAIN (c) = list;
- list = c;
+ OMP_CLAUSE_CHAIN (c) = omp_clauses;
+ omp_clauses = c;
}
-
- omp_clauses = list;
}
if (clauses->num_gangs_expr)
{
@@ -2726,9 +2717,6 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
case EXEC_OACC_EXIT_DATA:
construct_code = OACC_EXIT_DATA;
break;
- case EXEC_OACC_WAIT:
- construct_code = OACC_WAIT;
- break;
case EXEC_OACC_CACHE:
construct_code = OACC_CACHE;
break;
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index 861529e..c8f978d 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1335,6 +1335,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
case GF_OMP_TARGET_KIND_OACC_DATA:
kind = " oacc_data";
break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ kind = " oacc_enter_exit_data";
+ break;
case GF_OMP_TARGET_KIND_OACC_UPDATE:
kind = " oacc_update";
break;
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 8eb3993..7bc673a 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -109,6 +109,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_UPDATE = 2,
GF_OMP_TARGET_KIND_OACC_DATA = 3,
GF_OMP_TARGET_KIND_OACC_UPDATE = 4,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 5,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 0ba1b23..9a5d85c 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -7364,6 +7364,14 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
+ case OACC_ENTER_DATA:
+ clauses = OACC_ENTER_DATA_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ break;
+ case OACC_EXIT_DATA:
+ clauses = OACC_EXIT_DATA_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ break;
case OACC_UPDATE:
clauses = OACC_UPDATE_CLAUSES (expr);
kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
@@ -8305,8 +8313,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
case OACC_HOST_DATA:
case OACC_DECLARE:
- case OACC_ENTER_DATA:
- case OACC_EXIT_DATA:
case OACC_CACHE:
sorry ("directive not yet implemented");
ret = GS_ALL_DONE;
@@ -8359,6 +8365,8 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
case OACC_UPDATE:
case OMP_TARGET_UPDATE:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
gimplify_omp_target_update (expr_p, pre_p);
ret = GS_ALL_DONE;
break;
diff --git a/gcc/oacc-builtins.def b/gcc/oacc-builtins.def
index ec60612..0ac97f2 100644
--- a/gcc/oacc-builtins.def
+++ b/gcc/oacc-builtins.def
@@ -31,6 +31,9 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
BT_FN_VOID, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_KERNELS, "GOACC_kernels",
BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT_INT_INT_VAR,
ATTR_NOTHROW_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ba94f80..b219008 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -9436,6 +9436,9 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_DATA:
start_ix = BUILT_IN_GOACC_DATA_START;
break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
+ break;
case GF_OMP_TARGET_KIND_OACC_UPDATE:
start_ix = BUILT_IN_GOACC_UPDATE;
break;
@@ -9570,6 +9573,7 @@ expand_omp_target (struct omp_region *region)
args->quick_push (t4);
if (kind == GF_OMP_TARGET_KIND_OACC_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA
|| kind == GF_OMP_TARGET_KIND_OACC_UPDATE)
{
int idx;
@@ -9582,9 +9586,9 @@ expand_omp_target (struct omp_region *region)
t1 = fold_convert_loc (gimple_location (entry_stmt),
integer_type_node, build_int_cst (integer_type_node, -2));
- args->quick_push (t1);
+ args->safe_push (t1);
idx = args->length ();
- args->quick_push (fold_convert_loc (gimple_location (entry_stmt),
+ args->safe_push (fold_convert_loc (gimple_location (entry_stmt),
integer_type_node, integer_minus_one_node));
c = find_omp_clause (clauses, OMP_CLAUSE_WAIT);
@@ -9596,7 +9600,7 @@ expand_omp_target (struct omp_region *region)
{
if (OMP_CLAUSE_CODE (t) == OMP_CLAUSE_WAIT)
{
- args->quick_push (fold_convert (integer_type_node,
+ args->safe_push (fold_convert (integer_type_node,
OMP_CLAUSE_WAIT_EXPR (t)));
n++;
}
@@ -11864,6 +11868,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_MAP_FORCE_DEALLOC:
case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA
|| kind == GF_OMP_TARGET_KIND_OACC_UPDATE);
break;
default:
@@ -11963,6 +11968,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
talign_shift = 3;
break;
case GF_OMP_TARGET_KIND_OACC_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
tkind_type = short_unsigned_type_node;
talign_shift = 8;
@@ -12169,7 +12175,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|| kind == GF_OMP_TARGET_KIND_OACC_DATA)
new_body = tgt_body;
if (kind != GF_OMP_TARGET_KIND_UPDATE
- && kind != GF_OMP_TARGET_KIND_OACC_UPDATE)
+ && kind != GF_OMP_TARGET_KIND_OACC_UPDATE
+ && kind != GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)
{
gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
gimple_omp_set_body (stmt, new_body);
diff --git a/gcc/testsuite/c-c++-common/goacc/data-2.c b/gcc/testsuite/c-c++-common/goacc/data-2.c
new file mode 100644
index 0000000..9c0a185
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/data-2.c
@@ -0,0 +1,21 @@
+void
+foo (void)
+{
+ int a, b[100];
+ int n;
+#pragma acc enter data copyin (a, b) async wait
+#pragma acc enter data create (b[20:30]) async wait
+#pragma acc enter data (a) /* { dg-error "expected clause before '\\\(' token" } */
+#pragma acc enter data create (b(1:10)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc exit data delete (a) if (0)
+#pragma acc exit data copyout (b) if (a)
+#pragma acc exit data delete (b)
+#pragma acc enter /* { dg-error "expected 'data' in" } */
+#pragma acc exit /* { dg-error "expected 'data' in" } */
+#pragma acc enter data /* { dg-error "has no data movement clause" } */
+#pragma acc exit data /* { dg-error "has no data movement clause" } */
+#pragma acc enter Data /* { dg-error "invalid pragma before" } */
+#pragma acc exit copyout (b) /* { dg-error "invalid pragma before" } */
+}
+
+/* { dg-error "has no data movement clause" "" { target *-*-* } 8 } */
diff --git a/gcc/testsuite/c-c++-common/goacc/update-1.c b/gcc/testsuite/c-c++-common/goacc/update-1.c
index 970fdca..2a3a910 100644
--- a/gcc/testsuite/c-c++-common/goacc/update-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/update-1.c
@@ -4,7 +4,9 @@ f (void)
#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host/self' clause" } */
int i = 0;
+ int a[10];
#pragma acc update device(i)
#pragma acc update host(i)
#pragma acc update self(i)
+#pragma acc update host(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
}
diff --git a/libgfortran/Makefile.in b/libgfortran/Makefile.in
index 2eac2e8..7a231bb 100644
--- a/libgfortran/Makefile.in
+++ b/libgfortran/Makefile.in
@@ -1,9 +1,9 @@
-# Makefile.in generated by automake 1.11.6 from Makefile.am.
+# Makefile.in generated by automake 1.11.1 from Makefile.am.
# @configure_input@
# Copyright (C) 1994, 1995, 1996, 1997, 1998, 1999, 2000, 2001, 2002,
-# 2003, 2004, 2005, 2006, 2007, 2008, 2009, 2010, 2011 Free Software
-# Foundation, Inc.
+# 2003, 2004, 2005, 2006, 2007, 2008, 2009 Free Software Foundation,
+# Inc.
# This Makefile.in is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
@@ -18,23 +18,6 @@
VPATH = @srcdir@
-am__make_dryrun = \
- { \
- am__dry=no; \
- case $$MAKEFLAGS in \
- *\\[\ \ ]*) \
- echo 'am--echo: ; @echo "AM" OK' | $(MAKE) -f - 2>/dev/null \
- | grep '^AM OK$$' >/dev/null || am__dry=yes;; \
- *) \
- for am__flg in $$MAKEFLAGS; do \
- case $$am__flg in \
- *=*|--*) ;; \
- *n*) am__dry=yes; break;; \
- esac; \
- done;; \
- esac; \
- test $$am__dry = yes; \
- }
pkgdatadir = $(datadir)/@PACKAGE@
pkgincludedir = $(includedir)/@PACKAGE@
pkglibdir = $(libdir)/@PACKAGE@
@@ -106,12 +89,6 @@ am__nobase_list = $(am__nobase_strip_setup); \
am__base_list = \
sed '$$!N;$$!N;$$!N;$$!N;$$!N;$$!N;$$!N;s/\n/ /g' | \
sed '$$!N;$$!N;$$!N;$$!N;s/\n/ /g'
-am__uninstall_files_from_dir = { \
- test -z "$$files" \
- || { test ! -d "$$dir" && test ! -f "$$dir" && test ! -r "$$dir"; } \
- || { echo " ( cd '$$dir' && rm -f" $$files ")"; \
- $(am__cd) "$$dir" && rm -f $$files; }; \
- }
am__installdirs = "$(DESTDIR)$(cafexeclibdir)" \
"$(DESTDIR)$(myexeclibdir)" "$(DESTDIR)$(toolexeclibdir)" \
"$(DESTDIR)$(toolexeclibdir)" "$(DESTDIR)$(fincludedir)"
@@ -358,11 +335,6 @@ MULTIDIRS =
MULTISUBDIR =
MULTIDO = true
MULTICLEAN = true
-am__can_run_installinfo = \
- case $$AM_UPDATE_INFO_DIR in \
- n|no|NO) false;; \
- *) (install-info --version) >/dev/null 2>&1;; \
- esac
DATA = $(toolexeclib_DATA)
HEADERS = $(nodist_finclude_HEADERS)
ETAGS = etags
@@ -1283,7 +1255,7 @@ all: $(BUILT_SOURCES) config.h
.SUFFIXES:
.SUFFIXES: .F90 .c .f90 .lo .o .obj
-am--refresh: Makefile
+am--refresh:
@:
$(srcdir)/Makefile.in: @MAINTAINER_MODE_TRUE@ $(srcdir)/Makefile.am $(am__configure_deps)
@for dep in $?; do \
@@ -1319,8 +1291,10 @@ $(ACLOCAL_M4): @MAINTAINER_MODE_TRUE@ $(am__aclocal_m4_deps)
$(am__aclocal_m4_deps):
config.h: stamp-h1
- @if test ! -f $@; then rm -f stamp-h1; else :; fi
- @if test ! -f $@; then $(MAKE) $(AM_MAKEFLAGS) stamp-h1; else :; fi
+ @if test ! -f $@; then \
+ rm -f stamp-h1; \
+ $(MAKE) $(AM_MAKEFLAGS) stamp-h1; \
+ else :; fi
stamp-h1: $(srcdir)/config.h.in $(top_builddir)/config.status
@rm -f stamp-h1
@@ -1336,6 +1310,7 @@ libgfortran.spec: $(top_builddir)/config.status $(srcdir)/libgfortran.spec.in
cd $(top_builddir) && $(SHELL) ./config.status $@
install-cafexeclibLTLIBRARIES: $(cafexeclib_LTLIBRARIES)
@$(NORMAL_INSTALL)
+ test -z "$(cafexeclibdir)" || $(MKDIR_P) "$(DESTDIR)$(cafexeclibdir)"
@list='$(cafexeclib_LTLIBRARIES)'; test -n "$(cafexeclibdir)" || list=; \
list2=; for p in $$list; do \
if test -f $$p; then \
@@ -1343,8 +1318,6 @@ install-cafexeclibLTLIBRARIES: $(cafexeclib_LTLIBRARIES)
else :; fi; \
done; \
test -z "$$list2" || { \
- echo " $(MKDIR_P) '$(DESTDIR)$(cafexeclibdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(cafexeclibdir)" || exit 1; \
echo " $(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 '$(DESTDIR)$(cafexeclibdir)'"; \
$(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 "$(DESTDIR)$(cafexeclibdir)"; \
}
@@ -1368,6 +1341,7 @@ clean-cafexeclibLTLIBRARIES:
done
install-myexeclibLTLIBRARIES: $(myexeclib_LTLIBRARIES)
@$(NORMAL_INSTALL)
+ test -z "$(myexeclibdir)" || $(MKDIR_P) "$(DESTDIR)$(myexeclibdir)"
@list='$(myexeclib_LTLIBRARIES)'; test -n "$(myexeclibdir)" || list=; \
list2=; for p in $$list; do \
if test -f $$p; then \
@@ -1375,8 +1349,6 @@ install-myexeclibLTLIBRARIES: $(myexeclib_LTLIBRARIES)
else :; fi; \
done; \
test -z "$$list2" || { \
- echo " $(MKDIR_P) '$(DESTDIR)$(myexeclibdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(myexeclibdir)" || exit 1; \
echo " $(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 '$(DESTDIR)$(myexeclibdir)'"; \
$(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 "$(DESTDIR)$(myexeclibdir)"; \
}
@@ -1400,6 +1372,7 @@ clean-myexeclibLTLIBRARIES:
done
install-toolexeclibLTLIBRARIES: $(toolexeclib_LTLIBRARIES)
@$(NORMAL_INSTALL)
+ test -z "$(toolexeclibdir)" || $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)"
@list='$(toolexeclib_LTLIBRARIES)'; test -n "$(toolexeclibdir)" || list=; \
list2=; for p in $$list; do \
if test -f $$p; then \
@@ -1407,8 +1380,6 @@ install-toolexeclibLTLIBRARIES: $(toolexeclib_LTLIBRARIES)
else :; fi; \
done; \
test -z "$$list2" || { \
- echo " $(MKDIR_P) '$(DESTDIR)$(toolexeclibdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)" || exit 1; \
echo " $(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 '$(DESTDIR)$(toolexeclibdir)'"; \
$(LIBTOOL) $(AM_LIBTOOLFLAGS) $(LIBTOOLFLAGS) --mode=install $(INSTALL) $(INSTALL_STRIP_FLAG) $$list2 "$(DESTDIR)$(toolexeclibdir)"; \
}
@@ -1430,11 +1401,11 @@ clean-toolexeclibLTLIBRARIES:
echo "rm -f \"$${dir}/so_locations\""; \
rm -f "$${dir}/so_locations"; \
done
-libcaf_single.la: $(libcaf_single_la_OBJECTS) $(libcaf_single_la_DEPENDENCIES) $(EXTRA_libcaf_single_la_DEPENDENCIES)
+libcaf_single.la: $(libcaf_single_la_OBJECTS) $(libcaf_single_la_DEPENDENCIES)
$(libcaf_single_la_LINK) -rpath $(cafexeclibdir) $(libcaf_single_la_OBJECTS) $(libcaf_single_la_LIBADD) $(LIBS)
-libgfortran.la: $(libgfortran_la_OBJECTS) $(libgfortran_la_DEPENDENCIES) $(EXTRA_libgfortran_la_DEPENDENCIES)
+libgfortran.la: $(libgfortran_la_OBJECTS) $(libgfortran_la_DEPENDENCIES)
$(libgfortran_la_LINK) -rpath $(toolexeclibdir) $(libgfortran_la_OBJECTS) $(libgfortran_la_LIBADD) $(LIBS)
-libgfortranbegin.la: $(libgfortranbegin_la_OBJECTS) $(libgfortranbegin_la_DEPENDENCIES) $(EXTRA_libgfortranbegin_la_DEPENDENCIES)
+libgfortranbegin.la: $(libgfortranbegin_la_OBJECTS) $(libgfortranbegin_la_DEPENDENCIES)
$(libgfortranbegin_la_LINK) -rpath $(myexeclibdir) $(libgfortranbegin_la_OBJECTS) $(libgfortranbegin_la_LIBADD) $(LIBS)
mostlyclean-compile:
@@ -5693,11 +5664,8 @@ maintainer-clean-multi:
$(MULTICLEAN) $(AM_MAKEFLAGS) DO=maintainer-clean multi-clean # $(MAKE)
install-toolexeclibDATA: $(toolexeclib_DATA)
@$(NORMAL_INSTALL)
+ test -z "$(toolexeclibdir)" || $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)"
@list='$(toolexeclib_DATA)'; test -n "$(toolexeclibdir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(toolexeclibdir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(toolexeclibdir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
@@ -5711,14 +5679,13 @@ uninstall-toolexeclibDATA:
@$(NORMAL_UNINSTALL)
@list='$(toolexeclib_DATA)'; test -n "$(toolexeclibdir)" || list=; \
files=`for p in $$list; do echo $$p; done | sed -e 's|^.*/||'`; \
- dir='$(DESTDIR)$(toolexeclibdir)'; $(am__uninstall_files_from_dir)
+ test -n "$$files" || exit 0; \
+ echo " ( cd '$(DESTDIR)$(toolexeclibdir)' && rm -f" $$files ")"; \
+ cd "$(DESTDIR)$(toolexeclibdir)" && rm -f $$files
install-nodist_fincludeHEADERS: $(nodist_finclude_HEADERS)
@$(NORMAL_INSTALL)
+ test -z "$(fincludedir)" || $(MKDIR_P) "$(DESTDIR)$(fincludedir)"
@list='$(nodist_finclude_HEADERS)'; test -n "$(fincludedir)" || list=; \
- if test -n "$$list"; then \
- echo " $(MKDIR_P) '$(DESTDIR)$(fincludedir)'"; \
- $(MKDIR_P) "$(DESTDIR)$(fincludedir)" || exit 1; \
- fi; \
for p in $$list; do \
if test -f "$$p"; then d=; else d="$(srcdir)/"; fi; \
echo "$$d$$p"; \
@@ -5732,7 +5699,9 @@ uninstall-nodist_fincludeHEADERS:
@$(NORMAL_UNINSTALL)
@list='$(nodist_finclude_HEADERS)'; test -n "$(fincludedir)" || list=; \
files=`for p in $$list; do echo $$p; done | sed -e 's|^.*/||'`; \
- dir='$(DESTDIR)$(fincludedir)'; $(am__uninstall_files_from_dir)
+ test -n "$$files" || exit 0; \
+ echo " ( cd '$(DESTDIR)$(fincludedir)' && rm -f" $$files ")"; \
+ cd "$(DESTDIR)$(fincludedir)" && rm -f $$files
ID: $(HEADERS) $(SOURCES) $(LISP) $(TAGS_FILES)
list='$(SOURCES) $(HEADERS) $(LISP) $(TAGS_FILES)'; \
@@ -5804,15 +5773,10 @@ install-am: all-am
installcheck: installcheck-am
install-strip:
- if test -z '$(STRIP)'; then \
- $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
- install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
- install; \
- else \
- $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
- install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
- "INSTALL_PROGRAM_ENV=STRIPPROG='$(STRIP)'" install; \
- fi
+ $(MAKE) $(AM_MAKEFLAGS) INSTALL_PROGRAM="$(INSTALL_STRIP_PROGRAM)" \
+ install_sh_PROGRAM="$(INSTALL_STRIP_PROGRAM)" INSTALL_STRIP_FLAG=-s \
+ `test -z '$(STRIP)' || \
+ echo "INSTALL_PROGRAM_ENV=STRIPPROG='$(STRIP)'"` install
mostlyclean-generic:
clean-generic:
diff --git a/libgfortran/aclocal.m4 b/libgfortran/aclocal.m4
index 0ec2c8f..8673daa 100644
--- a/libgfortran/aclocal.m4
+++ b/libgfortran/aclocal.m4
@@ -1,8 +1,7 @@
-# generated automatically by aclocal 1.11.6 -*- Autoconf -*-
+# generated automatically by aclocal 1.11.1 -*- Autoconf -*-
# Copyright (C) 1996, 1997, 1998, 1999, 2000, 2001, 2002, 2003, 2004,
-# 2005, 2006, 2007, 2008, 2009, 2010, 2011 Free Software Foundation,
-# Inc.
+# 2005, 2006, 2007, 2008, 2009 Free Software Foundation, Inc.
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
@@ -20,15 +19,12 @@ You have another version of autoconf. It may work, but is not guaranteed to.
If you have problems, you may need to regenerate the build system entirely.
To do so, use the procedure documented by the package, typically `autoreconf'.])])
-# Copyright (C) 2002, 2003, 2005, 2006, 2007, 2008, 2011 Free Software
-# Foundation, Inc.
+# Copyright (C) 2002, 2003, 2005, 2006, 2007, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_AUTOMAKE_VERSION(VERSION)
# ----------------------------
# Automake X.Y traces this macro to ensure aclocal.m4 has been
@@ -38,7 +34,7 @@ AC_DEFUN([AM_AUTOMAKE_VERSION],
[am__api_version='1.11'
dnl Some users find AM_AUTOMAKE_VERSION and mistake it for a way to
dnl require some minimum version. Point them to the right macro.
-m4_if([$1], [1.11.6], [],
+m4_if([$1], [1.11.1], [],
[AC_FATAL([Do not call $0, use AM_INIT_AUTOMAKE([$1]).])])dnl
])
@@ -54,21 +50,19 @@ m4_define([_AM_AUTOCONF_VERSION], [])
# Call AM_AUTOMAKE_VERSION and AM_AUTOMAKE_VERSION so they can be traced.
# This function is AC_REQUIREd by AM_INIT_AUTOMAKE.
AC_DEFUN([AM_SET_CURRENT_AUTOMAKE_VERSION],
-[AM_AUTOMAKE_VERSION([1.11.6])dnl
+[AM_AUTOMAKE_VERSION([1.11.1])dnl
m4_ifndef([AC_AUTOCONF_VERSION],
[m4_copy([m4_PACKAGE_VERSION], [AC_AUTOCONF_VERSION])])dnl
_AM_AUTOCONF_VERSION(m4_defn([AC_AUTOCONF_VERSION]))])
# AM_AUX_DIR_EXPAND -*- Autoconf -*-
-# Copyright (C) 2001, 2003, 2005, 2011 Free Software Foundation, Inc.
+# Copyright (C) 2001, 2003, 2005 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# For projects using AC_CONFIG_AUX_DIR([foo]), Autoconf sets
# $ac_aux_dir to `$srcdir/foo'. In other projects, it is set to
# `$srcdir', `$srcdir/..', or `$srcdir/../..'.
@@ -150,14 +144,14 @@ AC_CONFIG_COMMANDS_PRE(
Usually this means the macro was only invoked conditionally.]])
fi])])
-# Copyright (C) 1999, 2000, 2001, 2002, 2003, 2004, 2005, 2006, 2009,
-# 2010, 2011 Free Software Foundation, Inc.
+# Copyright (C) 1999, 2000, 2001, 2002, 2003, 2004, 2005, 2006, 2009
+# Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 12
+# serial 10
# There are a few dirty hacks below to avoid letting `AC_PROG_CC' be
# written in clear, in which case automake, when reading aclocal.m4,
@@ -197,7 +191,6 @@ AC_CACHE_CHECK([dependency style of $depcc],
# instance it was reported that on HP-UX the gcc test will end up
# making a dummy file named `D' -- because `-MD' means `put the output
# in D'.
- rm -rf conftest.dir
mkdir conftest.dir
# Copy depcomp to subdir because otherwise we won't find it if we're
# using a relative directory.
@@ -262,7 +255,7 @@ AC_CACHE_CHECK([dependency style of $depcc],
break
fi
;;
- msvc7 | msvc7msys | msvisualcpp | msvcmsys)
+ msvisualcpp | msvcmsys)
# This compiler won't grok `-c -o', but also, the minuso test has
# not run yet. These depmodes are late enough in the game, and
# so weak that their functioning should not be impacted.
@@ -327,13 +320,10 @@ AC_DEFUN([AM_DEP_TRACK],
if test "x$enable_dependency_tracking" != xno; then
am_depcomp="$ac_aux_dir/depcomp"
AMDEPBACKSLASH='\'
- am__nodep='_no'
fi
AM_CONDITIONAL([AMDEP], [test "x$enable_dependency_tracking" != xno])
AC_SUBST([AMDEPBACKSLASH])dnl
_AM_SUBST_NOTMAKE([AMDEPBACKSLASH])dnl
-AC_SUBST([am__nodep])dnl
-_AM_SUBST_NOTMAKE([am__nodep])dnl
])
# Generate code to set up dependency tracking. -*- Autoconf -*-
@@ -555,15 +545,12 @@ for _am_header in $config_headers :; do
done
echo "timestamp for $_am_arg" >`AS_DIRNAME(["$_am_arg"])`/stamp-h[]$_am_stamp_count])
-# Copyright (C) 2001, 2003, 2005, 2008, 2011 Free Software Foundation,
-# Inc.
+# Copyright (C) 2001, 2003, 2005, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_PROG_INSTALL_SH
# ------------------
# Define $install_sh.
@@ -582,8 +569,8 @@ AC_SUBST(install_sh)])
# Add --enable-maintainer-mode option to configure. -*- Autoconf -*-
# From Jim Meyering
-# Copyright (C) 1996, 1998, 2000, 2001, 2002, 2003, 2004, 2005, 2008,
-# 2011 Free Software Foundation, Inc.
+# Copyright (C) 1996, 1998, 2000, 2001, 2002, 2003, 2004, 2005, 2008
+# Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
@@ -603,7 +590,7 @@ AC_DEFUN([AM_MAINTAINER_MODE],
[disable], [m4_define([am_maintainer_other], [enable])],
[m4_define([am_maintainer_other], [enable])
m4_warn([syntax], [unexpected argument to AM@&t@_MAINTAINER_MODE: $1])])
-AC_MSG_CHECKING([whether to enable maintainer-specific portions of Makefiles])
+AC_MSG_CHECKING([whether to am_maintainer_other maintainer-specific portions of Makefiles])
dnl maintainer-mode's default is 'disable' unless 'enable' is passed
AC_ARG_ENABLE([maintainer-mode],
[ --][am_maintainer_other][-maintainer-mode am_maintainer_other make rules and dependencies not useful
@@ -749,15 +736,12 @@ else
fi
])
-# Copyright (C) 2003, 2004, 2005, 2006, 2011 Free Software Foundation,
-# Inc.
+# Copyright (C) 2003, 2004, 2005, 2006 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_PROG_MKDIR_P
# ---------------
# Check for `mkdir -p'.
@@ -780,14 +764,13 @@ esac
# Helper functions for option handling. -*- Autoconf -*-
-# Copyright (C) 2001, 2002, 2003, 2005, 2008, 2010 Free Software
-# Foundation, Inc.
+# Copyright (C) 2001, 2002, 2003, 2005, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 5
+# serial 4
# _AM_MANGLE_OPTION(NAME)
# -----------------------
@@ -795,13 +778,13 @@ AC_DEFUN([_AM_MANGLE_OPTION],
[[_AM_OPTION_]m4_bpatsubst($1, [[^a-zA-Z0-9_]], [_])])
# _AM_SET_OPTION(NAME)
-# --------------------
+# ------------------------------
# Set option NAME. Presently that only means defining a flag for this option.
AC_DEFUN([_AM_SET_OPTION],
[m4_define(_AM_MANGLE_OPTION([$1]), 1)])
# _AM_SET_OPTIONS(OPTIONS)
-# ------------------------
+# ----------------------------------
# OPTIONS is a space-separated list of Automake options.
AC_DEFUN([_AM_SET_OPTIONS],
[m4_foreach_w([_AM_Option], [$1], [_AM_SET_OPTION(_AM_Option)])])
@@ -877,14 +860,12 @@ Check your system clock])
fi
AC_MSG_RESULT(yes)])
-# Copyright (C) 2001, 2003, 2005, 2011 Free Software Foundation, Inc.
+# Copyright (C) 2001, 2003, 2005 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 1
-
# AM_PROG_INSTALL_STRIP
# ---------------------
# One issue with vendor `install' (even GNU) is that you can't
@@ -907,13 +888,13 @@ fi
INSTALL_STRIP_PROGRAM="\$(install_sh) -c -s"
AC_SUBST([INSTALL_STRIP_PROGRAM])])
-# Copyright (C) 2006, 2008, 2010 Free Software Foundation, Inc.
+# Copyright (C) 2006, 2008 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
# with or without modifications, as long as this notice is preserved.
-# serial 3
+# serial 2
# _AM_SUBST_NOTMAKE(VARIABLE)
# ---------------------------
@@ -922,13 +903,13 @@ AC_SUBST([INSTALL_STRIP_PROGRAM])])
AC_DEFUN([_AM_SUBST_NOTMAKE])
# AM_SUBST_NOTMAKE(VARIABLE)
-# --------------------------
+# ---------------------------
# Public sister of _AM_SUBST_NOTMAKE.
AC_DEFUN([AM_SUBST_NOTMAKE], [_AM_SUBST_NOTMAKE($@)])
# Check how to create a tarball. -*- Autoconf -*-
-# Copyright (C) 2004, 2005, 2012 Free Software Foundation, Inc.
+# Copyright (C) 2004, 2005 Free Software Foundation, Inc.
#
# This file is free software; the Free Software Foundation
# gives unlimited permission to copy and/or distribute it,
@@ -950,11 +931,10 @@ AC_DEFUN([AM_SUBST_NOTMAKE], [_AM_SUBST_NOTMAKE($@)])
# a tarball read from stdin.
# $(am__untar) < result.tar
AC_DEFUN([_AM_PROG_TAR],
-[# Always define AMTAR for backward compatibility. Yes, it's still used
-# in the wild :-( We should find a proper way to deprecate it ...
-AC_SUBST([AMTAR], ['$${TAR-tar}'])
+[# Always define AMTAR for backward compatibility.
+AM_MISSING_PROG([AMTAR], [tar])
m4_if([$1], [v7],
- [am__tar='$${TAR-tar} chof - "$$tardir"' am__untar='$${TAR-tar} xf -'],
+ [am__tar='${AMTAR} chof - "$$tardir"'; am__untar='${AMTAR} xf -'],
[m4_case([$1], [ustar],, [pax],,
[m4_fatal([Unknown tar format])])
AC_MSG_CHECKING([how to create a $1 tar archive])
diff --git a/libgfortran/configure b/libgfortran/configure
index df17829..e22a8be 100755
--- a/libgfortran/configure
+++ b/libgfortran/configure
@@ -657,7 +657,6 @@ CPP
am__fastdepCC_FALSE
am__fastdepCC_TRUE
CCDEPMODE
-am__nodep
AMDEPBACKSLASH
AMDEP_FALSE
AMDEP_TRUE
@@ -3395,11 +3394,11 @@ MAKEINFO=${MAKEINFO-"${am_missing_run}makeinfo"}
# We need awk for the "check" target. The system "awk" is bad on
# some platforms.
-# Always define AMTAR for backward compatibility. Yes, it's still used
-# in the wild :-( We should find a proper way to deprecate it ...
-AMTAR='$${TAR-tar}'
+# Always define AMTAR for backward compatibility.
-am__tar='$${TAR-tar} chof - "$$tardir"' am__untar='$${TAR-tar} xf -'
+AMTAR=${AMTAR-"${am_missing_run}tar"}
+
+am__tar='${AMTAR} chof - "$$tardir"'; am__untar='${AMTAR} xf -'
@@ -3532,7 +3531,6 @@ fi
if test "x$enable_dependency_tracking" != xno; then
am_depcomp="$ac_aux_dir/depcomp"
AMDEPBACKSLASH='\'
- am__nodep='_no'
fi
if test "x$enable_dependency_tracking" != xno; then
AMDEP_TRUE=
@@ -4350,7 +4348,6 @@ else
# instance it was reported that on HP-UX the gcc test will end up
# making a dummy file named `D' -- because `-MD' means `put the output
# in D'.
- rm -rf conftest.dir
mkdir conftest.dir
# Copy depcomp to subdir because otherwise we won't find it if we're
# using a relative directory.
@@ -4410,7 +4407,7 @@ else
break
fi
;;
- msvc7 | msvc7msys | msvisualcpp | msvcmsys)
+ msvisualcpp | msvcmsys)
# This compiler won't grok `-c -o', but also, the minuso test has
# not run yet. These depmodes are late enough in the game, and
# so weak that their functioning should not be impacted.
@@ -5526,7 +5523,6 @@ else
# instance it was reported that on HP-UX the gcc test will end up
# making a dummy file named `D' -- because `-MD' means `put the output
# in D'.
- rm -rf conftest.dir
mkdir conftest.dir
# Copy depcomp to subdir because otherwise we won't find it if we're
# using a relative directory.
@@ -5586,7 +5582,7 @@ else
break
fi
;;
- msvc7 | msvc7msys | msvisualcpp | msvcmsys)
+ msvisualcpp | msvcmsys)
# This compiler won't grok `-c -o', but also, the minuso test has
# not run yet. These depmodes are late enough in the game, and
# so weak that their functioning should not be impacted.
@@ -12350,7 +12346,7 @@ else
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
-#line 12353 "configure"
+#line 12349 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
@@ -12456,7 +12452,7 @@ else
lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2
lt_status=$lt_dlunknown
cat > conftest.$ac_ext <<_LT_EOF
-#line 12459 "configure"
+#line 12455 "configure"
#include "confdefs.h"
#if HAVE_DLFCN_H
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index d879851..ccdad71 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -317,6 +317,7 @@ GOACC_2.0 {
global:
GOACC_data_end;
GOACC_data_start;
+ GOACC_enter_exit_data;
GOACC_kernels;
GOACC_parallel;
GOACC_update;
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index 35b0627..4455be1 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -219,6 +219,10 @@ extern void GOMP_teams (unsigned int, unsigned int);
extern void GOACC_data_start (int, const void *,
size_t, void **, size_t *, unsigned short *);
extern void GOACC_data_end (void);
+extern void GOACC_enter_exit_data (int device, const void *openmp_target,
+ size_t mapnum, void **hostaddrs,
+ size_t *sizes, unsigned short *kinds,
+ int async, int num_waits, ...);
extern void GOACC_kernels (int, void (*) (void *), const void *,
size_t, void **, size_t *, unsigned short *,
int, int, int, int, int, ...);
@@ -231,4 +235,11 @@ extern void GOACC_update (int device, const void *openmp_target, size_t mapnum,
int num_waits, ...);
extern void GOACC_wait (int, int, ...);
+/* oacc-mem.c */
+
+extern void gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs,
+ size_t *sizes, void *kinds);
+extern void gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async,
+ int mapnum);
+
#endif /* LIBGOMP_G_H */
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 582a1e0..0c45d19 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -332,7 +332,7 @@ acc_unmap_data (void *h)
gomp_mutex_unlock (&acc_dev->mem_map.lock);
}
-
+
gomp_unmap_vars (t, true);
}
@@ -393,7 +393,7 @@ present_create_copy (unsigned f, void *h, size_t s)
gomp_mutex_unlock (&acc_dev->mem_map.lock);
}
-
+
return d;
}
@@ -502,3 +502,80 @@ acc_update_self (void *h, size_t s)
{
update_dev_host (0, h, s);
}
+
+void
+gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
+ void *kinds)
+{
+ struct target_mem_desc *tgt;
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+
+ gomp_notify (" %s: prepare mappings\n", __FUNCTION__);
+ tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev, mapnum, hostaddrs,
+ NULL, sizes, kinds, true, false);
+ gomp_notify (" %s: mappings prepared\n", __FUNCTION__);
+ tgt->prev = acc_dev->openacc.data_environ;
+ acc_dev->openacc.data_environ = tgt;
+}
+
+void
+gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
+{
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+ splay_tree_key n;
+ struct target_mem_desc *t;
+ int minrefs = (mapnum == 1) ? 2 : 3;
+
+ n = lookup_host (&acc_dev->mem_map, h, 1);
+
+ if (!n)
+ gomp_fatal ("%p is not a mapped block", (void *)h);
+
+ gomp_notify (" %s: restore mappings\n", __FUNCTION__);
+
+ t = n->tgt;
+
+ struct target_mem_desc *tp;
+
+ gomp_mutex_lock (&acc_dev->mem_map.lock);
+
+ if (t->refcount == minrefs)
+ {
+ /* This is the last reference, so pull the descriptor off the
+ chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from
+ freeing the device memory. */
+ t->tgt_end = 0;
+ t->to_free = 0;
+
+ for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
+ tp = t, t = t->prev)
+ {
+ if (n->tgt == t)
+ {
+ if (tp)
+ tp->prev = t->prev;
+ else
+ acc_dev->openacc.data_environ = t->prev;
+ break;
+ }
+ }
+ }
+
+ if (force_copyfrom)
+ t->list[0]->copy_from = 1;
+
+ gomp_mutex_unlock (&acc_dev->mem_map.lock);
+
+ /* If running synchronously, unmap immediately. */
+ if (async < acc_async_noval)
+ gomp_unmap_vars (t, true);
+ else
+ {
+ gomp_copy_from_async (t);
+ acc_dev->openacc.register_async_cleanup_func (t);
+ }
+
+ gomp_notify (" %s: mappings restored\n", __FUNCTION__);
+}
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 1639244..6dcab05 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -72,6 +72,18 @@ dump_var (char *s, size_t idx, void *hostaddr, size_t size, unsigned char kind)
return;
}
+static int
+find_pset (int pos, size_t mapnum, unsigned short *kinds)
+{
+ if (pos + 1 >= mapnum)
+ return 0;
+
+ unsigned char kind = kinds[pos+1] & 0xff;
+
+ return kind == GOMP_MAP_TO_PSET;
+}
+
+
/* Ensure that the target device for DEVICE_TYPE is initialised (and that
plugins have been loaded if appropriate). The ACC_dev variable for the
current thread will be set appropriately for the given device type on
@@ -243,6 +255,143 @@ GOACC_data_end (void)
gomp_notify (" %s: mappings restored\n", __FUNCTION__);
}
+void
+GOACC_enter_exit_data (int device, const void *openmp_target, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds,
+ int async, int num_waits, ...)
+{
+ struct goacc_thread *thr;
+ struct gomp_device_descr *acc_dev;
+ bool if_clause_condition_value = device != GOMP_IF_CLAUSE_FALSE;
+ bool data_enter = false;
+ size_t i;
+
+ select_acc_device (device);
+
+ thr = goacc_thread ();
+ acc_dev = thr->dev;
+
+ if ((acc_dev->capabilities & TARGET_CAP_SHARED_MEM)
+ || !if_clause_condition_value)
+ return;
+
+ if (num_waits > 0)
+ {
+ va_list ap;
+
+ va_start (ap, num_waits);
+
+ goacc_wait (async, num_waits, ap);
+
+ va_end (ap);
+ }
+
+ acc_dev->openacc.async_set_async_func (async);
+
+ /* Determine if this is an "acc enter data". */
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+ continue;
+
+ if (kind == GOMP_MAP_FORCE_ALLOC || kind == GOMP_MAP_FORCE_PRESENT
+ || kind == GOMP_MAP_FORCE_TO)
+ {
+ data_enter = true;
+ break;
+ }
+
+ if (kind == GOMP_MAP_FORCE_DEALLOC || kind == GOMP_MAP_FORCE_FROM)
+ break;
+
+ gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
+ kind);
+ }
+
+ if (data_enter)
+ {
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ /* Scan for PSETs. */
+ int psets = find_pset (i, mapnum, kinds);
+
+ if (!psets)
+ {
+ switch (kind)
+ {
+ case GOMP_MAP_POINTER:
+ gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i],
+ &kinds[i]);
+ break;
+ case GOMP_MAP_FORCE_ALLOC:
+ acc_create (hostaddrs[i], sizes[i]);
+ break;
+ case GOMP_MAP_FORCE_PRESENT:
+ acc_present_or_copyin (hostaddrs[i], sizes[i]);
+ break;
+ case GOMP_MAP_FORCE_TO:
+ acc_present_or_copyin (hostaddrs[i], sizes[i]);
+ break;
+ default:
+ gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
+ kind);
+ break;
+ }
+ }
+ else
+ {
+ gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]);
+ /* Increment 'i' by two because OpenACC requires fortran
+ arrays to be contiguous, so each PSET is associated with
+ one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and
+ one MAP_POINTER. */
+ i += 2;
+ }
+ }
+ }
+ else
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ int psets = find_pset (i, mapnum, kinds);
+
+ if (!psets)
+ {
+ switch (kind)
+ {
+ case GOMP_MAP_POINTER:
+ gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
+ == GOMP_MAP_FORCE_FROM,
+ async, 1);
+ break;
+ case GOMP_MAP_FORCE_DEALLOC:
+ acc_delete (hostaddrs[i], sizes[i]);
+ break;
+ case GOMP_MAP_FORCE_FROM:
+ acc_copyout (hostaddrs[i], sizes[i]);
+ break;
+ default:
+ gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
+ kind);
+ break;
+ }
+ }
+ else
+ {
+ gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
+ == GOMP_MAP_FORCE_FROM, async, 3);
+ /* See the above comment. */
+ i += 2;
+ }
+ }
+
+ acc_dev->openacc.async_set_async_func (acc_async_sync);
+}
void
GOACC_kernels (int device, void (*fn) (void *), const void *openmp_target,
@@ -359,6 +508,7 @@ GOACC_update (int device, const void *openmp_target, size_t mapnum,
switch (kind)
{
case GOMP_MAP_POINTER:
+ case GOMP_MAP_TO_PSET:
break;
case GOMP_MAP_FORCE_TO:
diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
index 53b3c23..b8b3e85 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -14,6 +14,10 @@ if [info exists lang_include_flags] then {
unset lang_include_flags
}
+proc check_effective_target_oacc_c { } {
+ return 0
+}
+
# Initialize dg.
dg-init
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
new file mode 100644
index 0000000..b990ade
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
@@ -0,0 +1,163 @@
+/* { dg-do run } */
+/* { dg-additional-options "-std=c99" { target oacc_c } } */
+
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+ int N = 128; //1024 * 1024;
+ float *a, *b, *c, *d, *e;
+ int i;
+ int nbytes;
+
+ nbytes = N * sizeof (float);
+
+ a = (float *) malloc (nbytes);
+ b = (float *) malloc (nbytes);
+ c = (float *) malloc (nbytes);
+ d = (float *) malloc (nbytes);
+ e = (float *) malloc (nbytes);
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
+#pragma acc parallel async wait
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
+#pragma acc wait
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
+#pragma acc parallel async (1)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 2.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1)
+
+#pragma acc parallel async (1) wait (1)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel async (2) wait (1)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel async (3) wait (1)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ e[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
+
+#pragma acc parallel async (1) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel async (2) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel async (3) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel wait (1) async (4)
+ for (int ii = 0; ii < N; ii++)
+ e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc wait (1)
+
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 4.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+
+ if (e[i] != 11.0)
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
new file mode 100644
index 0000000..f8f1b3b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c
@@ -0,0 +1,167 @@
+/* { dg-do run } */
+/* { dg-additional-options "-std=c99" { target oacc_c } } */
+
+#include <stdlib.h>
+
+int
+main (int argc, char **argv)
+{
+ int N = 128; //1024 * 1024;
+ float *a, *b, *c, *d, *e;
+ int i;
+ int nbytes;
+
+ nbytes = N * sizeof (float);
+
+ a = (float *) malloc (nbytes);
+ b = (float *) malloc (nbytes);
+ c = (float *) malloc (nbytes);
+ d = (float *) malloc (nbytes);
+ e = (float *) malloc (nbytes);
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
+#pragma acc parallel async wait
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+
+#pragma acc update host (a[0:N], b[0:N]) async wait
+#pragma acc wait
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N]) async (1)
+#pragma acc parallel async (1)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+
+#pragma acc update host (a[0:N], b[0:N]) async (1) wait (1)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 2.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N]) async (1)
+#pragma acc update device (b[0:N]) async (2)
+#pragma acc enter data copyin (c[0:N], d[0:N]) async (3)
+
+#pragma acc parallel async (1) wait (1,2)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel async (2) wait (1,3)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel async (3) wait (1,3)
+#pragma acc loop
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc update host (a[0:N], b[0:N], c[0:N], d[0:N]) async (1) wait (1,2,3)
+#pragma acc wait (1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 2.0;
+ b[i] = 0.0;
+ c[i] = 0.0;
+ d[i] = 0.0;
+ e[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N], c[0:N], d[0:N]) async (1)
+#pragma acc enter data copyin (e[0:N]) async (5)
+
+#pragma acc parallel async (1) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
+
+#pragma acc parallel async (2) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
+
+#pragma acc parallel async (3) wait (1)
+ for (int ii = 0; ii < N; ii++)
+ d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
+
+#pragma acc parallel wait (1,5) async (4)
+ for (int ii = 0; ii < N; ii++)
+ e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
+
+#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
+#pragma acc delete (N)
+#pragma acc wait (1)
+
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 2.0)
+ abort ();
+
+ if (b[i] != 4.0)
+ abort ();
+
+ if (c[i] != 4.0)
+ abort ();
+
+ if (d[i] != 1.0)
+ abort ();
+
+ if (e[i] != 11.0)
+ abort ();
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp
index ea79ed0..5558ec8 100644
--- a/libgomp/testsuite/libgomp.oacc-c/c.exp
+++ b/libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -19,6 +19,10 @@ if ![info exists DEFAULT_CFLAGS] then {
set DEFAULT_CFLAGS "-O2"
}
+proc check_effective_target_oacc_c { } {
+ return 1
+}
+
# Initialize dg.
dg-init
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-1.f90
new file mode 100644
index 0000000..5e94e2d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-1.f90
@@ -0,0 +1,45 @@
+! { dg-do run }
+
+program test
+ integer, parameter :: N = 8
+ real, allocatable :: a(:), b(:)
+
+ allocate (a(N))
+ allocate (b(N))
+
+ a(:) = 3.0
+ b(:) = 0.0
+
+ !$acc enter data copyin (a(1:N), b(1:N))
+
+ !$acc parallel
+ do i = 1, n
+ b(i) = a (i)
+ end do
+ !$acc end parallel
+
+ !$acc exit data copyout (a(1:N), b(1:N))
+
+ do i = 1, n
+ if (a(i) .ne. 3.0) call abort
+ if (b(i) .ne. 3.0) call abort
+ end do
+
+ a(:) = 5.0
+ b(:) = 1.0
+
+ !$acc enter data copyin (a(1:N), b(1:N))
+
+ !$acc parallel
+ do i = 1, n
+ b(i) = a (i)
+ end do
+ !$acc end parallel
+
+ !$acc exit data copyout (a(1:N), b(1:N))
+
+ do i = 1, n
+ if (a(i) .ne. 5.0) call abort
+ if (b(i) .ne. 5.0) call abort
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90
new file mode 100644
index 0000000..8736c2a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90
@@ -0,0 +1,31 @@
+! { dg-do run }
+
+program test
+ integer, parameter :: N = 8
+ real, allocatable :: a(:,:), b(:,:)
+
+ allocate (a(N,N))
+ allocate (b(N,N))
+
+ a(:,:) = 3.0
+ b(:,:) = 0.0
+
+ !$acc enter data copyin (a(1:N,1:N), b(1:N,1:N))
+
+ !$acc parallel
+ do i = 1, n
+ do j = 1, n
+ b(j,i) = a (j,i)
+ end do
+ end do
+ !$acc end parallel
+
+ !$acc exit data copyout (a(1:N,1:N), b(1:N,1:N))
+
+ do i = 1, n
+ do j = 1, n
+ if (a(j,i) .ne. 3.0) call abort
+ if (b(j,i) .ne. 3.0) call abort
+ end do
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-3.f90
new file mode 100644
index 0000000..9868cb0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-3.f90
@@ -0,0 +1,131 @@
+! { dg-do run }
+
+program asyncwait
+ real, allocatable :: a(:), b(:), c(:), d(:), e(:)
+ integer i, N
+
+ N = 64
+
+ allocate (a(N))
+ allocate (b(N))
+ allocate (c(N))
+ allocate (d(N))
+ allocate (e(N))
+
+ a(:) = 3.0
+ b(:) = 0.0
+
+ !$acc enter data copyin (a(1:N)) copyin (b(1:N)) copyin (N) async
+
+ !$acc parallel async wait
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc wait
+ !$acc exit data copyout (a(1:N)) copyout (b(1:N))
+
+ do i = 1, N
+ if (a(i) .ne. 3.0) call abort
+ if (b(i) .ne. 3.0) call abort
+ end do
+
+ a(:) = 2.0
+ b(:) = 0.0
+
+ !$acc enter data copyin (a(1:N)) copyin (b(1:N)) async (1)
+
+ !$acc parallel async (1) wait (1)
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc wait (1)
+ !$acc exit data copyout (a(1:N)) copyout (b(1:N))
+
+ do i = 1, N
+ if (a(i) .ne. 2.0) call abort
+ if (b(i) .ne. 2.0) call abort
+ end do
+
+ a(:) = 3.0
+ b(:) = 0.0
+ c(:) = 0.0
+ d(:) = 0.0
+
+ !$acc enter data copyin (a(1:N)) create (b(1:N)) create (c(1:N)) create (d(1:N))
+
+ !$acc parallel async (1)
+ do i = 1, N
+ b(i) = (a(i) * a(i) * a(i)) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ c(i) = (a(i) * 4) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
+ end do
+ !$acc end parallel
+
+ !$acc wait (1)
+ !$acc exit data copyout (a(1:N)) copyout (b(1:N)) copyout (c(1:N)) copyout (d(1:N))
+
+ do i = 1, N
+ if (a(i) .ne. 3.0) call abort
+ if (b(i) .ne. 9.0) call abort
+ if (c(i) .ne. 4.0) call abort
+ if (d(i) .ne. 1.0) call abort
+ end do
+
+ a(:) = 2.0
+ b(:) = 0.0
+ c(:) = 0.0
+ d(:) = 0.0
+ e(:) = 0.0
+
+ !$acc enter data copyin (a(1:N)) create (b(1:N)) create (c(1:N)) create (d(1:N)) copyin (e(1:N))
+
+ !$acc parallel async (1)
+ do i = 1, N
+ b(i) = (a(i) * a(i) * a(i)) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ c(i) = (a(i) * 4) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel wait (1) async (1)
+ do i = 1, N
+ e(i) = a(i) + b(i) + c(i) + d(i)
+ end do
+ !$acc end parallel
+
+ !$acc wait (1)
+ !$acc exit data copyout (a(1:N)) copyout (b(1:N)) copyout (c(1:N)) copyout (d(1:N)) copyout (e(1:N))
+ !$acc exit data delete (N)
+
+ do i = 1, N
+ if (a(i) .ne. 2.0) call abort
+ if (b(i) .ne. 4.0) call abort
+ if (c(i) .ne. 4.0) call abort
+ if (d(i) .ne. 1.0) call abort
+ if (e(i) .ne. 11.0) call abort
+ end do
+end program asyncwait
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
new file mode 100644
index 0000000..41c45fb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
@@ -0,0 +1,136 @@
+! { dg-do run }
+
+program asyncwait
+ real, allocatable :: a(:), b(:), c(:), d(:), e(:)
+ integer i, N
+
+ N = 64
+
+ allocate (a(N))
+ allocate (b(N))
+ allocate (c(N))
+ allocate (d(N))
+ allocate (e(N))
+
+ a(:) = 3.0
+ b(:) = 0.0
+
+ !$acc enter data copyin (a(1:N)) copyin (b(1:N)) copyin (N) async
+
+ !$acc parallel async wait
+ !$acc loop
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc update host (a(1:N), b(1:N)) async wait
+ !$acc wait
+
+ do i = 1, N
+ if (a(i) .ne. 3.0) call abort
+ if (b(i) .ne. 3.0) call abort
+ end do
+
+ a(:) = 2.0
+ b(:) = 0.0
+
+ !$acc update device (a(1:N), b(1:N)) async (1)
+
+ !$acc parallel async (1) wait (1)
+ !$acc loop
+ do i = 1, N
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc update host (a(1:N), b(1:N)) async (1) wait (1)
+ !$acc wait (1)
+
+ do i = 1, N
+ if (a(i) .ne. 2.0) call abort
+ if (b(i) .ne. 2.0) call abort
+ end do
+
+ a(:) = 3.0
+ b(:) = 0.0
+ c(:) = 0.0
+ d(:) = 0.0
+
+ !$acc enter data copyin (c(1:N), d(1:N)) async (1)
+ !$acc update device (a(1:N), b(1:N)) async (1)
+
+ !$acc parallel async (1)
+ do i = 1, N
+ b(i) = (a(i) * a(i) * a(i)) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ c(i) = (a(i) * 4) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
+ end do
+ !$acc end parallel
+
+ !$acc update host (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) wait (1)
+
+ !$acc wait (1)
+
+ do i = 1, N
+ if (a(i) .ne. 3.0) call abort
+ if (b(i) .ne. 9.0) call abort
+ if (c(i) .ne. 4.0) call abort
+ if (d(i) .ne. 1.0) call abort
+ end do
+
+ a(:) = 2.0
+ b(:) = 0.0
+ c(:) = 0.0
+ d(:) = 0.0
+ e(:) = 0.0
+
+ !$acc enter data copyin (e(1:N)) async (1)
+ !$acc update device (a(1:N), b(1:N), c(1:N), d(1:N)) async (1)
+
+ !$acc parallel async (1)
+ do i = 1, N
+ b(i) = (a(i) * a(i) * a(i)) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ c(i) = (a(i) * 4) / a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel async (1)
+ do i = 1, N
+ d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel wait (1) async (1)
+ do i = 1, N
+ e(i) = a(i) + b(i) + c(i) + d(i)
+ end do
+ !$acc end parallel
+
+ !$acc update host (a(1:N), b(1:N), c(1:N), d(1:N), e(1:N)) async (1) wait (1)
+ !$acc wait (1)
+ !$acc exit data delete (N, a(1:N), b(1:N), c(1:N), d(1:N), e(1:N))
+
+ do i = 1, N
+ if (a(i) .ne. 2.0) call abort
+ if (b(i) .ne. 4.0) call abort
+ if (c(i) .ne. 4.0) call abort
+ if (d(i) .ne. 1.0) call abort
+ if (e(i) .ne. 11.0) call abort
+ end do
+end program asyncwait