@@ -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)
@@ -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
@@ -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 },
@@ -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,
@@ -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):
@@ -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
@@ -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 ();
}
@@ -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;
}
@@ -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;
@@ -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)
@@ -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;
@@ -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
new file mode 100644
@@ -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)
+}
new file mode 100644
@@ -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)
+}
@@ -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.
@@ -242,4 +242,5 @@ GOACC_2.0 {
GOACC_data_start;
GOACC_kernels;
GOACC_parallel;
+ GOACC_update;
};
@@ -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_);
+}
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