* [gomp4] acc enter/exit data
@ 2014-10-31 0:51 Cesar Philippidis
2014-11-05 16:56 ` [gomp4] OpenACC update host/self maintenance (was: acc enter/exit data) Thomas Schwinge
` (2 more replies)
0 siblings, 3 replies; 6+ messages in thread
From: Cesar Philippidis @ 2014-10-31 0:51 UTC (permalink / raw)
To: gcc-patches, fortran
[-- Attachment #1: Type: text/plain, Size: 1458 bytes --]
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
[-- Attachment #2: data-internal-v2b.diff --]
[-- Type: text/x-patch, Size: 71650 bytes --]
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
^ permalink raw reply [flat|nested] 6+ messages in thread
* [gomp4] OpenACC update host/self maintenance (was: acc enter/exit data)
2014-10-31 0:51 [gomp4] acc enter/exit data Cesar Philippidis
@ 2014-11-05 16:56 ` Thomas Schwinge
2014-11-06 14:04 ` [gomp4] acc enter/exit data Thomas Schwinge
2014-12-10 9:54 ` Thomas Schwinge
2 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2014-11-05 16:56 UTC (permalink / raw)
To: gcc-patches, fortran; +Cc: Cesar Philippidis
[-- Attachment #1: Type: text/plain, Size: 28473 bytes --]
Hi!
On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> 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.
Applied to gomp-4_0-branch in r217148:
commit a7bba5ecc7c62a022616f55ff1d8fb48266fcb67
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Nov 5 16:54:07 2014 +0000
OpenACC update host/self maintenance.
gcc/c/
* c-parser.c (c_parser_omp_clause_name) <"host">: Return
PRAGMA_OMP_CLAUSE_HOST.
gcc/c/
(c_parser_oacc_data_clause): Group PRAGMA_OMP_CLAUSE_SELF next to
PRAGMA_OMP_CLAUSE_HOST.
gcc/cp/
* parser.c (cp_parser_oacc_data_clause): Group
PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST.
gcc/fortran/
* openmp.c (OMP_CLAUSE_HOST, OMP_CLAUSE_SELF): Merge into the new
OMP_CLAUSE_HOST_SELF. Update all users.
gcc/
* tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and
OMP_CLAUSE_OACC_DEVICE. Update all users.
gcc/testsuite/
* c-c++-common/goacc/update-1.c: Extend.
* gfortran.dg/goacc/assumed.f95: Likewise.
* gfortran.dg/goacc/coarray.f95: Likewise.
* gfortran.dg/goacc/cray.f95: Likewise.
* gfortran.dg/goacc/literal.f95: Likewise.
* gfortran.dg/goacc/parameter.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/update-1-2.c: New file.
* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/data-4.f90: In one instance, use
the self clause instead of host clause.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217148 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 3 +
gcc/c/ChangeLog.gomp | 6 +
gcc/c/c-parser.c | 10 +-
gcc/cp/ChangeLog.gomp | 3 +
gcc/cp/parser.c | 8 +-
gcc/fortran/ChangeLog.gomp | 3 +
gcc/fortran/openmp.c | 21 +-
gcc/gimplify.c | 4 -
gcc/omp-low.c | 4 -
gcc/testsuite/ChangeLog.gomp | 7 +
gcc/testsuite/c-c++-common/goacc/update-1.c | 5 +
gcc/testsuite/gfortran.dg/goacc/assumed.f95 | 8 +-
gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 3 +-
gcc/testsuite/gfortran.dg/goacc/cray.f95 | 6 +-
gcc/testsuite/gfortran.dg/goacc/literal.f95 | 5 +-
gcc/testsuite/gfortran.dg/goacc/parameter.f95 | 3 +-
gcc/tree-core.h | 10 +-
gcc/tree-pretty-print.c | 6 -
gcc/tree.c | 6 -
libgomp/ChangeLog.gomp | 5 +
.../libgomp.oacc-c-c++-common/update-1-2.c | 282 +++++++++++++++++++++
.../{data-4.f90 => data-4-2.f90} | 8 +-
libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 | 2 +-
23 files changed, 353 insertions(+), 65 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 2c2b349..d140a35 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * tree-core.h (enum omp_clause_code): Remove OMP_CLAUSE_HOST and
+ OMP_CLAUSE_OACC_DEVICE. Update all users.
+
* gimplify.c (gimplify_oacc_cache): New function.
(gimplify_expr): Use it for OACC_CACHE.
(gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses): Handle
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 70278b9..a223a17 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,5 +1,11 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * c-parser.c (c_parser_omp_clause_name) <"host">: Return
+ PRAGMA_OMP_CLAUSE_HOST.
+
+ * c-parser.c (c_parser_oacc_data_clause): Group
+ PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST.
+
* c-parser.c (c_parser_oacc_cache): Generate OACC_CACHE.
* c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 40d4314..bd2864f 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -9832,7 +9832,7 @@ c_parser_omp_clause_name (c_parser *parser)
break;
case 'h':
if (!strcmp ("host", p))
- result = PRAGMA_OMP_CLAUSE_SELF;
+ result = PRAGMA_OMP_CLAUSE_HOST;
break;
case 'i':
if (!strcmp ("inbranch", p))
@@ -10187,8 +10187,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
enum omp_clause_map_kind kind;
switch (c_kind)
{
- default:
- gcc_unreachable ();
case PRAGMA_OMP_CLAUSE_COPY:
kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
break;
@@ -10208,6 +10206,7 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
kind = OMP_CLAUSE_MAP_FORCE_TO;
break;
case PRAGMA_OMP_CLAUSE_HOST:
+ case PRAGMA_OMP_CLAUSE_SELF:
kind = OMP_CLAUSE_MAP_FORCE_FROM;
break;
case PRAGMA_OMP_CLAUSE_PRESENT:
@@ -10225,9 +10224,8 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
kind = OMP_CLAUSE_MAP_ALLOC;
break;
- case PRAGMA_OMP_CLAUSE_SELF:
- kind = OMP_CLAUSE_MAP_FORCE_FROM;
- break;
+ default:
+ gcc_unreachable ();
}
tree nl, c;
nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index 46d4912..f5d400f 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * parser.c (cp_parser_oacc_data_clause): Group
+ PRAGMA_OMP_CLAUSE_SELF next to PRAGMA_OMP_CLAUSE_HOST.
+
* parser.c (cp_parser_oacc_cache): Generate OACC_CACHE.
* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE__CACHE_.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index ea4ad2f..9c2c3ca 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -27812,8 +27812,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
enum omp_clause_map_kind kind;
switch (c_kind)
{
- default:
- gcc_unreachable ();
case PRAGMA_OMP_CLAUSE_COPY:
kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
break;
@@ -27833,6 +27831,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
kind = OMP_CLAUSE_MAP_FORCE_TO;
break;
case PRAGMA_OMP_CLAUSE_HOST:
+ case PRAGMA_OMP_CLAUSE_SELF:
kind = OMP_CLAUSE_MAP_FORCE_FROM;
break;
case PRAGMA_OMP_CLAUSE_PRESENT:
@@ -27850,9 +27849,8 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
kind = OMP_CLAUSE_MAP_ALLOC;
break;
- case PRAGMA_OMP_CLAUSE_SELF:
- kind = OMP_CLAUSE_MAP_FORCE_FROM;
- break;
+ default:
+ gcc_unreachable ();
}
tree nl, c;
nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list);
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index 98e3971..d10560e 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * openmp.c (OMP_CLAUSE_HOST, OMP_CLAUSE_SELF): Merge into the new
+ OMP_CLAUSE_HOST_SELF. Update all users.
+
* gfortran.texi: Update for OpenACC.
* intrinsic.texi: Likewise.
* invoke.texi: Likewise.
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index c7af004..959798a 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -445,13 +445,12 @@ match_oacc_clause_gang (gfc_omp_clauses *cp)
#define OMP_CLAUSE_INDEPENDENT (1ULL << 49)
#define OMP_CLAUSE_USE_DEVICE (1ULL << 50)
#define OMP_CLAUSE_DEVICE_RESIDENT (1ULL << 51)
-#define OMP_CLAUSE_HOST (1ULL << 52)
+#define OMP_CLAUSE_HOST_SELF (1ULL << 52)
#define OMP_CLAUSE_OACC_DEVICE (1ULL << 53)
#define OMP_CLAUSE_WAIT (1ULL << 54)
#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,24 +681,20 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, unsigned long long mask,
true)
== MATCH_YES)
continue;
- if ((mask & OMP_CLAUSE_HOST)
- && 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 ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FORCE_TO))
continue;
+ if ((mask & OMP_CLAUSE_HOST_SELF)
+ && (gfc_match ("host ( ") == MATCH_YES
+ || gfc_match ("self ( ") == MATCH_YES)
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_FORCE_FROM))
+ 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)
{
@@ -1170,7 +1165,7 @@ 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_SELF \
+ (OMP_CLAUSE_IF | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_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 \
diff --git gcc/gimplify.c gcc/gimplify.c
index d58876f..5aeb726 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6288,8 +6288,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
@@ -6692,8 +6690,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
case OMP_CLAUSE_VECTOR_LENGTH:
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
diff --git gcc/omp-low.c gcc/omp-low.c
index 1c9d942..0b45e69 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1977,8 +1977,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
@@ -2125,8 +2123,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_WAIT:
break;
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 1faf0fa..a02f58a 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,12 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * c-c++-common/goacc/update-1.c: Extend.
+ * gfortran.dg/goacc/assumed.f95: Likewise.
+ * gfortran.dg/goacc/coarray.f95: Likewise.
+ * gfortran.dg/goacc/cray.f95: Likewise.
+ * gfortran.dg/goacc/literal.f95: Likewise.
+ * gfortran.dg/goacc/parameter.f95: Likewise.
+
* c-c++-common/goacc/cache-1.c: New file.
* gfortran.dg/goacc/data-tree.f95: Remove dg-prune-output directive.
diff --git gcc/testsuite/c-c++-common/goacc/update-1.c gcc/testsuite/c-c++-common/goacc/update-1.c
index 2a3a910..97e9379 100644
--- gcc/testsuite/c-c++-common/goacc/update-1.c
+++ gcc/testsuite/c-c++-common/goacc/update-1.c
@@ -8,5 +8,10 @@ f (void)
#pragma acc update device(i)
#pragma acc update host(i)
#pragma acc update self(i)
+#pragma acc update device(a[1:3])
+#pragma acc update host(a[1:3])
+#pragma acc update self(a[1:3])
+#pragma acc update device(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
#pragma acc update host(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
+#pragma acc update self(a(1:3)) /* { dg-error "expected '\\\)' before '\\\(' token" } */
}
diff --git gcc/testsuite/gfortran.dg/goacc/assumed.f95 gcc/testsuite/gfortran.dg/goacc/assumed.f95
index 15bfa0c..3287241 100644
--- gcc/testsuite/gfortran.dg/goacc/assumed.f95
+++ gcc/testsuite/gfortran.dg/goacc/assumed.f95
@@ -19,8 +19,9 @@ contains
do i = 1,5
enddo
!$acc end parallel loop
- !$acc update host (a) ! { dg-error "Assumed size" }
!$acc update device (a) ! { dg-error "Assumed size" }
+ !$acc update host (a) ! { dg-error "Assumed size" }
+ !$acc update self (a) ! { dg-error "Assumed size" }
end subroutine assumed_size
subroutine assumed_rank(a)
implicit none
@@ -39,7 +40,8 @@ contains
do i = 1,5
enddo
!$acc end parallel loop
- !$acc update host (a) ! { dg-error "Assumed rank" }
!$acc update device (a) ! { dg-error "Assumed rank" }
+ !$acc update host (a) ! { dg-error "Assumed rank" }
+ !$acc update self (a) ! { dg-error "Assumed rank" }
end subroutine assumed_rank
-end module test
\ No newline at end of file
+end module test
diff --git gcc/testsuite/gfortran.dg/goacc/coarray.f95 gcc/testsuite/gfortran.dg/goacc/coarray.f95
index ab13157..4f1224e 100644
--- gcc/testsuite/gfortran.dg/goacc/coarray.f95
+++ gcc/testsuite/gfortran.dg/goacc/coarray.f95
@@ -27,8 +27,9 @@ contains
!$acc cache (a)
enddo
!$acc end parallel loop
- !$acc update host (a)
!$acc update device (a)
+ !$acc update host (a)
+ !$acc update self (a)
end subroutine oacc1
end module test
! { dg-prune-output "ACC cache unimplemented" }
diff --git gcc/testsuite/gfortran.dg/goacc/cray.f95 gcc/testsuite/gfortran.dg/goacc/cray.f95
index 3225b28..8f2c077 100644
--- gcc/testsuite/gfortran.dg/goacc/cray.f95
+++ gcc/testsuite/gfortran.dg/goacc/cray.f95
@@ -28,8 +28,9 @@ contains
!$acc cache (pointee) ! TODO: This must fail, as in openacc-1_0-branch
enddo
!$acc end parallel loop
- !$acc update host (pointee) ! { dg-error "Cray pointee" }
!$acc update device (pointee) ! { dg-error "Cray pointee" }
+ !$acc update host (pointee) ! { dg-error "Cray pointee" }
+ !$acc update self (pointee) ! { dg-error "Cray pointee" }
!$acc data copy (ptr)
!$acc end data
!$acc data deviceptr (ptr) ! { dg-error "Cray pointer" }
@@ -47,8 +48,9 @@ contains
!$acc cache (ptr) ! TODO: This must fail, as in openacc-1_0-branch
enddo
!$acc end parallel loop
- !$acc update host (ptr)
!$acc update device (ptr)
+ !$acc update host (ptr)
+ !$acc update self (ptr)
end subroutine oacc1
end module test
! { dg-prune-output "unimplemented" }
diff --git gcc/testsuite/gfortran.dg/goacc/literal.f95 gcc/testsuite/gfortran.dg/goacc/literal.f95
index bdbf66d..e6760d0 100644
--- gcc/testsuite/gfortran.dg/goacc/literal.f95
+++ gcc/testsuite/gfortran.dg/goacc/literal.f95
@@ -23,7 +23,8 @@ contains
!$acc cache (10) ! { dg-error "Syntax error" }
enddo
!$acc end parallel loop
- !$acc update host (10) ! { dg-error "Syntax error" }
!$acc update device (10) ! { dg-error "Syntax error" }
+ !$acc update host (10) ! { dg-error "Syntax error" }
+ !$acc update self (10) ! { dg-error "Syntax error" }
end subroutine oacc1
-end module test
\ No newline at end of file
+end module test
diff --git gcc/testsuite/gfortran.dg/goacc/parameter.f95 gcc/testsuite/gfortran.dg/goacc/parameter.f95
index 785d7f9..1364181 100644
--- gcc/testsuite/gfortran.dg/goacc/parameter.f95
+++ gcc/testsuite/gfortran.dg/goacc/parameter.f95
@@ -24,8 +24,9 @@ contains
!$acc cache (a) ! TODO: This must fail, as in openacc-1_0-branch
enddo
!$acc end parallel loop
- !$acc update host (a) ! { dg-error "not a variable" }
!$acc update device (a) ! { dg-error "not a variable" }
+ !$acc update host (a) ! { dg-error "not a variable" }
+ !$acc update self (a) ! { dg-error "not a variable" }
end subroutine oacc1
end module test
! { dg-prune-output "unimplemented" }
diff --git gcc/tree-core.h gcc/tree-core.h
index 42ad6a0..cd38861 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -259,8 +259,8 @@ enum omp_clause_code {
OMP_CLAUSE_TO,
/* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr,
- present, present_or_copy (pcopy), present_or_copyin (pcopyin),
- present_or_copyout (pcopyout), present_or_create (pcreate)}
+ device, host (self), present, present_or_copy (pcopy), present_or_copyin
+ (pcopyin), present_or_copyout (pcopyout), present_or_create (pcreate)}
(variable-list).
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
@@ -270,12 +270,6 @@ enum omp_clause_code {
#pragma acc cache (variable-list). */
OMP_CLAUSE__CACHE_,
- /* OpenACC clause: host (variable_list). */
- OMP_CLAUSE_HOST,
-
- /* OpenACC clause: device (variable_list). */
- OMP_CLAUSE_OACC_DEVICE,
-
/* OpenACC clause: device_resident (variable_list). */
OMP_CLAUSE_DEVICE_RESIDENT,
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index d678f36..1458913 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -335,12 +335,6 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE__LOOPTEMP_:
name = "_looptemp_";
goto print_remap;
- case OMP_CLAUSE_HOST:
- name = "host";
- goto print_remap;
- case OMP_CLAUSE_OACC_DEVICE:
- name = "device";
- goto print_remap;
case OMP_CLAUSE_DEVICE_RESIDENT:
name = "device_resident";
goto print_remap;
diff --git gcc/tree.c gcc/tree.c
index f39c63f..16d156b 100644
--- gcc/tree.c
+++ gcc/tree.c
@@ -271,8 +271,6 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
2, /* OMP_CLAUSE__CACHE_ */
- 1, /* OMP_CLAUSE_HOST */
- 1, /* OMP_CLAUSE_OACC_DEVICE */
1, /* OMP_CLAUSE_DEVICE_RESIDENT */
1, /* OMP_CLAUSE_USE_DEVICE */
1, /* OMP_CLAUSE_GANG */
@@ -330,8 +328,6 @@ const char * const omp_clause_code_name[] =
"to",
"map",
"_cache_",
- "host",
- "device",
"device_resident",
"use_device",
"gang",
@@ -11120,8 +11116,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
switch (OMP_CLAUSE_CODE (*tp))
{
- case OMP_CLAUSE_HOST:
- case OMP_CLAUSE_OACC_DEVICE:
case OMP_CLAUSE_DEVICE_RESIDENT:
case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_GANG:
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 096a2a9..d9b92fe 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,10 @@
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: New file.
+ * testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-4.f90: In one instance, use
+ the self clause instead of host clause.
+
* testsuite/libgomp.oacc-c/cache-1.c: Remove directives that are
expected to fail, and rename the file to...
* testsuite/libgomp.oacc-c-c++-common/cache-1.c: ... this.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
new file mode 100644
index 0000000..c7e7257
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/update-1-2.c
@@ -0,0 +1,282 @@
+/* Copy of update-1.c with self exchanged with host for #pragma acc update. */
+
+/* { dg-do run } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+#include <string.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdbool.h>
+
+int
+main (int argc, char **argv)
+{
+ int N = 8;
+ float *a, *b, *c;
+ float *d_a, *d_b, *d_c;
+ int i;
+
+ a = (float *) malloc (N * sizeof (float));
+ b = (float *) malloc (N * sizeof (float));
+ c = (float *) malloc (N * sizeof (float));
+
+ d_a = (float *) acc_malloc (N * sizeof (float));
+ d_b = (float *) acc_malloc (N * sizeof (float));
+ d_c = (float *) acc_malloc (N * sizeof (float));
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 3.0;
+ b[i] = 0.0;
+ }
+
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ acc_map_data (c, d_c, N * sizeof (float));
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 3.0)
+ abort ();
+
+ if (b[i] != 3.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ b[i] = 1.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update host (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ b[i] = 0.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 7.0;
+ b[i] = 2.0;
+ }
+
+#pragma acc update device (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 7.0)
+ abort ();
+
+ if (b[i] != 7.0)
+ abort ();
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 9.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ if (a[i] != 9.0)
+ abort ();
+
+ if (b[i] != 9.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 5.0;
+ }
+
+#pragma acc update device (a[0:N])
+
+ for (i = 0; i < N; i++)
+ {
+ a[i] = 6.0;
+ }
+
+#pragma acc update device (a[0:N >> 1])
+
+#pragma acc parallel present (a[0:N], b[0:N])
+ {
+ int ii;
+
+ for (ii = 0; ii < N; ii++)
+ b[ii] = a[ii];
+ }
+
+#pragma acc update self (a[0:N], b[0:N])
+
+ for (i = 0; i < (N >> 1); i++)
+ {
+ if (a[i] != 6.0)
+ abort ();
+
+ if (b[i] != 6.0)
+ abort ();
+ }
+
+ for (i = (N >> 1); i < N; i++)
+ {
+ if (a[i] != 5.0)
+ abort ();
+
+ if (b[i] != 5.0)
+ abort ();
+ }
+
+ if (!acc_is_present (&a[0], (N * sizeof (float))))
+ abort ();
+
+ if (!acc_is_present (&b[0], (N * sizeof (float))))
+ abort ();
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
similarity index 91%
copy from libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
copy to libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
index 41c45fb..16a8598 100644
--- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
@@ -1,3 +1,5 @@
+! Copy of data-4.f90 with self exchanged with host for !acc update.
+
! { dg-do run }
program asyncwait
@@ -24,7 +26,7 @@ program asyncwait
end do
!$acc end parallel
- !$acc update host (a(1:N), b(1:N)) async wait
+ !$acc update self (a(1:N), b(1:N)) async wait
!$acc wait
do i = 1, N
@@ -78,7 +80,7 @@ program asyncwait
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 update self (a(1:N), b(1:N), c(1:N), d(1:N)) async (1) wait (1)
!$acc wait (1)
@@ -122,7 +124,7 @@ program asyncwait
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 update self (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))
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-4.f90 libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
index 41c45fb..f6886b0 100644
--- libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/data-4.f90
@@ -44,7 +44,7 @@ program asyncwait
end do
!$acc end parallel
- !$acc update host (a(1:N), b(1:N)) async (1) wait (1)
+ !$acc update self (a(1:N), b(1:N)) async (1) wait (1)
!$acc wait (1)
do i = 1, N
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [gomp4] acc enter/exit data
2014-10-31 0:51 [gomp4] acc enter/exit data Cesar Philippidis
2014-11-05 16:56 ` [gomp4] OpenACC update host/self maintenance (was: acc enter/exit data) Thomas Schwinge
@ 2014-11-06 14:04 ` Thomas Schwinge
2014-12-10 9:54 ` Thomas Schwinge
2 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2014-11-06 14:04 UTC (permalink / raw)
To: gcc-patches; +Cc: Cesar Philippidis
[-- Attachment #1: Type: text/plain, Size: 1805 bytes --]
Hi!
On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch add support for OpenACC's enter/exit data directive.
> gcc/
> * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
Applied to gomp-4_0-branch in r217190:
commit 81c45b54c802bd76efc941750c73ca5410e83420
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Nov 6 14:02:27 2014 +0000
Correctly classify OpenACC's enter/exit data directives.
gcc/
* gimple.h (is_gimple_omp_oacc_specifically): Return true for
GIMPLE_OMP_TARGET's GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217190 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 3 +++
gcc/gimple.h | 2 +-
2 files changed, 4 insertions(+), 1 deletion(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 0bc9080..9c997ce 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-11-06 Thomas Schwinge <thomas@codesourcery.com>
+ * gimple.h (is_gimple_omp_oacc_specifically): Return true for
+ GIMPLE_OMP_TARGET's GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
+
* omp-low.c (lower_reduction_clauses): Initialize tid.
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
diff --git gcc/gimple.h gcc/gimple.h
index 7bc673a..4faeaaa 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -5867,8 +5867,8 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
switch (gimple_omp_target_kind (stmt))
{
case GF_OMP_TARGET_KIND_OACC_DATA:
- return true;
case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
return true;
default:
return false;
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [gomp4] acc enter/exit data
2014-10-31 0:51 [gomp4] acc enter/exit data Cesar Philippidis
2014-11-05 16:56 ` [gomp4] OpenACC update host/self maintenance (was: acc enter/exit data) Thomas Schwinge
2014-11-06 14:04 ` [gomp4] acc enter/exit data Thomas Schwinge
@ 2014-12-10 9:54 ` Thomas Schwinge
2014-12-10 9:59 ` Jakub Jelinek
2 siblings, 1 reply; 6+ messages in thread
From: Thomas Schwinge @ 2014-12-10 9:54 UTC (permalink / raw)
To: gcc-patches; +Cc: Cesar Philippidis, fortran
[-- Attachment #1: Type: text/plain, Size: 3684 bytes --]
Hi!
On Thu, 30 Oct 2014 17:11:04 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch add support for OpenACC's enter/exit data directive. [...]
> gcc/
> * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA.
In r218567, I committed the following to gomp-4_0-branch:
commit 86724db93ad780106102573f2cfadd6f884e8650
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Dec 10 09:52:14 2014 +0000
Fix OpenACC enter/exit data ICE.
[...]: In function 'f_acc_data':
[...]:4:1: internal compiler error: in expand_gimple_stmt_1, at cfgexpand.c:3413
f_acc_data (void)
^
0x70cad3 expand_gimple_stmt_1
[...]/source-gcc/gcc/cfgexpand.c:3413
0x70cad3 expand_gimple_stmt
[...]/source-gcc/gcc/cfgexpand.c:3440
0x712b3d expand_gimple_basic_block
[...]/source-gcc/gcc/cfgexpand.c:5273
0x71479e execute
[...]/source-gcc/gcc/cfgexpand.c:5882
gcc/
* omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
<GIMPLE_OMP_TARGET>: Handle
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA the same as
GF_OMP_TARGET_KIND_OACC_UPDATE.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218567 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 7 +++++++
gcc/omp-low.c | 8 ++++++--
gcc/testsuite/c-c++-common/goacc/nesting-2.c | 11 +++++++++++
3 files changed, 24 insertions(+), 2 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index af59ada..bece7c1 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
+
+ * omp-low.c (build_omp_regions_1, make_gimple_omp_edges)
+ <GIMPLE_OMP_TARGET>: Handle
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA the same as
+ GF_OMP_TARGET_KIND_OACC_UPDATE.
+
2014-11-13 Cesar Philippidis <cesar@codesourcery.com>
* omp-low.c (oacc_get_reduction_array_id): Fix whitespace.
diff --git gcc/omp-low.c gcc/omp-low.c
index 9af3b8a..6fed38f 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -9404,7 +9404,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
else if (code == GIMPLE_OMP_TARGET
&& (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE
|| (gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_OACC_UPDATE)))
+ == GF_OMP_TARGET_KIND_OACC_UPDATE)
+ || (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)))
new_omp_region (bb, code, parent);
else
{
@@ -12270,7 +12272,9 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
cur_region = new_omp_region (bb, code, cur_region);
fallthru = true;
if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
- || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE)
+ || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE
+ || (gimple_omp_target_kind (last)
+ == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA))
cur_region = cur_region->outer;
break;
diff --git gcc/testsuite/c-c++-common/goacc/nesting-2.c gcc/testsuite/c-c++-common/goacc/nesting-2.c
new file mode 100644
index 0000000..0d350c6
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-2.c
@@ -0,0 +1,11 @@
+int i;
+
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma acc update host(i)
+#pragma acc enter data copyin(i)
+ }
+}
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [gomp4] acc enter/exit data
2014-12-10 9:54 ` Thomas Schwinge
@ 2014-12-10 9:59 ` Jakub Jelinek
2014-12-17 22:18 ` Thomas Schwinge
0 siblings, 1 reply; 6+ messages in thread
From: Jakub Jelinek @ 2014-12-10 9:59 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Cesar Philippidis, fortran
On Wed, Dec 10, 2014 at 10:54:13AM +0100, Thomas Schwinge wrote:
> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -9404,7 +9404,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
> else if (code == GIMPLE_OMP_TARGET
> && (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE
> || (gimple_omp_target_kind (stmt)
> - == GF_OMP_TARGET_KIND_OACC_UPDATE)))
> + == GF_OMP_TARGET_KIND_OACC_UPDATE)
> + || (gimple_omp_target_kind (stmt)
> + == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)))
> new_omp_region (bb, code, parent);
> else
> {
> @@ -12270,7 +12272,9 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
> cur_region = new_omp_region (bb, code, cur_region);
> fallthru = true;
> if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
> - || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE)
> + || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE
> + || (gimple_omp_target_kind (last)
> + == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA))
I'd say that at this point a
switch (gimple_omp_target_kind (last))
{
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
...
default:
...
}
would be cleaner. The first hunk is more questionable, because there is
else and it would require duplicating of the else body in default:, goto
or similar, but perhaps it would be better that way too.
Jakub
^ permalink raw reply [flat|nested] 6+ messages in thread
* Re: [gomp4] acc enter/exit data
2014-12-10 9:59 ` Jakub Jelinek
@ 2014-12-17 22:18 ` Thomas Schwinge
0 siblings, 0 replies; 6+ messages in thread
From: Thomas Schwinge @ 2014-12-17 22:18 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches, Cesar Philippidis, fortran
[-- Attachment #1: Type: text/plain, Size: 5732 bytes --]
Hi!
On Wed, 10 Dec 2014 10:59:34 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Wed, Dec 10, 2014 at 10:54:13AM +0100, Thomas Schwinge wrote:
> > --- gcc/omp-low.c
> > +++ gcc/omp-low.c
> > @@ -9404,7 +9404,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
> > else if (code == GIMPLE_OMP_TARGET
> > && (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE
> > || (gimple_omp_target_kind (stmt)
> > - == GF_OMP_TARGET_KIND_OACC_UPDATE)))
> > + == GF_OMP_TARGET_KIND_OACC_UPDATE)
> > + || (gimple_omp_target_kind (stmt)
> > + == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)))
> > new_omp_region (bb, code, parent);
> > else
> > {
> > @@ -12270,7 +12272,9 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
> > cur_region = new_omp_region (bb, code, cur_region);
> > fallthru = true;
> > if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
> > - || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE)
> > + || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE
> > + || (gimple_omp_target_kind (last)
> > + == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA))
>
> I'd say that at this point a
> switch (gimple_omp_target_kind (last))
> {
> case GF_OMP_TARGET_KIND_UPDATE:
> case GF_OMP_TARGET_KIND_OACC_UPDATE:
> case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
> ...
> default:
> ...
> }
> would be cleaner. The first hunk is more questionable, because there is
> else and it would require duplicating of the else body in default:, goto
> or similar, but perhaps it would be better that way too.
Thanks for the suggestion. I found a way to express the first one
differently; committed to gomp-4_0-branch in r218837:
commit c9c55fd5c318f0ed6b866930d445a3df4aa058e8
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Dec 17 22:14:58 2014 +0000
Simplify multi-line if conditions.
gcc/
* omp-low.c (build_omp_regions_1, make_gimple_omp_edges): Simplify
multi-line if conditions.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218837 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 5 +++++
gcc/omp-low.c | 57 ++++++++++++++++++++++++++++++++++++++----------------
2 files changed, 45 insertions(+), 17 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 970e744..f925902 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+
+ * omp-low.c (build_omp_regions_1, make_gimple_omp_edges): Simplify
+ multi-line if conditions.
+
2014-12-10 Thomas Schwinge <thomas@codesourcery.com>
* omp-low.c (scan_omp_target): Remove taskreg_nesting_level and
diff --git gcc/omp-low.c gcc/omp-low.c
index a1fbccf..fd117dc 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -9379,7 +9379,6 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
region->exit = bb;
parent = parent->outer;
}
-
else if (code == GIMPLE_OMP_CONTINUE)
{
gcc_assert (parent);
@@ -9389,21 +9388,34 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
{
/* GIMPLE_OMP_SECTIONS_SWITCH is part of
GIMPLE_OMP_SECTIONS, and we do nothing for it. */
- ;
}
- else if (code == GIMPLE_OMP_TARGET
- && (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE
- || (gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_OACC_UPDATE)
- || (gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA)))
- new_omp_region (bb, code, parent);
else
{
- /* Otherwise, this directive becomes the parent for a new
- region. */
region = new_omp_region (bb, code, parent);
- parent = region;
+ /* Otherwise... */
+ if (code == GIMPLE_OMP_TARGET)
+ {
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ /* ..., other than for those stand-alone directives... */
+ region = NULL;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ }
+ /* ..., this directive becomes the parent for a new region. */
+ if (region)
+ parent = region;
}
}
@@ -12259,11 +12271,22 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GIMPLE_OMP_TARGET:
cur_region = new_omp_region (bb, code, cur_region);
fallthru = true;
- if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
- || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE
- || (gimple_omp_target_kind (last)
- == GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA))
- cur_region = cur_region->outer;
+ switch (gimple_omp_target_kind (last))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ cur_region = cur_region->outer;
+ break;
+ default:
+ gcc_unreachable ();
+ }
break;
case GIMPLE_OMP_SECTIONS:
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 6+ messages in thread
end of thread, other threads:[~2014-12-17 22:17 UTC | newest]
Thread overview: 6+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2014-10-31 0:51 [gomp4] acc enter/exit data Cesar Philippidis
2014-11-05 16:56 ` [gomp4] OpenACC update host/self maintenance (was: acc enter/exit data) Thomas Schwinge
2014-11-06 14:04 ` [gomp4] acc enter/exit data Thomas Schwinge
2014-12-10 9:54 ` Thomas Schwinge
2014-12-10 9:59 ` Jakub Jelinek
2014-12-17 22:18 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).