This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[GOMP4, COMMITTED] OpenACC update directive.
- From: Thomas Schwinge <thomas at codesourcery dot com>
- To: <gcc-patches at gcc dot gnu dot org>
- Cc: <jnorris at codesourcery dot com>
- Date: Fri, 6 Jun 2014 07:54:41 +0200
- Subject: [GOMP4, COMMITTED] OpenACC update directive.
- Authentication-results: sourceware.org; auth=none
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/c-family/
* c-pragma.c (oacc_pragmas): Add "update".
* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_UPDATE.
(enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_HOST and
PRAGMA_OMP_CLAUSE_SELF.
gcc/c/
* c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_UPDATE.
(c_parser_omp_clause_name, c_parser_oacc_data_clause)
(c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_DEVICE,
PRAGMA_OMP_CLAUSE_HOST, PRAGMA_OMP_CLAUSE_SELF.
(c_parser_oacc_update): New function.
gcc/
* gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_UPDATE, and
extend mask size, GF_OMP_TARGET_KIND_MASK.
(is_gimple_omp_oacc_specifically): Handle
GF_OMP_TARGET_KIND_OACC_UPDATE.
* gimplify.c (gimplify_omp_target_update, gimplify_expr):
Likewise.
* gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
* omp-low.c (scan_omp_target, expand_omp_target)
(build_omp_regions_1, lower_omp_target, lower_omp_1)
(make_gimple_omp_edges): Likewise.
* oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin.
gcc/testsuite/
* c-c++-common/goacc/pragma_context.c: New file.
* c-c++-common/goacc/update-1.c: Likewise.
libgomp/
* libgomp.map (GOACC_2.0): Add GOACC_update.
* oacc-parallel.c (GOACC_update): New function.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211299 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 15 +++++
gcc/c-family/ChangeLog.gomp | 8 +++
gcc/c-family/c-pragma.c | 1 +
gcc/c-family/c-pragma.h | 3 +
gcc/c/ChangeLog.gomp | 9 +++
gcc/c/c-parser.c | 80 +++++++++++++++++++++++
gcc/gimple-pretty-print.c | 3 +
gcc/gimple.h | 5 +-
gcc/gimplify.c | 29 +++++---
gcc/oacc-builtins.def | 2 +
gcc/omp-low.c | 30 ++++++---
gcc/testsuite/ChangeLog.gomp | 5 ++
gcc/testsuite/c-c++-common/goacc/pragma_context.c | 32 +++++++++
gcc/testsuite/c-c++-common/goacc/update-1.c | 10 +++
libgomp/ChangeLog.gomp | 6 ++
libgomp/libgomp.map | 1 +
libgomp/oacc-parallel.c | 24 +++++++
17 files changed, 246 insertions(+), 17 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/pragma_context.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/update-1.c
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 88f09b3..be1aa16 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,18 @@
+2014-06-06 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_UPDATE, and
+ extend mask size, GF_OMP_TARGET_KIND_MASK.
+ (is_gimple_omp_oacc_specifically): Handle
+ GF_OMP_TARGET_KIND_OACC_UPDATE.
+ * gimplify.c (gimplify_omp_target_update, gimplify_expr):
+ Likewise.
+ * gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
+ * omp-low.c (scan_omp_target, expand_omp_target)
+ (build_omp_regions_1, lower_omp_target, lower_omp_1)
+ (make_gimple_omp_edges): Likewise.
+ * oacc-builtins.def (BUILT_IN_GOACC_UPDATE): New builtin.
+
2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
* gimplify.c (gimplify_scan_omp_clauses)
diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp
index 37ebfe9..0894675 100644
--- gcc/c-family/ChangeLog.gomp
+++ gcc/c-family/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2014-06-06 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * c-pragma.c (oacc_pragmas): Add "update".
+ * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_UPDATE.
+ (enum pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_HOST and
+ PRAGMA_OMP_CLAUSE_SELF.
+
2014-03-20 Thomas Schwinge <thomas@codesourcery.com>
* c-omp.c (check_omp_for_incr_expr, c_finish_omp_for): Update
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index 85da7ac..6a9d8e0 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1177,6 +1177,7 @@ static const struct omp_pragma_def oacc_pragmas[] = {
{ "kernels", PRAGMA_OACC_KERNELS },
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },
+ { "update", PRAGMA_OACC_UPDATE },
};
static const struct omp_pragma_def omp_pragmas[] = {
{ "atomic", PRAGMA_OMP_ATOMIC },
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index bb9c367..f3dfc63 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -31,6 +31,7 @@ typedef enum pragma_kind {
PRAGMA_OACC_KERNELS,
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
+ PRAGMA_OACC_UPDATE,
PRAGMA_OMP_ATOMIC,
PRAGMA_OMP_BARRIER,
PRAGMA_OMP_CANCEL,
@@ -88,6 +89,7 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
PRAGMA_OMP_CLAUSE_FOR,
PRAGMA_OMP_CLAUSE_FROM,
+ PRAGMA_OMP_CLAUSE_HOST,
PRAGMA_OMP_CLAUSE_IF,
PRAGMA_OMP_CLAUSE_INBRANCH,
PRAGMA_OMP_CLAUSE_LASTPRIVATE,
@@ -113,6 +115,7 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_SAFELEN,
PRAGMA_OMP_CLAUSE_SCHEDULE,
PRAGMA_OMP_CLAUSE_SECTIONS,
+ PRAGMA_OMP_CLAUSE_SELF,
PRAGMA_OMP_CLAUSE_SHARED,
PRAGMA_OMP_CLAUSE_SIMDLEN,
PRAGMA_OMP_CLAUSE_TASKGROUP,
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 1e80031..f1e45f3 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,12 @@
+2014-06-06 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_UPDATE.
+ (c_parser_omp_clause_name, c_parser_oacc_data_clause)
+ (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_DEVICE,
+ PRAGMA_OMP_CLAUSE_HOST, PRAGMA_OMP_CLAUSE_SELF.
+ (c_parser_oacc_update): New function.
+
2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
* c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index e20348e..bf4bad62 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -1207,6 +1207,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_update (c_parser *);
static void c_parser_omp_barrier (c_parser *);
static void c_parser_omp_flush (c_parser *);
static tree c_parser_omp_for_loop (location_t, c_parser *, enum tree_code,
@@ -4457,6 +4458,14 @@ c_parser_initval (c_parser *parser, struct c_expr *after,
Although they are erroneous if the labels declared aren't defined,
is it useful for the syntax to be this way?
+ OpenACC:
+
+ block-item:
+ openacc-directive
+
+ openacc-directive:
+ update-directive
+
OpenMP:
block-item:
@@ -9433,6 +9442,17 @@ c_parser_pragma (c_parser *parser, enum pragma_context context)
switch (id)
{
+ case PRAGMA_OACC_UPDATE:
+ if (context != pragma_compound)
+ {
+ if (context == pragma_stmt)
+ c_parser_error (parser, "%<#pragma acc update%> may only be "
+ "used in compound statements");
+ goto bad_stmt;
+ }
+ c_parser_oacc_update (parser);
+ return false;
+
case PRAGMA_OMP_BARRIER:
if (context != pragma_compound)
{
@@ -9680,6 +9700,10 @@ c_parser_omp_clause_name (c_parser *parser)
else if (!strcmp ("from", p))
result = PRAGMA_OMP_CLAUSE_FROM;
break;
+ case 'h':
+ if (!strcmp ("host", p))
+ result = PRAGMA_OMP_CLAUSE_SELF;
+ break;
case 'i':
if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
@@ -9755,6 +9779,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_SHARED;
else if (!strcmp ("simdlen", p))
result = PRAGMA_OMP_CLAUSE_SIMDLEN;
+ else if (!strcmp ("self", p))
+ result = PRAGMA_OMP_CLAUSE_SELF;
break;
case 't':
if (!strcmp ("taskgroup", p))
@@ -9960,6 +9986,12 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OMP_CLAUSE_DELETE:
kind = OMP_CLAUSE_MAP_FORCE_DEALLOC;
break;
+ case PRAGMA_OMP_CLAUSE_DEVICE:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_HOST:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
case PRAGMA_OMP_CLAUSE_PRESENT:
kind = OMP_CLAUSE_MAP_FORCE_PRESENT;
break;
@@ -9975,6 +10007,9 @@ 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;
}
tree nl, c;
nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
@@ -11248,10 +11283,18 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "delete";
break;
+ case PRAGMA_OMP_CLAUSE_DEVICE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "device";
+ break;
case PRAGMA_OMP_CLAUSE_DEVICEPTR:
clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
c_name = "deviceptr";
break;
+ case PRAGMA_OMP_CLAUSE_HOST:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "host";
+ break;
case PRAGMA_OMP_CLAUSE_NUM_GANGS:
clauses = c_parser_omp_clause_num_gangs (parser, clauses);
c_name = "num_gangs";
@@ -11280,6 +11323,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present_or_create";
break;
+ case PRAGMA_OMP_CLAUSE_SELF:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "self";
+ break;
case PRAGMA_OMP_CLAUSE_VECTOR_LENGTH:
clauses = c_parser_omp_clause_vector_length (parser, clauses);
c_name = "vector_length";
@@ -11721,6 +11768,39 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
return stmt;
}
+/* OpenACC 2.0:
+ # pragma acc update oacc-update-clause[optseq] new-line
+*/
+
+#define OACC_UPDATE_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HOST) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SELF) )
+
+static void
+c_parser_oacc_update (c_parser *parser)
+{
+ location_t loc = c_parser_peek_token (parser)->location;
+
+ c_parser_consume_pragma (parser);
+
+ tree clauses = c_parser_oacc_all_clauses (parser, OACC_UPDATE_CLAUSE_MASK,
+ "#pragma acc update");
+ if (find_omp_clause (clauses, OMP_CLAUSE_MAP) == NULL_TREE)
+ {
+ error_at (loc,
+ "%<#pragma acc update%> must contain at least one "
+ "%<device%> or %<host/self%> clause");
+ return;
+ }
+
+ tree stmt = make_node (OACC_UPDATE);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_UPDATE_CLAUSES (stmt) = clauses;
+ SET_EXPR_LOCATION (stmt, loc);
+ add_stmt (stmt);
+}
+
/* OpenMP 2.5:
# pragma omp atomic new-line
expression-stmt
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index 53cb138..ab06551 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1299,6 +1299,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_UPDATE:
+ kind = " oacc_update";
+ break;
default:
gcc_unreachable ();
}
diff --git gcc/gimple.h gcc/gimple.h
index 60b4896..1a1b952 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -101,11 +101,12 @@ enum gf_mask {
GF_OMP_FOR_KIND_CILKSIMD = GF_OMP_FOR_SIMD | 1,
GF_OMP_FOR_COMBINED = 1 << 3,
GF_OMP_FOR_COMBINED_INTO = 1 << 4,
- GF_OMP_TARGET_KIND_MASK = (1 << 2) - 1,
+ GF_OMP_TARGET_KIND_MASK = (1 << 3) - 1,
GF_OMP_TARGET_KIND_REGION = 0,
GF_OMP_TARGET_KIND_DATA = 1,
GF_OMP_TARGET_KIND_UPDATE = 2,
GF_OMP_TARGET_KIND_OACC_DATA = 3,
+ GF_OMP_TARGET_KIND_OACC_UPDATE = 4,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -5830,6 +5831,8 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
{
case GF_OMP_TARGET_KIND_OACC_DATA:
return true;
+ case GF_OMP_TARGET_KIND_OACC_UPDATE:
+ return true;
default:
return false;
}
diff --git gcc/gimplify.c gcc/gimplify.c
index a1b6be6..6b8e376 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7229,19 +7229,32 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
*expr_p = NULL_TREE;
}
-/* Gimplify the gross structure of OpenMP target update construct. */
+/* Gimplify the gross structure of OpenACC update and OpenMP target update
+ constructs. */
static void
gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
{
- tree expr = *expr_p;
+ tree expr = *expr_p, clauses;
+ int kind;
gimple stmt;
- gimplify_scan_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr), pre_p,
- ORT_WORKSHARE);
- gimplify_adjust_omp_clauses (&OMP_TARGET_UPDATE_CLAUSES (expr));
- stmt = gimple_build_omp_target (NULL, GF_OMP_TARGET_KIND_UPDATE,
- OMP_TARGET_UPDATE_CLAUSES (expr));
+ switch (TREE_CODE (expr))
+ {
+ case OACC_UPDATE:
+ clauses = OACC_UPDATE_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_OACC_UPDATE;
+ break;
+ case OMP_TARGET_UPDATE:
+ clauses = OMP_TARGET_UPDATE_CLAUSES (expr);
+ kind = GF_OMP_TARGET_KIND_UPDATE;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ gimplify_scan_omp_clauses (&clauses, pre_p, ORT_WORKSHARE);
+ gimplify_adjust_omp_clauses (&clauses);
+ stmt = gimple_build_omp_target (NULL, kind, clauses);
gimplify_seq_add_stmt (pre_p, stmt);
*expr_p = NULL_TREE;
@@ -8169,7 +8182,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
case OACC_HOST_DATA:
case OACC_DECLARE:
- case OACC_UPDATE:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
case OACC_WAIT:
@@ -8222,6 +8234,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_UPDATE:
case OMP_TARGET_UPDATE:
gimplify_omp_target_update (expr_p, pre_p);
ret = GS_ALL_DONE;
diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def
index 3935da4..dfb688c 100644
--- gcc/oacc-builtins.def
+++ gcc/oacc-builtins.def
@@ -37,3 +37,5 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_KERNELS, "GOACC_kernels",
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR_INT_INT_INT,
ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
diff --git gcc/omp-low.c gcc/omp-low.c
index 39f0598..6fd0f30 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -2419,7 +2419,8 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
tree name;
int kind = gimple_omp_target_kind (stmt);
- if (kind == GF_OMP_TARGET_KIND_OACC_DATA)
+ if (kind == GF_OMP_TARGET_KIND_OACC_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_UPDATE)
{
gcc_assert (taskreg_nesting_level == 0);
gcc_assert (target_nesting_level == 0);
@@ -8647,6 +8648,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_UPDATE:
+ start_ix = BUILT_IN_GOACC_UPDATE;
+ break;
default:
gcc_unreachable ();
}
@@ -8662,9 +8666,11 @@ expand_omp_target (struct omp_region *region)
cond = OMP_CLAUSE_IF_EXPR (c);
c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
- gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
if (c)
{
+ gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA
+ && kind != GF_OMP_TARGET_KIND_OACC_UPDATE);
+
device = OMP_CLAUSE_DEVICE_ID (c);
clause_loc = OMP_CLAUSE_LOCATION (c);
}
@@ -8920,7 +8926,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_UPDATE
+ || (gimple_omp_target_kind (stmt)
+ == GF_OMP_TARGET_KIND_OACC_UPDATE)))
new_omp_region (bb, code, parent);
else
{
@@ -10604,7 +10612,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_MAP_FORCE_PRESENT:
case OMP_CLAUSE_MAP_FORCE_DEALLOC:
case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
- gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA);
+ gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_UPDATE);
break;
default:
gcc_unreachable ();
@@ -10615,7 +10624,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
- gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
+ gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA
+ && kind != GF_OMP_TARGET_KIND_OACC_UPDATE);
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -10702,6 +10712,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_UPDATE:
tkind_type = short_unsigned_type_node;
talign_shift = 8;
break;
@@ -10904,7 +10915,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
else if (kind == GF_OMP_TARGET_KIND_DATA
|| kind == GF_OMP_TARGET_KIND_OACC_DATA)
new_body = tgt_body;
- if (kind != GF_OMP_TARGET_KIND_UPDATE)
+ if (kind != GF_OMP_TARGET_KIND_UPDATE
+ && kind != GF_OMP_TARGET_KIND_OACC_UPDATE)
{
gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));
gimple_omp_set_body (stmt, new_body);
@@ -11124,7 +11136,8 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GIMPLE_OMP_TARGET:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
- if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA)
+ if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA
+ || gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_UPDATE)
gcc_assert (!ctx->cancellable);
lower_omp_target (gsi_p, ctx);
break;
@@ -11584,7 +11597,8 @@ 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)
+ if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE
+ || gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_OACC_UPDATE)
cur_region = cur_region->outer;
break;
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 08ec907..17493ce 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-06-06 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc/pragma_context.c: New file.
+ * c-c++-common/goacc/update-1.c: Likewise.
+
2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
* c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
diff --git gcc/testsuite/c-c++-common/goacc/pragma_context.c gcc/testsuite/c-c++-common/goacc/pragma_context.c
new file mode 100644
index 0000000..ad33d92
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/pragma_context.c
@@ -0,0 +1,32 @@
+// pragma_external
+#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */
+
+// pragma_struct
+struct s_pragma_struct
+{
+#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */
+};
+
+// pragma_param
+void
+f_pragma_param (
+#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */
+ void)
+{
+}
+
+// pragma_stmt
+void
+f2 (void)
+{
+ if (0)
+#pragma acc update /* { dg-error "'#pragma acc update' may only be used in compound statements before '#pragma'" } */
+}
+
+// pragma_compound
+void
+f3 (void)
+{
+ int i = 0;
+#pragma acc update device(i)
+}
diff --git gcc/testsuite/c-c++-common/goacc/update-1.c gcc/testsuite/c-c++-common/goacc/update-1.c
new file mode 100644
index 0000000..970fdca
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/update-1.c
@@ -0,0 +1,10 @@
+void
+f (void)
+{
+#pragma acc update /* { dg-error "'#pragma acc update' must contain at least one 'device' or 'host/self' clause" } */
+
+ int i = 0;
+#pragma acc update device(i)
+#pragma acc update host(i)
+#pragma acc update self(i)
+}
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index d348672..8876a2e 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2014-06-06 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * libgomp.map (GOACC_2.0): Add GOACC_update.
+ * oacc-parallel.c (GOACC_update): New function.
+
2014-03-18 Ilya Verbin <ilya.verbin@intel.com>
* libgomp.map (GOMP_4.0.1): New symbol version.
diff --git libgomp/libgomp.map libgomp/libgomp.map
index f169f46..c575be3 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -242,4 +242,5 @@ GOACC_2.0 {
GOACC_data_start;
GOACC_kernels;
GOACC_parallel;
+ GOACC_update;
};
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index b1fd898..79b6254 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -104,3 +104,27 @@ GOACC_kernels (int device, void (*fn) (void *), const void *openmp_target,
GOACC_parallel (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds,
num_gangs, num_workers, vector_length);
}
+
+
+void
+GOACC_update (int device, const void *openmp_target, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+ unsigned char kinds_[mapnum];
+ size_t i;
+
+ /* TODO. Eventually, we'll be interpreting all mapping kinds according to
+ the OpenACC semantics; for now we're re-using what is implemented for
+ OpenMP. */
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i];
+ unsigned char align = kinds[i] >> 8;
+ if (kind > 4)
+ gomp_fatal ("memory mapping kind %x for %zd is not yet supported",
+ kind, i);
+
+ kinds_[i] = kind | align << 3;
+ }
+ GOMP_target_update (device, openmp_target, mapnum, hostaddrs, sizes, kinds_);
+}
--
1.9.1