* [gomp4] Initial support for OpenACC data clauses
@ 2014-01-14 15:09 Thomas Schwinge
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
` (3 more replies)
0 siblings, 4 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-01-14 15:09 UTC (permalink / raw)
To: jakub, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 532 bytes --]
Hi!
Here is a patch series that adds initial support for OpenACC data
clauses. It is not yet complete, but I thought I might as well already
now strive to get this integrated upstream instead of "hoarding" the
patches locally.
Would it be a good idea to also commit to trunk the (portions of the)
patches that don't directly relate with OpenACC stuff? That way, trunk
and gomp-4_0-branch would diverge a little less? Or, would you first
like to see all of this stabilitize on gomp-4_0-branch?
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread* [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET. 2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge @ 2014-01-14 15:10 ` thomas 2014-01-14 15:10 ` [gomp4 2/6] Prepare for extending omp_clause_map_kind thomas 2014-01-28 9:44 ` [gomp4] Initial support for OpenACC data clauses Thomas Schwinge ` (2 subsequent siblings) 3 siblings, 1 reply; 22+ messages in thread From: thomas @ 2014-01-14 15:10 UTC (permalink / raw) To: jakub, gcc-patches; +Cc: Thomas Schwinge From: Thomas Schwinge <thomas@codesourcery.com> gcc/ * gimplify.c (gimplify_call_expr, gimplify_modify_expr) (omp_firstprivatize_variable, omp_notice_threadprivate_variable) (omp_notice_variable, gimplify_adjust_omp_clauses) (gimplify_omp_workshare): Treat ORT_TARGET as a flag, not as a value. --- gcc/gimplify.c | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git gcc/gimplify.c gcc/gimplify.c index e45bed2..90507c2 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -2363,7 +2363,7 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) during omplower pass instead. */ struct gimplify_omp_ctx *ctx; for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context) - if (ctx->region_type == ORT_TARGET) + if (ctx->region_type & ORT_TARGET) break; if (ctx == NULL) fold_stmt (&gsi); @@ -4534,7 +4534,7 @@ gimplify_modify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, during omplower pass instead. */ struct gimplify_omp_ctx *ctx; for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context) - if (ctx->region_type == ORT_TARGET) + if (ctx->region_type & ORT_TARGET) break; if (ctx == NULL) fold_stmt (&gsi); @@ -5317,7 +5317,7 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl) else return; } - else if (ctx->region_type == ORT_TARGET) + else if (ctx->region_type & ORT_TARGET) omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY); else if (ctx->region_type != ORT_WORKSHARE && ctx->region_type != ORT_SIMD @@ -5499,7 +5499,7 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, struct gimplify_omp_ctx *octx; for (octx = ctx; octx; octx = octx->outer_context) - if (octx->region_type == ORT_TARGET) + if (octx->region_type & ORT_TARGET) { n = splay_tree_lookup (octx->variables, (splay_tree_key)decl); if (n == NULL) @@ -5560,7 +5560,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) } n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); - if (ctx->region_type == ORT_TARGET) + if (ctx->region_type & ORT_TARGET) { if (n == NULL) { @@ -6285,7 +6285,7 @@ gimplify_adjust_omp_clauses (tree *list_p) if (!DECL_P (decl)) break; n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); - if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN)) + if ((ctx->region_type & ORT_TARGET) && !(n->value & GOVD_SEEN)) remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST @@ -6857,7 +6857,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) gcc_unreachable (); } gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort); - if (ort == ORT_TARGET || ort == ORT_TARGET_DATA) + if ((ort & ORT_TARGET) || ort == ORT_TARGET_DATA) { push_gimplify_context (); gimple g = gimplify_and_return_first (OMP_BODY (expr), &body); -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 2/6] Prepare for extending omp_clause_map_kind. 2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas @ 2014-01-14 15:10 ` thomas 2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas 0 siblings, 1 reply; 22+ messages in thread From: thomas @ 2014-01-14 15:10 UTC (permalink / raw) To: jakub, gcc-patches; +Cc: Thomas Schwinge From: Thomas Schwinge <thomas@codesourcery.com> gcc/ * tree-core.h (omp_clause_map_kind): Make the identifiers' bit patterns more obvious. Add comments. * omp-low.c (lower_oacc_parallel, lower_omp_target): Test for omp_clause_map_kind flags set instead of for values. --- gcc/omp-low.c | 22 ++++++++++++++-------- gcc/tree-core.h | 16 +++++++++++----- 2 files changed, 25 insertions(+), 13 deletions(-) diff --git gcc/omp-low.c gcc/omp-low.c index eb755c3..899e970 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -8855,13 +8855,16 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree avar = create_tmp_var (TREE_TYPE (var), NULL); mark_addressable (avar); - if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC - && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM) + enum omp_clause_map_kind map_kind + = OMP_CLAUSE_MAP_KIND (c); + if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_TO)) + || map_kind == OMP_CLAUSE_MAP_POINTER) gimplify_assign (avar, var, &ilist); avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); - if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM - || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM) + if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_FROM)) && !TYPE_READONLY (TREE_TYPE (var))) { x = build_sender_ref (ovar, ctx); @@ -10331,13 +10334,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_assert (kind == GF_OMP_TARGET_KIND_REGION); tree avar = create_tmp_var (TREE_TYPE (var), NULL); mark_addressable (avar); - if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC - && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM) + enum omp_clause_map_kind map_kind + = OMP_CLAUSE_MAP_KIND (c); + if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_TO)) + || map_kind == OMP_CLAUSE_MAP_POINTER) gimplify_assign (avar, var, &ilist); avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); - if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM - || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM) + if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_FROM)) && !TYPE_READONLY (TREE_TYPE (var))) { x = build_sender_ref (ovar, ctx); diff --git gcc/tree-core.h gcc/tree-core.h index e2750e0..3602b5f 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -1112,14 +1112,20 @@ enum omp_clause_depend_kind enum omp_clause_map_kind { - OMP_CLAUSE_MAP_ALLOC, - OMP_CLAUSE_MAP_TO, - OMP_CLAUSE_MAP_FROM, - OMP_CLAUSE_MAP_TOFROM, + /* If not already present, allocate. */ + OMP_CLAUSE_MAP_ALLOC = 0, + /* ..., and copy to device. */ + OMP_CLAUSE_MAP_TO = 1 << 0, + /* ..., and copy from device. */ + OMP_CLAUSE_MAP_FROM = 1 << 1, + /* ..., and copy to and from device. */ + OMP_CLAUSE_MAP_TOFROM = OMP_CLAUSE_MAP_TO | OMP_CLAUSE_MAP_FROM, + /* Special map kinds. */ + OMP_CLAUSE_MAP_SPECIAL = 1 << 2, /* The following kind is an internal only map kind, used for pointer based array sections. OMP_CLAUSE_SIZE for these is not the pointer size, which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */ - OMP_CLAUSE_MAP_POINTER + OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL }; enum omp_clause_proc_bind_kind -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 3/6] Initial support for OpenACC memory mapping semantics. 2014-01-14 15:10 ` [gomp4 2/6] Prepare for extending omp_clause_map_kind thomas @ 2014-01-14 15:10 ` thomas 2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas ` (2 more replies) 0 siblings, 3 replies; 22+ messages in thread From: thomas @ 2014-01-14 15:10 UTC (permalink / raw) To: jakub, gcc-patches; +Cc: Thomas Schwinge From: Thomas Schwinge <thomas@codesourcery.com> gcc/ * tree-core.h (omp_clause_map_kind): Add OMP_CLAUSE_MAP_FORCE, OMP_CLAUSE_MAP_FORCE_ALLOC, OMP_CLAUSE_MAP_FORCE_TO, OMP_CLAUSE_MAP_FORCE_FROM, OMP_CLAUSE_MAP_FORCE_TOFROM, OMP_CLAUSE_MAP_FORCE_PRESENT, OMP_CLAUSE_MAP_FORCE_DEALLOC, and OMP_CLAUSE_MAP_FORCE_DEVICEPTR. * tree-pretty-print.c (dump_omp_clause): Handle these. * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_FORCE. (omp_region_type): Add ORT_TARGET_MAP_FORCE. (omp_add_variable, omp_notice_threadprivate_variable) (omp_notice_variable, gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses_1): Extend accordingly. (gimplify_oacc_parallel): Add ORT_TARGET_MAP_FORCE to ORT_TARGET usage. * omp-low.c (install_var_field, scan_sharing_clauses) (lower_oacc_parallel, lower_omp_target): Extend accordingly. --- gcc/gimplify.c | 92 ++++++++++++++++++++++++++++++++++++++++++------- gcc/omp-low.c | 33 +++++++++++------- gcc/tree-core.h | 19 +++++++++- gcc/tree-pretty-print.c | 21 +++++++++++ 4 files changed, 140 insertions(+), 25 deletions(-) diff --git gcc/gimplify.c gcc/gimplify.c index 90507c2..633784f 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -69,7 +69,13 @@ enum gimplify_omp_var_data GOVD_PRIVATE_OUTER_REF = 1024, GOVD_LINEAR = 2048, GOVD_ALIGNED = 4096, + + /* Flags for GOVD_MAP. */ + /* Don't copy back. */ GOVD_MAP_TO_ONLY = 8192, + /* Force a specific behavior (or else, a run-time error). */ + GOVD_MAP_FORCE = 16384, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -86,7 +92,11 @@ enum omp_region_type ORT_UNTIED_TASK = 5, ORT_TEAMS = 8, ORT_TARGET_DATA = 16, - ORT_TARGET = 32 + ORT_TARGET = 32, + + /* Flags for ORT_TARGET. */ + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */ + ORT_TARGET_MAP_FORCE = 64 }; /* Gimplify hashtable helper. */ @@ -5430,9 +5440,20 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) copy into or out of the context. */ if (!(flags & GOVD_LOCAL)) { - nflags = flags & GOVD_MAP - ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT - : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE; + if (flags & GOVD_MAP) + { + nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT; +#if 0 + /* Not sure if this is actually needed; haven't found a case + where this would change anything; TODO. */ + if (flags & GOVD_MAP_FORCE) + nflags |= OMP_CLAUSE_MAP_FORCE; +#endif + } + else if (flags & GOVD_PRIVATE) + nflags = GOVD_PRIVATE; + else + nflags = GOVD_FIRSTPRIVATE; nflags |= flags & GOVD_SEEN; t = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (t) == INDIRECT_REF); @@ -5501,6 +5522,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, for (octx = ctx; octx; octx = octx->outer_context) if (octx->region_type & ORT_TARGET) { + gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE)); + n = splay_tree_lookup (octx->variables, (splay_tree_key)decl); if (n == NULL) { @@ -5562,19 +5585,45 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); if (ctx->region_type & ORT_TARGET) { + unsigned map_force; + if (ctx->region_type & ORT_TARGET_MAP_FORCE) + map_force = GOVD_MAP_FORCE; + else + map_force = 0; if (n == NULL) { if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl))) { error ("%qD referenced in target region does not have " "a mappable type", decl); - omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags); + omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT | flags); } else - omp_add_variable (ctx, decl, GOVD_MAP | flags); + omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags); } else - n->value |= flags; + { +#if 0 + /* The following fails for: + + int l = 10; + float c[l]; + #pragma acc parallel copy(c[2:4]) + { + #pragma acc parallel + { + int t = sizeof c; + } + } + + ..., which we currently don't have to care about (nesting + disabled), but eventually will have to; TODO. */ + if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT)) + gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force); +#endif + + n->value |= flags; + } ret = lang_hooks.decls.omp_disregard_value_expr (decl, true); goto do_outer; } @@ -5858,6 +5907,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add; case OMP_CLAUSE_MAP: + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case OMP_CLAUSE_MAP_FORCE_PRESENT: + case OMP_CLAUSE_MAP_FORCE_DEALLOC: + case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: + input_location = OMP_CLAUSE_LOCATION (c); + /* TODO. */ + sorry ("data clause not yet implemented"); + remove = true; + break; + default: + break; + } if (OMP_CLAUSE_SIZE (c) && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) @@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; else if (code == OMP_CLAUSE_MAP) { - OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY - ? OMP_CLAUSE_MAP_TO - : OMP_CLAUSE_MAP_TOFROM; + unsigned map_kind; + map_kind = (flags & GOVD_MAP_TO_ONLY + ? OMP_CLAUSE_MAP_TO + : OMP_CLAUSE_MAP_TOFROM); + if (flags & GOVD_MAP_FORCE) + map_kind |= OMP_CLAUSE_MAP_FORCE; + OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind; + if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { @@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p) tree expr = *expr_p; gimple g; gimple_seq body = NULL; + enum omp_region_type ort = + (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE); - gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, - ORT_TARGET); + gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort); push_gimplify_context (); diff --git gcc/omp-low.c gcc/omp-low.c index 899e970..8c7df1b 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1064,6 +1064,8 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) || !splay_tree_lookup (ctx->field_map, (splay_tree_key) var)); gcc_assert ((mask & 2) == 0 || !ctx->sfield_map || !splay_tree_lookup (ctx->sfield_map, (splay_tree_key) var)); + gcc_assert ((mask & 3) == 3 + || gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); type = TREE_TYPE (var); if (mask & 4) @@ -1611,6 +1613,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); case OMP_CLAUSE_MAP: if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); @@ -1630,11 +1633,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) { - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in #pragma omp target data, there is nothing to map for those. */ - if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA + if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } @@ -8709,8 +8712,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_MAP: - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: var = OMP_CLAUSE_DECL (c); if (!DECL_P (var)) { @@ -8797,8 +8798,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_MAP: - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: nc = c; ovar = OMP_CLAUSE_DECL (c); if (!DECL_P (ovar)) @@ -8893,12 +8892,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_MAP: tkind = OMP_CLAUSE_MAP_KIND (c); break; - case OMP_CLAUSE_TO: - tkind = OMP_CLAUSE_MAP_TO; - break; - case OMP_CLAUSE_FROM: - tkind = OMP_CLAUSE_MAP_FROM; - break; default: gcc_unreachable (); } @@ -10179,6 +10172,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_MAP: +#ifdef ENABLE_CHECKING + /* First check what we're prepared to handle in the following. */ + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case OMP_CLAUSE_MAP_ALLOC: + case OMP_CLAUSE_MAP_TO: + case OMP_CLAUSE_MAP_FROM: + case OMP_CLAUSE_MAP_TOFROM: + case OMP_CLAUSE_MAP_POINTER: + break; + default: + gcc_unreachable (); + } +#endif + /* FALLTHRU */ + case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: var = OMP_CLAUSE_DECL (c); diff --git gcc/tree-core.h gcc/tree-core.h index 3602b5f..0aedea3 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -1125,7 +1125,24 @@ enum omp_clause_map_kind /* The following kind is an internal only map kind, used for pointer based array sections. OMP_CLAUSE_SIZE for these is not the pointer size, which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */ - OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL + OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL, + /* The following are only valid for OpenACC. */ + /* Flag to force a specific behavior (or else, a run-time error). */ + OMP_CLAUSE_MAP_FORCE = 1 << 3, + /* Allocate. */ + OMP_CLAUSE_MAP_FORCE_ALLOC = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_ALLOC, + /* ..., and copy to device. */ + OMP_CLAUSE_MAP_FORCE_TO = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TO, + /* ..., and copy from device. */ + OMP_CLAUSE_MAP_FORCE_FROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_FROM, + /* ..., and copy to and from device. */ + OMP_CLAUSE_MAP_FORCE_TOFROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TOFROM, + /* Must already be present. */ + OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL, + /* Deallocate a mapping, without copying from device. */ + OMP_CLAUSE_MAP_FORCE_DEALLOC, + /* Is a device pointer. */ + OMP_CLAUSE_MAP_FORCE_DEVICEPTR }; enum omp_clause_proc_bind_kind diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index 320c35b..f75f181 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -506,6 +506,27 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) case OMP_CLAUSE_MAP_TOFROM: pp_string (buffer, "tofrom"); break; + case OMP_CLAUSE_MAP_FORCE_ALLOC: + pp_string (buffer, "force_alloc"); + break; + case OMP_CLAUSE_MAP_FORCE_TO: + pp_string (buffer, "force_to"); + break; + case OMP_CLAUSE_MAP_FORCE_FROM: + pp_string (buffer, "force_from"); + break; + case OMP_CLAUSE_MAP_FORCE_TOFROM: + pp_string (buffer, "force_tofrom"); + break; + case OMP_CLAUSE_MAP_FORCE_PRESENT: + pp_string (buffer, "force_present"); + break; + case OMP_CLAUSE_MAP_FORCE_DEALLOC: + pp_string (buffer, "force_dealloc"); + break; + case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: + pp_string (buffer, "force_deviceptr"); + break; default: gcc_unreachable (); } -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing. 2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas @ 2014-01-14 15:10 ` thomas 2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas 2014-02-21 19:48 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge [not found] ` <538DF785.3050206@mentor.com> 2 siblings, 1 reply; 22+ messages in thread From: thomas @ 2014-01-14 15:10 UTC (permalink / raw) To: jakub, gcc-patches; +Cc: Thomas Schwinge From: Thomas Schwinge <thomas@codesourcery.com> gcc/c/ * c-parser.c (c_parser_oacc_all_clauses): New function. (c_parser_oacc_parallel): Use it. * c-typeck.c (c_finish_omp_clauses): Update comment. Remove duplicated variable initialization. --- gcc/c/c-parser.c | 59 +++++++++++++++++++++++++++++++++++++++++++++++++++----- gcc/c/c-typeck.c | 4 ++-- 2 files changed, 56 insertions(+), 7 deletions(-) diff --git gcc/c/c-parser.c gcc/c/c-parser.c index ce46f31..c8b80db 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -9465,7 +9465,7 @@ c_parser_pragma_pch_preprocess (c_parser *parser) c_common_pch_pragma (parse_in, TREE_STRING_POINTER (name)); } \f -/* OpenMP 2.5 / 3.0 / 3.1 / 4.0 parsing routines. */ +/* OpenACC and OpenMP parsing routines. */ /* Returns name of the next clause. If the clause is not recognized PRAGMA_OMP_CLAUSE_NONE is returned and @@ -10767,9 +10767,58 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list) return list; } +/* Parse all OpenACC clauses. The set clauses allowed by the directive + is a bitmask in MASK. Return the list of clauses found. */ + +static tree +c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, + const char *where, bool finish_p = true) +{ + tree clauses = NULL; + bool first = true; + + while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL)) + { + location_t here; + pragma_omp_clause c_kind; + const char *c_name; + tree prev = clauses; + + if (!first && c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + + here = c_parser_peek_token (parser)->location; + c_kind = c_parser_omp_clause_name (parser); + + switch (c_kind) + { + default: + c_parser_error (parser, "expected clause"); + goto saw_error; + } + + first = false; + + if (((mask >> c_kind) & 1) == 0 && !parser->error) + { + /* Remove the invalid clause(s) from the list to avoid + confusing the rest of the compiler. */ + clauses = prev; + error_at (here, "%qs is not valid for %qs", c_name, where); + } + } + + saw_error: + c_parser_skip_to_pragma_eol (parser); + + if (finish_p) + return c_finish_omp_clauses (clauses); + + return clauses; +} + /* Parse all OpenMP clauses. The set clauses allowed by the directive - is a bitmask in MASK. Return the list of clauses found; the result - of clause default goes in *pdefault. */ + is a bitmask in MASK. Return the list of clauses found. */ static tree c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, @@ -11019,8 +11068,8 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser) { tree stmt, clauses, block; - clauses = c_parser_omp_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, - "#pragma acc parallel"); + clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, + "#pragma acc parallel"); gcc_assert (clauses == NULL); block = c_begin_omp_parallel (); diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 854e149..81f0c5c 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -11661,7 +11661,7 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data) return NULL_TREE; } -/* For all elements of CLAUSES, validate them vs OpenMP constraints. +/* For all elements of CLAUSES, validate them against their constraints. Remove any elements from the list that are invalid. */ tree @@ -11669,7 +11669,7 @@ c_finish_omp_clauses (tree clauses) { bitmap_head generic_head, firstprivate_head, lastprivate_head; bitmap_head aligned_head; - tree c, t, *pc = &clauses; + tree c, t, *pc; bool branch_seen = false; bool copyprivate_seen = false; tree *nowait_clause = NULL; -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 5/6] Initial support in the C front end for OpenACC data clauses. 2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas @ 2014-01-14 15:10 ` thomas 2014-01-14 15:10 ` [gomp4 6/6] Enable initial " thomas 2014-02-12 11:17 ` [gomp4 5/6] Initial " Thomas Schwinge 0 siblings, 2 replies; 22+ messages in thread From: thomas @ 2014-01-14 15:10 UTC (permalink / raw) To: jakub, gcc-patches; +Cc: Thomas Schwinge From: Thomas Schwinge <thomas@codesourcery.com> gcc/c-family/ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Handle these. (c_parser_oacc_data_clause, c_parser_oacc_data_clause_deviceptr): New functions. (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. gcc/ * tree-core.h (omp_clause_code): Update description for OMP_CLAUSE_MAP. --- gcc/c-family/c-pragma.h | 12 +++- gcc/c/c-parser.c | 171 +++++++++++++++++++++++++++++++++++++++++++++++- gcc/tree-core.h | 6 +- 3 files changed, 184 insertions(+), 5 deletions(-) diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index 64eed11..2c8af67 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -63,18 +63,23 @@ typedef enum pragma_kind { } pragma_kind; -/* All clauses defined by OpenMP 2.5, 3.0, 3.1 and 4.0. +/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, and 4.0. Used internally by both C and C++ parsers. */ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_NONE = 0, PRAGMA_OMP_CLAUSE_ALIGNED, PRAGMA_OMP_CLAUSE_COLLAPSE, + PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYIN, + PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_COPYPRIVATE, + PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DEFAULT, + PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEPEND, PRAGMA_OMP_CLAUSE_DEVICE, + PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_DIST_SCHEDULE, PRAGMA_OMP_CLAUSE_FINAL, PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, @@ -92,6 +97,11 @@ typedef enum pragma_omp_clause { PRAGMA_OMP_CLAUSE_NUM_THREADS, PRAGMA_OMP_CLAUSE_ORDERED, PRAGMA_OMP_CLAUSE_PARALLEL, + PRAGMA_OMP_CLAUSE_PRESENT, + PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, + PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, + PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, + PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OMP_CLAUSE_PRIVATE, PRAGMA_OMP_CLAUSE_PROC_BIND, PRAGMA_OMP_CLAUSE_REDUCTION, diff --git gcc/c/c-parser.c gcc/c/c-parser.c index c8b80db..48c55e6 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -9496,16 +9496,26 @@ c_parser_omp_clause_name (c_parser *parser) case 'c': if (!strcmp ("collapse", p)) result = PRAGMA_OMP_CLAUSE_COLLAPSE; + else if (!strcmp ("copy", p)) + result = PRAGMA_OMP_CLAUSE_COPY; else if (!strcmp ("copyin", p)) result = PRAGMA_OMP_CLAUSE_COPYIN; + else if (!strcmp ("copyout", p)) + result = PRAGMA_OMP_CLAUSE_COPYOUT; else if (!strcmp ("copyprivate", p)) result = PRAGMA_OMP_CLAUSE_COPYPRIVATE; + else if (!strcmp ("create", p)) + result = PRAGMA_OMP_CLAUSE_CREATE; break; case 'd': - if (!strcmp ("depend", p)) + if (!strcmp ("delete", p)) + result = PRAGMA_OMP_CLAUSE_DELETE; + else if (!strcmp ("depend", p)) result = PRAGMA_OMP_CLAUSE_DEPEND; else if (!strcmp ("device", p)) result = PRAGMA_OMP_CLAUSE_DEVICE; + else if (!strcmp ("deviceptr", p)) + result = PRAGMA_OMP_CLAUSE_DEVICEPTR; else if (!strcmp ("dist_schedule", p)) result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE; break; @@ -9550,6 +9560,16 @@ c_parser_omp_clause_name (c_parser *parser) case 'p': if (!strcmp ("parallel", p)) result = PRAGMA_OMP_CLAUSE_PARALLEL; + else if (!strcmp ("present", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT; + else if (!strcmp ("present_or_copy", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY; + else if (!strcmp ("present_or_copyin", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN; + else if (!strcmp ("present_or_copyout", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT; + else if (!strcmp ("present_or_create", p)) + result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE; else if (!strcmp ("private", p)) result = PRAGMA_OMP_CLAUSE_PRIVATE; else if (!strcmp ("proc_bind", p)) @@ -9611,7 +9631,7 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code, } } -/* OpenMP 2.5: +/* OpenACC 2.0, OpenMP 2.5: variable-list: identifier variable-list , identifier @@ -9712,7 +9732,7 @@ c_parser_omp_variable_list (c_parser *parser, } /* Similarly, but expect leading and trailing parenthesis. This is a very - common case for omp clauses. */ + common case for OpenACC and OpenMP clauses. */ static tree c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, @@ -9729,6 +9749,107 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, return list; } +/* OpenACC 2.0: + copy ( variable-list ) + copyin ( variable-list ) + copyout ( variable-list ) + create ( variable-list ) + delete ( variable-list ) + present ( variable-list ) + present_or_copy ( variable-list ) + present_or_copyin ( variable-list ) + present_or_copyout ( variable-list ) + present_or_create ( variable-list ) */ + +static tree +c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, + tree list) +{ + enum omp_clause_map_kind kind; + switch (c_kind) + { + default: + gcc_unreachable (); + case PRAGMA_OMP_CLAUSE_COPY: + kind = OMP_CLAUSE_MAP_FORCE_TOFROM; + break; + case PRAGMA_OMP_CLAUSE_COPYIN: + kind = OMP_CLAUSE_MAP_FORCE_TO; + break; + case PRAGMA_OMP_CLAUSE_COPYOUT: + kind = OMP_CLAUSE_MAP_FORCE_FROM; + break; + case PRAGMA_OMP_CLAUSE_CREATE: + kind = OMP_CLAUSE_MAP_FORCE_ALLOC; + break; + case PRAGMA_OMP_CLAUSE_DELETE: + kind = OMP_CLAUSE_MAP_FORCE_DEALLOC; + break; + case PRAGMA_OMP_CLAUSE_PRESENT: + kind = OMP_CLAUSE_MAP_FORCE_PRESENT; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY: + kind = OMP_CLAUSE_MAP_TOFROM; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN: + kind = OMP_CLAUSE_MAP_TO; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT: + kind = OMP_CLAUSE_MAP_FROM; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE: + kind = OMP_CLAUSE_MAP_ALLOC; + break; + } + tree nl, c; + nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list); + + for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_MAP_KIND (c) = kind; + + return nl; +} + +/* OpenACC 2.0: + deviceptr ( variable-list ) */ + +static tree +c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list) +{ + location_t loc = c_parser_peek_token (parser)->location; + tree vars, t; + + /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic + c_parser_oacc_data_clause), as for PRAGMA_OMP_CLAUSE_DEVICEPTR, + variable-list must only allow for pointer variables. */ + vars = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_ERROR, NULL); + for (t = vars; t && t; t = TREE_CHAIN (t)) + { + tree v = TREE_PURPOSE (t); + + /* FIXME diagnostics: Ideally we should keep individual + locations for all the variables in the var list to make the + following errors more precise. Perhaps + c_parser_omp_var_list_parens() should construct a list of + locations to go along with the var list. */ + + if (TREE_CODE (v) != VAR_DECL) + error_at (loc, "%qD is not a variable", v); + else if (TREE_TYPE (v) == error_mark_node) + ; + else if (!POINTER_TYPE_P (TREE_TYPE (v))) + error_at (loc, "%qD is not a pointer variable", v); + + tree u = build_omp_clause (loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_MAP_KIND (u) = OMP_CLAUSE_MAP_FORCE_DEVICEPTR; + OMP_CLAUSE_DECL (u) = v; + OMP_CLAUSE_CHAIN (u) = list; + list = u; + } + + return list; +} + /* OpenMP 3.0: collapse ( constant-expression ) */ @@ -10792,6 +10913,50 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, switch (c_kind) { + case PRAGMA_OMP_CLAUSE_COPY: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copy"; + break; + case PRAGMA_OMP_CLAUSE_COPYIN: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copyin"; + break; + case PRAGMA_OMP_CLAUSE_COPYOUT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "copyout"; + break; + case PRAGMA_OMP_CLAUSE_CREATE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "create"; + break; + case PRAGMA_OMP_CLAUSE_DELETE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "delete"; + break; + case PRAGMA_OMP_CLAUSE_DEVICEPTR: + clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses); + c_name = "deviceptr"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copy"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copyin"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_copyout"; + break; + case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE: + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); + c_name = "present_or_create"; + break; default: c_parser_error (parser, "expected clause"); goto saw_error; diff --git gcc/tree-core.h gcc/tree-core.h index 0aedea3..bfe4943 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -258,7 +258,11 @@ enum omp_clause_code { /* OpenMP clause: to (variable-list). */ OMP_CLAUSE_TO, - /* OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ + /* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr, + present, present_or_copy, present_or_copyin, present_or_copyout, + present_or_create} (variable-list). + + OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, /* Internal clause: temporary for combined loops expansion. */ -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 6/6] Enable initial support in the C front end for OpenACC data clauses. 2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas @ 2014-01-14 15:10 ` thomas 2014-02-12 11:17 ` [gomp4 5/6] Initial " Thomas Schwinge 1 sibling, 0 replies; 22+ messages in thread From: thomas @ 2014-01-14 15:10 UTC (permalink / raw) To: jakub, gcc-patches; +Cc: Thomas Schwinge From: Thomas Schwinge <thomas@codesourcery.com> gcc/c/ * c-parser.c (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: New file. * c-c++-common/goacc/deviceptr-1.c: New file. libgomp/ * testsuite/libgomp.oacc-c/parallel-1.c: Extend. --- gcc/c/c-parser.c | 14 +- .../c-c++-common/goacc/data-clause-duplicate-1.c | 13 ++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 64 +++++++++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c | 150 +++++++++++++++++++-- 4 files changed, 228 insertions(+), 13 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-1.c diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 48c55e6..d6a2af0 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -11225,8 +11225,17 @@ c_parser_omp_structured_block (c_parser *parser) LOC is the location of the #pragma token. */ -#define OACC_PARALLEL_CLAUSE_MASK \ - (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE) +#define OACC_PARALLEL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) ) static tree c_parser_oacc_parallel (location_t loc, c_parser *parser) @@ -11235,7 +11244,6 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser) clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK, "#pragma acc parallel"); - gcc_assert (clauses == NULL); block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser)); diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c new file mode 100644 index 0000000..1bcf5be --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -0,0 +1,13 @@ +void +fun (void) +{ + float *fp; +#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel create(fp[:10]) deviceptr(fp) + /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */ + /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */ + ; +} diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c new file mode 100644 index 0000000..0f0cf0c --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -0,0 +1,64 @@ +void +fun1 (void) +{ +#pragma acc parallel deviceptr(u) /* { dg-error "'u' undeclared" } */ + ; +#pragma acc parallel deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + ; + +#pragma acc parallel deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */ + ; +#pragma acc parallel deviceptr(fun1[2:5]) + /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 11 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 11 } */ + ; + + int i; +#pragma acc parallel deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */ + ; +#pragma acc parallel deviceptr(i[0:4]) + /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 19 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 19 } */ + ; + + float fa[10]; +#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */ + ; +#pragma acc parallel deviceptr(fa[1:5]) + /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 27 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 27 } */ + ; + + float *fp; +#pragma acc parallel deviceptr(fp) + ; +#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */ + ; +} + +void +fun2 (void) +{ + int i; + float *fp; +#pragma acc parallel deviceptr(fp,u,fun2,i,fp) + /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 44 } */ + /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 44 } */ + /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 44 } */ + /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 44 } */ + ; +} + +void +fun3 (void) +{ + float *fp; +#pragma acc parallel deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +#pragma acc parallel copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; +} + +/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */ diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c index b40545d..ff54b9d 100644 --- libgomp/testsuite/libgomp.oacc-c/parallel-1.c +++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c @@ -2,25 +2,155 @@ extern void abort (); -volatile int i; +int i; int main(void) { - volatile int j; + int j, v; - i = -0x42; - j = -42; -#pragma acc parallel +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) copyin (i, j) { - if (i != -0x42 || j != -42) + if (i != -1 || j != -2) abort (); - i = 42; - j = 0x42; - if (i != 42 || j != 0x42) + i = 2; + j = 1; + if (i != 2 || j != 1) abort (); + v = 1; } - if (i != 42 || j != 0x42) + if (v != 1 || i != -1 || j != -2) abort (); + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) copyout (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) copy (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) create (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != -1 || j != -2) + abort (); +#endif + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyin (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != -1 || j != -2) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyout (i, j) + { + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copy (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); + +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) present (i, j) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); +#endif + +#if 0 + i = -1; + j = -2; + v = 0; +#pragma acc parallel /* copyout */ present_or_copyout (v) + { + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + v = 1; + } + if (v != 1 || i != 2 || j != 1) + abort (); +#endif + return 0; } -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 5/6] Initial support in the C front end for OpenACC data clauses. 2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas 2014-01-14 15:10 ` [gomp4 6/6] Enable initial " thomas @ 2014-02-12 11:17 ` Thomas Schwinge 1 sibling, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-02-12 11:17 UTC (permalink / raw) To: gcc-patches, Ilmir Usmanov; +Cc: jakub [-- Attachment #1: Type: text/plain, Size: 6906 bytes --] Hi! On Tue, 14 Jan 2014 16:10:07 +0100, I wrote: > gcc/c-family/ > * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY, > PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE, > PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEVICEPTR, > PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and > PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. > gcc/c/ > * c-parser.c (c_parser_omp_clause_name): Handle these. > (c_parser_oacc_data_clause, c_parser_oacc_data_clause_deviceptr): > New functions. > (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_COPY, > PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT, > PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DELETE, > PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN, > PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and > PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE. > gcc/ > * tree-core.h (omp_clause_code): Update description for > OMP_CLAUSE_MAP. This I committed to gomp-4_0-branch as r207177. In <http://news.gmane.org/find-root.php?message_id=%3C52E68AE3.9030706%40samsung.com%3E>, Ilmir mentioned that I'm missing to handle the »short names: pcopy, pcopyin, pcopyout and pcreate (see 2.6.5.9 - 2.6.5.12 of OpenACC 2.0«. Unless there are any comments, I'll soon commit the following to gomp-4_0-branch: commit 9a1f6c075f6198c9ae3281387b875e6012e4387e Author: Thomas Schwinge <thomas@codesourcery.com> Date: Wed Feb 12 11:59:51 2014 +0100 OpenACC: pcopy, pcopyin, pcopyout, pcreate clauses. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin, pcopyout, pcreate clauses. (c_parser_oacc_data_clause): Update comment. gcc/ * tree-core.h (omp_clause_code) <map>: Mention pcopy, pcopyin, pcopyout, pcreate OpenACC clauses. gcc/testsuite/ * c-c++-common/goacc/pcopy.c: New file. * c-c++-common/goacc/pcopyin.c: Likewise. * c-c++-common/goacc/pcopyout.c: Likewise. * c-c++-common/goacc/pcreate.c: Likewise. diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 6e89471..f401cef 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -9671,13 +9671,17 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_PARALLEL; else if (!strcmp ("present", p)) result = PRAGMA_OMP_CLAUSE_PRESENT; - else if (!strcmp ("present_or_copy", p)) + else if (!strcmp ("present_or_copy", p) + || !strcmp ("pcopy", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY; - else if (!strcmp ("present_or_copyin", p)) + else if (!strcmp ("present_or_copyin", p) + || !strcmp ("pcopyin", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN; - else if (!strcmp ("present_or_copyout", p)) + else if (!strcmp ("present_or_copyout", p) + || !strcmp ("pcopyout", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT; - else if (!strcmp ("present_or_create", p)) + else if (!strcmp ("present_or_create", p) + || !strcmp ("pcreate", p)) result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE; else if (!strcmp ("private", p)) result = PRAGMA_OMP_CLAUSE_PRIVATE; @@ -9870,9 +9874,13 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind, delete ( variable-list ) present ( variable-list ) present_or_copy ( variable-list ) + pcopy ( variable-list ) present_or_copyin ( variable-list ) + pcopyin ( variable-list ) present_or_copyout ( variable-list ) - present_or_create ( variable-list ) */ + pcopyout ( variable-list ) + present_or_create ( variable-list ) + pcreate ( variable-list ) */ static tree c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, diff --git gcc/testsuite/c-c++-common/goacc/pcopy.c gcc/testsuite/c-c++-common/goacc/pcopy.c new file mode 100644 index 0000000..fd16525 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcopy.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcopy(cp[3:5]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/testsuite/c-c++-common/goacc/pcopyin.c gcc/testsuite/c-c++-common/goacc/pcopyin.c new file mode 100644 index 0000000..c009d24 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcopyin.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcopyin(cp[4:6]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/testsuite/c-c++-common/goacc/pcopyout.c gcc/testsuite/c-c++-common/goacc/pcopyout.c new file mode 100644 index 0000000..6099eff --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcopyout.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcopyout(cp[5:7]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/testsuite/c-c++-common/goacc/pcreate.c gcc/testsuite/c-c++-common/goacc/pcreate.c new file mode 100644 index 0000000..2f6e836 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/pcreate.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel pcreate(cp[6:8]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ diff --git gcc/tree-core.h gcc/tree-core.h index a5a95cd..2d9bf0c 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -259,8 +259,9 @@ enum omp_clause_code { OMP_CLAUSE_TO, /* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr, - present, present_or_copy, present_or_copyin, present_or_copyout, - present_or_create} (variable-list). + 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). */ OMP_CLAUSE_MAP, Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 3/6] Initial support for OpenACC memory mapping semantics. 2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas 2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas @ 2014-02-21 19:48 ` Thomas Schwinge 2014-02-21 20:32 ` [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE Thomas Schwinge [not found] ` <538DF785.3050206@mentor.com> 2 siblings, 1 reply; 22+ messages in thread From: Thomas Schwinge @ 2014-02-21 19:48 UTC (permalink / raw) To: gcc-patches; +Cc: jakub [-- Attachment #1: Type: text/plain, Size: 8797 bytes --] Hi! On Tue, 14 Jan 2014 16:10:05 +0100, I wrote: > --- gcc/gimplify.c > +++ gcc/gimplify.c > @@ -86,7 +92,11 @@ enum omp_region_type > ORT_UNTIED_TASK = 5, > ORT_TEAMS = 8, > ORT_TARGET_DATA = 16, > - ORT_TARGET = 32 > + ORT_TARGET = 32, > + > + /* Flags for ORT_TARGET. */ > + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */ > + ORT_TARGET_MAP_FORCE = 64 > }; Continuing on that route, I have now applied the following to gomp-4_0-branch in r208014: commit dee2965ae547af0bc90d618e7fa40fbf2f5292b4 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Fri Feb 21 19:45:12 2014 +0000 Gimplification: New flag ORT_TARGET_OFFLOAD replaces !ORT_TARGET_DATA. gcc/ * gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA. Update all users. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208014 138bc75d-0d04-0410-961f-82ee72b054a4 diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 1ce952d..bf8ec96 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,9 @@ 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + * gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a + flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA. + Update all users. + * omp-low.c (gimple_code_is_oacc): Move to... * gimple.h (is_gimple_omp_oacc_specifically): ... here. Update users, and also use it in more places where currently we've only diff --git gcc/gimplify.c gcc/gimplify.c index 51a1b73..9aa9301c 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -100,10 +100,11 @@ enum omp_region_type ORT_TASK = 4, ORT_UNTIED_TASK = 5, ORT_TEAMS = 8, - ORT_TARGET_DATA = 16, - ORT_TARGET = 32, + ORT_TARGET = 16, /* Flags for ORT_TARGET. */ + /* Prepare this region for offloading. */ + ORT_TARGET_OFFLOAD = 32, /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */ ORT_TARGET_MAP_FORCE = 64 }; @@ -2202,7 +2203,7 @@ gimplify_arg (tree *arg_p, gimple_seq *pre_p, location_t call_location) return gimplify_expr (arg_p, pre_p, NULL, test, fb); } -/* Don't fold STMT inside ORT_TARGET, because it can break code by adding decl +/* Don't fold inside offloading regsion: it can break code by adding decl references that weren't in the source. We'll do it during omplower pass instead. */ @@ -2211,7 +2212,8 @@ maybe_fold_stmt (gimple_stmt_iterator *gsi) { struct gimplify_omp_ctx *ctx; for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context) - if (ctx->region_type & ORT_TARGET) + if (ctx->region_type & ORT_TARGET + && ctx->region_type & ORT_TARGET_OFFLOAD) return false; return fold_stmt (gsi); } @@ -5388,10 +5390,12 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl) return; } else if (ctx->region_type & ORT_TARGET) - omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY); + { + if (ctx->region_type & ORT_TARGET_OFFLOAD) + omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY); + } else if (ctx->region_type != ORT_WORKSHARE - && ctx->region_type != ORT_SIMD - && ctx->region_type != ORT_TARGET_DATA) + && ctx->region_type != ORT_SIMD) omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE); ctx = ctx->outer_context; @@ -5580,7 +5584,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, struct gimplify_omp_ctx *octx; for (octx = ctx; octx; octx = octx->outer_context) - if (octx->region_type & ORT_TARGET) + if ((octx->region_type & ORT_TARGET) + && (octx->region_type & ORT_TARGET_OFFLOAD)) { gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE)); @@ -5643,7 +5648,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) } n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); - if (ctx->region_type & ORT_TARGET) + if ((ctx->region_type & ORT_TARGET) + && (ctx->region_type & ORT_TARGET_OFFLOAD)) { unsigned map_force; if (ctx->region_type & ORT_TARGET_MAP_FORCE) @@ -5695,7 +5701,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) if (ctx->region_type == ORT_WORKSHARE || ctx->region_type == ORT_SIMD - || ctx->region_type == ORT_TARGET_DATA) + || ((ctx->region_type & ORT_TARGET) + && !(ctx->region_type & ORT_TARGET_OFFLOAD))) goto do_outer; /* ??? Some compiler-generated variables (like SAVE_EXPRs) could be @@ -5746,7 +5753,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) { splay_tree_node n2; - if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0) + if (octx->region_type & ORT_TARGET) continue; n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl); if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED) @@ -5899,7 +5906,7 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate) || (!copyprivate && lang_hooks.decls.omp_privatize_by_reference (decl))); - if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0) + if (ctx->region_type & ORT_TARGET) continue; n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); @@ -6456,7 +6463,9 @@ gimplify_adjust_omp_clauses (tree *list_p) if (!DECL_P (decl)) break; n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); - if ((ctx->region_type & ORT_TARGET) && !(n->value & GOVD_SEEN)) + if ((ctx->region_type & ORT_TARGET) + && (ctx->region_type & ORT_TARGET_OFFLOAD) + && !(n->value & GOVD_SEEN)) remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST @@ -6574,8 +6583,9 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p) tree expr = *expr_p; gimple g; gimple_seq body = NULL; - enum omp_region_type ort = - (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE); + enum omp_region_type ort = (enum omp_region_type) (ORT_TARGET + | ORT_TARGET_OFFLOAD + | ORT_TARGET_MAP_FORCE); gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort); @@ -7031,11 +7041,11 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) case OMP_SINGLE: break; case OMP_TARGET: + ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD); + break; + case OMP_TARGET_DATA: ort = ORT_TARGET; break; - case OMP_TARGET_DATA: - ort = ORT_TARGET_DATA; - break; case OMP_TEAMS: ort = ORT_TEAMS; break; @@ -7043,7 +7053,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) gcc_unreachable (); } gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort); - if ((ort & ORT_TARGET) || ort == ORT_TARGET_DATA) + if (ort & ORT_TARGET) { push_gimplify_context (); gimple g = gimplify_and_return_first (OMP_BODY (expr), &body); @@ -7051,7 +7061,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) pop_gimplify_context (g); else pop_gimplify_context (NULL); - if (ort == ORT_TARGET_DATA) + if (!(ort & ORT_TARGET_OFFLOAD)) { gimple_seq cleanup = NULL; tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA); @@ -8697,7 +8707,9 @@ gimplify_body (tree fndecl, bool do_parms) { gcc_assert (gimplify_omp_ctxp == NULL); if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl))) - gimplify_omp_ctxp = new_omp_context (ORT_TARGET); + gimplify_omp_ctxp + = new_omp_context ((enum omp_region_type) (ORT_TARGET + | ORT_TARGET_OFFLOAD)); } /* Unshare most shared trees in the body and in that of any nested functions. diff --git gcc/omp-low.c gcc/omp-low.c index b975dad..9fef4c1 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -10858,8 +10858,8 @@ lower_omp (gimple_seq *body, omp_context *ctx) gimple_stmt_iterator gsi; for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi)) lower_omp_1 (&gsi, ctx); - /* During gimplification, we have not always invoked fold_stmt - (gimplify.c:maybe_fold_stmt); call it now. */ + /* During gimplification, we haven't folded statments inside offloading + regions (gimplify.c:maybe_fold_stmt); do that now. */ if (target_nesting_level) for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi)) fold_stmt (&gsi); Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE. 2014-02-21 19:48 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge @ 2014-02-21 20:32 ` Thomas Schwinge 2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge 0 siblings, 1 reply; 22+ messages in thread From: Thomas Schwinge @ 2014-02-21 20:32 UTC (permalink / raw) To: gcc-patches; +Cc: jakub From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/ * omp-low.c (scan_sharing_clauses): Catch unexpected occurrences of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208015 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 3 +++ gcc/omp-low.c | 25 +++++++++++++++++++++++++ 2 files changed, 28 insertions(+) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index bf8ec96..bd46f2e 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,8 @@ 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + * omp-low.c (scan_sharing_clauses): Catch unexpected occurrences + of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP. + * gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA. Update all users. diff --git gcc/omp-low.c gcc/omp-low.c index 9fef4c1..bca4599 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1630,6 +1630,26 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FROM: gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt)); case OMP_CLAUSE_MAP: + switch (OMP_CLAUSE_CODE (c)) + { + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: + /* The to and from clauses are only ever seen with OpenMP target + update constructs. */ + gcc_assert (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET + && (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_UPDATE)); + break; + case OMP_CLAUSE_MAP: + /* The map clause is never seen with OpenMP target update + constructs. */ + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET + || (gimple_omp_target_kind (ctx->stmt) + != GF_OMP_TARGET_KIND_UPDATE)); + break; + default: + gcc_unreachable (); + } if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); decl = OMP_CLAUSE_DECL (c); @@ -1799,6 +1819,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_MAP: + /* The map clause is never seen with OpenMP target update + constructs. */ + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET + || (gimple_omp_target_kind (ctx->stmt) + != GF_OMP_TARGET_KIND_UPDATE)); if (!gimple_code_is_oacc (ctx->stmt) && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) break; -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA. 2014-02-21 20:32 ` [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE Thomas Schwinge @ 2014-02-21 20:32 ` Thomas Schwinge 2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge ` (2 more replies) 0 siblings, 3 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-02-21 20:32 UTC (permalink / raw) To: gcc-patches; +Cc: jakub From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA. (is_gimple_omp_oacc_specifically): Handle it. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise. * omp-low.c (scan_sharing_clauses, scan_omp_target) (expand_omp_target, lower_omp_target, lower_omp_1): Likewise. * gimple.def (GIMPLE_OMP_TARGET): Update comment. * gimple.c (gimple_build_omp_target): Likewise. (gimple_copy): Catch unimplemented case. * tree-inline.c (remap_gimple_stmt): Likewise. * tree-nested.c (convert_nonlocal_reference_stmt) (convert_local_reference_stmt, convert_gimple_call): Likewise. * oacc-builtins.def (BUILT_IN_GOACC_DATA_START) (BUILT_IN_GOACC_DATA_END): New builtins. libgomp/ * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New functions. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208016 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 15 ++++++ gcc/gimple-pretty-print.c | 3 ++ gcc/gimple.c | 4 +- gcc/gimple.def | 1 + gcc/gimple.h | 9 ++++ gcc/gimplify.c | 33 +++++++++--- gcc/oacc-builtins.def | 6 ++- gcc/omp-low.c | 132 ++++++++++++++++++++++++++++++++++++---------- gcc/tree-inline.c | 1 + gcc/tree-nested.c | 3 ++ libgomp/ChangeLog.gomp | 7 +++ libgomp/libgomp.map | 2 + libgomp/libgomp_g.h | 3 ++ libgomp/oacc-parallel.c | 34 +++++++++++- 14 files changed, 213 insertions(+), 40 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index bd46f2e..824ec94 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,20 @@ 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA. + (is_gimple_omp_oacc_specifically): Handle it. + * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. + * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise. + * omp-low.c (scan_sharing_clauses, scan_omp_target) + (expand_omp_target, lower_omp_target, lower_omp_1): Likewise. + * gimple.def (GIMPLE_OMP_TARGET): Update comment. + * gimple.c (gimple_build_omp_target): Likewise. + (gimple_copy): Catch unimplemented case. + * tree-inline.c (remap_gimple_stmt): Likewise. + * tree-nested.c (convert_nonlocal_reference_stmt) + (convert_local_reference_stmt, convert_gimple_call): Likewise. + * oacc-builtins.def (BUILT_IN_GOACC_DATA_START) + (BUILT_IN_GOACC_DATA_END): New builtins. + * omp-low.c (scan_sharing_clauses): Catch unexpected occurrences of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP. diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index 91a3eb2..ad9369c 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1289,6 +1289,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) case GF_OMP_TARGET_KIND_UPDATE: kind = " update"; break; + case GF_OMP_TARGET_KIND_OACC_DATA: + kind = " oacc_data"; + break; default: gcc_unreachable (); } diff --git gcc/gimple.c gcc/gimple.c index 2a967aa..30561b1 100644 --- gcc/gimple.c +++ gcc/gimple.c @@ -1051,7 +1051,8 @@ gimple_build_omp_single (gimple_seq body, tree clauses) /* Build a GIMPLE_OMP_TARGET statement. BODY is the sequence of statements that will be executed. - CLAUSES are any of the OMP target construct's clauses. */ + KIND is the kind of target region. + CLAUSES are any of the construct's clauses. */ gimple gimple_build_omp_target (gimple_seq body, int kind, tree clauses) @@ -1747,6 +1748,7 @@ gimple_copy (gimple stmt) case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: copy_omp_body: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); new_seq = gimple_seq_copy (gimple_omp_body (stmt)); gimple_omp_set_body (copy, new_seq); break; diff --git gcc/gimple.def gcc/gimple.def index 2b78c06..ce800bd 100644 --- gcc/gimple.def +++ gcc/gimple.def @@ -360,6 +360,7 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE) DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT) /* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents + #pragma acc data #pragma omp target {,data,update} BODY is the sequence of statements inside the target construct (NULL for target update). diff --git gcc/gimple.h gcc/gimple.h index 0d250ef..b4ee9fa 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -102,6 +102,7 @@ enum gf_mask { GF_OMP_TARGET_KIND_REGION = 0 << 0, GF_OMP_TARGET_KIND_DATA = 1 << 0, GF_OMP_TARGET_KIND_UPDATE = 2 << 0, + GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -5684,6 +5685,14 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) { case GIMPLE_OACC_PARALLEL: return true; + case GIMPLE_OMP_TARGET: + switch (gimple_omp_target_kind (stmt)) + { + case GF_OMP_TARGET_KIND_OACC_DATA: + return true; + default: + return false; + } default: return false; } diff --git gcc/gimplify.c gcc/gimplify.c index 9aa9301c..fd4305c 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7023,9 +7023,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) return GS_ALL_DONE; } -/* Gimplify the gross structure of other OpenMP constructs. - In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA - and OMP_TEAMS. */ +/* Gimplify the gross structure of several OpenACC or OpenMP constructs. */ static void gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) @@ -7033,12 +7031,17 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) tree expr = *expr_p; gimple stmt; gimple_seq body = NULL; - enum omp_region_type ort = ORT_WORKSHARE; + enum omp_region_type ort; switch (TREE_CODE (expr)) { + case OACC_DATA: + ort = (enum omp_region_type) (ORT_TARGET + | ORT_TARGET_MAP_FORCE); + break; case OMP_SECTIONS: case OMP_SINGLE: + ort = ORT_WORKSHARE; break; case OMP_TARGET: ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD); @@ -7063,9 +7066,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) pop_gimplify_context (NULL); if (!(ort & ORT_TARGET_OFFLOAD)) { - gimple_seq cleanup = NULL; - tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA); + enum built_in_function end_ix; + switch (TREE_CODE (expr)) + { + case OACC_DATA: + end_ix = BUILT_IN_GOACC_DATA_END; + break; + case OMP_TARGET_DATA: + end_ix = BUILT_IN_GOMP_TARGET_END_DATA; + break; + default: + gcc_unreachable (); + } + tree fn = builtin_decl_explicit (end_ix); g = gimple_build_call (fn, 0); + gimple_seq cleanup = NULL; gimple_seq_add_stmt (&cleanup, g); g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY); body = NULL; @@ -7078,6 +7093,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) switch (TREE_CODE (expr)) { + case OACC_DATA: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA, + OACC_DATA_CLAUSES (expr)); + break; case OMP_SECTIONS: stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); break; @@ -8047,7 +8066,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, break; case OACC_KERNELS: - case OACC_DATA: case OACC_HOST_DATA: case OACC_DECLARE: case OACC_UPDATE: @@ -8076,6 +8094,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = gimplify_omp_for (expr_p, pre_p); break; + case OACC_DATA: case OMP_SECTIONS: case OMP_SINGLE: case OMP_TARGET: diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def index a75e42d..eaf3228 100644 --- gcc/oacc-builtins.def +++ gcc/oacc-builtins.def @@ -1,7 +1,7 @@ /* This file contains the definitions and documentation for the OpenACC builtins used in the GNU compiler. - Copyright (C) 2013 Free Software Foundation, Inc. + Copyright (C) 2013-2014 Free Software Foundation, Inc. Contributed by Thomas Schwinge <thomas@codesourcery.com>. @@ -29,3 +29,7 @@ along with GCC; see the file COPYING3. If not see DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel", BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) +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) diff --git gcc/omp-low.c gcc/omp-low.c index bca4599..6dec687 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { tree c, decl; bool scan_array_reductions = false; + bool offloaded; + switch (gimple_code (ctx->stmt)) + { + case GIMPLE_OACC_PARALLEL: + offloaded = true; + break; + case GIMPLE_OMP_TARGET: + switch (gimple_omp_target_kind (ctx->stmt)) + { + case GF_OMP_TARGET_KIND_REGION: + offloaded = true; + break; + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_UPDATE: + case GF_OMP_TARGET_KIND_OACC_DATA: + offloaded = false; + break; + default: + gcc_unreachable (); + } + break; + default: + offloaded = false; + } for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { @@ -1669,11 +1693,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) { /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in - #pragma omp target data, there is nothing to map for + target regions that are not offloaded; there is nothing to map for those. */ - if (!gimple_code_is_oacc (ctx->stmt) - && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA - && !POINTER_TYPE_P (TREE_TYPE (decl))) + if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } if (DECL_P (decl)) @@ -1698,9 +1720,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_field (decl, true, 7, ctx); else install_var_field (decl, true, 3, ctx); - if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL - || (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_REGION)) + if (offloaded) install_var_local (decl, ctx); } } @@ -1824,8 +1844,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET || (gimple_omp_target_kind (ctx->stmt) != GF_OMP_TARGET_KIND_UPDATE)); - if (!gimple_code_is_oacc (ctx->stmt) - && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) + if (!offloaded) break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) @@ -2340,7 +2359,7 @@ scan_omp_single (gimple stmt, omp_context *outer_ctx) layout_type (ctx->record_type); } -/* Scan an OpenMP target{, data, update} directive. */ +/* Scan a GIMPLE_OMP_TARGET. */ static void scan_omp_target (gimple stmt, omp_context *outer_ctx) @@ -2349,6 +2368,12 @@ 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) + { + gcc_assert (taskreg_nesting_level == 0); + gcc_assert (target_nesting_level == 0); + } + ctx = new_omp_context (stmt, outer_ctx); ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0); ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED; @@ -8218,7 +8243,7 @@ expand_omp_atomic (struct omp_region *region) } -/* Expand the OpenMP target{, data, update} directive starting at REGION. */ +/* Expand the GIMPLE_OMP_TARGET starting at REGION. */ static void expand_omp_target (struct omp_region *region) @@ -8401,12 +8426,23 @@ expand_omp_target (struct omp_region *region) clauses = gimple_omp_target_clauses (entry_stmt); - if (kind == GF_OMP_TARGET_KIND_REGION) - start_ix = BUILT_IN_GOMP_TARGET; - else if (kind == GF_OMP_TARGET_KIND_DATA) - start_ix = BUILT_IN_GOMP_TARGET_DATA; - else - start_ix = BUILT_IN_GOMP_TARGET_UPDATE; + switch (kind) + { + case GF_OMP_TARGET_KIND_REGION: + start_ix = BUILT_IN_GOMP_TARGET; + break; + case GF_OMP_TARGET_KIND_DATA: + start_ix = BUILT_IN_GOMP_TARGET_DATA; + break; + case GF_OMP_TARGET_KIND_UPDATE: + start_ix = BUILT_IN_GOMP_TARGET_UPDATE; + break; + case GF_OMP_TARGET_KIND_OACC_DATA: + start_ix = BUILT_IN_GOACC_DATA_START; + break; + default: + gcc_unreachable (); + } /* By default, the value of DEVICE is -1 (let runtime library choose) and there is no conditional. */ @@ -8414,10 +8450,12 @@ expand_omp_target (struct omp_region *region) device = build_int_cst (integer_type_node, -1); c = find_omp_clause (clauses, OMP_CLAUSE_IF); + gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA); if (c) 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) { device = OMP_CLAUSE_DEVICE_ID (c); @@ -8433,6 +8471,7 @@ expand_omp_target (struct omp_region *region) (cond ? device : -2). */ if (cond) { + gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA); cond = gimple_boolify (cond); basic_block cond_bb, then_bb, else_bb; @@ -8523,7 +8562,9 @@ expand_omp_target (struct omp_region *region) gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET); gsi_remove (&gsi, true); } - if (kind == GF_OMP_TARGET_KIND_DATA && region->exit) + if ((kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) + && region->exit) { gsi = gsi_last_bb (region->exit); g = gsi_stmt (gsi); @@ -10277,7 +10318,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) } } -/* Lower the OpenMP target directive in the current statement +/* Lower the GIMPLE_OMP_TARGET in the current statement in GSI_P. CTX holds context information for the directive. */ static void @@ -10298,7 +10339,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt)); tgt_body = gimple_bind_body (tgt_bind); } - else if (kind == GF_OMP_TARGET_KIND_DATA) + else if (kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) tgt_body = gimple_omp_body (stmt); child_fn = ctx->cb.dst_fn; @@ -10322,6 +10364,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_MAP_TOFROM: case OMP_CLAUSE_MAP_POINTER: break; + case OMP_CLAUSE_MAP_FORCE_ALLOC: + case OMP_CLAUSE_MAP_FORCE_TO: + case OMP_CLAUSE_MAP_FORCE_FROM: + case OMP_CLAUSE_MAP_FORCE_TOFROM: + 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); + break; default: gcc_unreachable (); } @@ -10330,6 +10381,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); var = OMP_CLAUSE_DECL (c); if (!DECL_P (var)) { @@ -10373,7 +10426,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) lower_omp (&tgt_body, ctx); target_nesting_level--; } - else if (kind == GF_OMP_TARGET_KIND_DATA) + else if (kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) lower_omp (&tgt_body, ctx); if (kind == GF_OMP_TARGET_KIND_REGION) @@ -10400,9 +10454,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1; TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1; TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1; + tree tkind_type; + int talign_shift; + switch (kind) + { + case GF_OMP_TARGET_KIND_REGION: + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_UPDATE: + tkind_type = unsigned_char_type_node; + talign_shift = 3; + break; + case GF_OMP_TARGET_KIND_OACC_DATA: + tkind_type = short_unsigned_type_node; + talign_shift = 8; + break; + default: + gcc_unreachable (); + } TREE_VEC_ELT (t, 2) - = create_tmp_var (build_array_type_nelts (unsigned_char_type_node, - map_cnt), + = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt), ".omp_data_kinds"); DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1; TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1; @@ -10515,7 +10585,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (TREE_CODE (s) != INTEGER_CST) TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0; - unsigned char tkind = 0; + unsigned HOST_WIDE_INT tkind; switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_MAP: @@ -10530,14 +10600,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: gcc_unreachable (); } - unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); + gcc_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift)); + unsigned HOST_WIDE_INT talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); talign = ceil_log2 (talign); - tkind |= talign << 3; + tkind |= talign << talign_shift; + gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); CONSTRUCTOR_APPEND_ELT (vkind, purpose, - build_int_cst (unsigned_char_type_node, - tkind)); + build_int_cstu (tkind_type, tkind)); if (nc && nc != c) c = nc; } @@ -10589,7 +10660,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_seq (&new_body, tgt_body); new_body = maybe_catch_exception (new_body); } - else if (kind == GF_OMP_TARGET_KIND_DATA) + 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) { @@ -10810,6 +10882,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) + gcc_assert (!ctx->cancellable); lower_omp_target (gsi_p, ctx); break; case GIMPLE_OMP_TEAMS: diff --git gcc/tree-inline.c gcc/tree-inline.c index 99903333..61c1cc8 100644 --- gcc/tree-inline.c +++ gcc/tree-inline.c @@ -1397,6 +1397,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id) break; case GIMPLE_OMP_TARGET: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); s1 = remap_gimple_seq (gimple_omp_body (stmt), id); copy = gimple_build_omp_target (s1, gimple_omp_target_kind (stmt), diff --git gcc/tree-nested.c gcc/tree-nested.c index 8933d02..afa7abb 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1307,6 +1307,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); save_suppress = info->suppress_expansion; convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, @@ -1769,6 +1770,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); save_suppress = info->suppress_expansion; convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); walk_body (convert_local_reference_stmt, convert_local_reference_op, @@ -2184,6 +2186,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p, case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_CRITICAL: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt)); break; diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 3dffde4..5c15656 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,10 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. + * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. + * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New + functions. + 2014-02-20 Thomas Schwinge <thomas@codesourcery.com> * target.c (gomp_load_plugin_for_device): Don't call dlcose if diff --git libgomp/libgomp.map libgomp/libgomp.map index 2b64d05..cb52e45 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -233,5 +233,7 @@ OACC_2.0 { GOACC_2.0 { global: + GOACC_data_end; + GOACC_data_start; GOACC_parallel; }; diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h index 7c24317..b9083a5 100644 --- libgomp/libgomp_g.h +++ libgomp/libgomp_g.h @@ -218,5 +218,8 @@ extern void GOMP_teams (unsigned int, unsigned int); extern void GOACC_parallel (int, void (*) (void *), const void *, size_t, void **, size_t *, unsigned short *); +extern void GOACC_data_start (int, const void *, + size_t, void **, size_t *, unsigned short *); +extern void GOACC_data_end (void); #endif /* LIBGOMP_G_H */ diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index bf7b74c..3ac7e39 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -1,4 +1,4 @@ -/* Copyright (C) 2013 Free Software Foundation, Inc. +/* Copyright (C) 2013-2014 Free Software Foundation, Inc. Contributed by Thomas Schwinge <thomas@codesourcery.com>. @@ -23,7 +23,7 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see <http://www.gnu.org/licenses/>. */ -/* This file handles the OpenACC parallel construct. */ +/* This file handles the OpenACC data and parallel constructs. */ #include "libgomp.h" #include "libgomp_g.h" @@ -51,3 +51,33 @@ GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target, } GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds_); } + + +void +GOACC_data_start (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_data (device, openmp_target, mapnum, hostaddrs, sizes, kinds_); +} + +void +GOACC_data_end (void) +{ + GOMP_target_end_data (); +} -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 3/3] OpenACC data construct support in the C front end. 2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge @ 2014-02-21 20:32 ` Thomas Schwinge 2014-03-12 13:48 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge 2014-03-20 14:39 ` [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) Thomas Schwinge 2 siblings, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-02-21 20:32 UTC (permalink / raw) To: gcc-patches; +Cc: jakub From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/c-family/ * c-pragma.c (oacc_pragmas): Add "data". * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA. gcc/c/ * c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition. (c_parser_oacc_data): New function. (c_parser_omp_construct): Handle PRAGMA_OACC_DATA. * c-tree.h (c_finish_oacc_data): New prototype. * c-typeck.c (c_finish_oacc_data): New function. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC data construct. * c-c++-common/goacc/nesting-fail-1.c: Likewise. * c-c++-common/goacc/parallel-fail-1.c: Rename to... * c-c++-common/goacc/clauses-fail.c: ... this new file. Extend for OpenACC data construct. * c-c++-common/goacc/data-1.c: New file. libgomp/ * testsuite/libgomp.oacc-c/data-1.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208017 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/c-family/ChangeLog.gomp | 5 + gcc/c-family/c-pragma.c | 1 + gcc/c-family/c-pragma.h | 1 + gcc/c/ChangeLog.gomp | 8 + gcc/c/c-parser.c | 42 +++++ gcc/c/c-tree.h | 1 + gcc/c/c-typeck.c | 19 +++ gcc/testsuite/ChangeLog.gomp | 10 ++ .../c-c++-common/goacc-gomp/nesting-fail-1.c | 92 ++++++++++- gcc/testsuite/c-c++-common/goacc/clauses-fail.c | 9 ++ gcc/testsuite/c-c++-common/goacc/data-1.c | 6 + gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 18 ++- gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c | 6 - libgomp/ChangeLog.gomp | 2 + libgomp/testsuite/libgomp.oacc-c/data-1.c | 170 +++++++++++++++++++++ 15 files changed, 380 insertions(+), 10 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/clauses-fail.c create mode 100644 gcc/testsuite/c-c++-common/goacc/data-1.c delete mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c/data-1.c diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp index e092d53..3da377f 100644 --- gcc/c-family/ChangeLog.gomp +++ gcc/c-family/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * c-pragma.c (oacc_pragmas): Add "data". + * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA. + 2014-01-28 Thomas Schwinge <thomas@codesourcery.com> * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY, diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c index f69486a..08374aa 100644 --- gcc/c-family/c-pragma.c +++ gcc/c-family/c-pragma.c @@ -1169,6 +1169,7 @@ static vec<pragma_ns_name> registered_pp_pragmas; struct omp_pragma_def { const char *name; unsigned int id; }; static const struct omp_pragma_def oacc_pragmas[] = { + { "data", PRAGMA_OACC_DATA }, { "parallel", PRAGMA_OACC_PARALLEL }, }; static const struct omp_pragma_def omp_pragmas[] = { diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h index 1ea5b1d..d092f9f 100644 --- gcc/c-family/c-pragma.h +++ gcc/c-family/c-pragma.h @@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see typedef enum pragma_kind { PRAGMA_NONE = 0, + PRAGMA_OACC_DATA, PRAGMA_OACC_PARALLEL, PRAGMA_OMP_ATOMIC, PRAGMA_OMP_BARRIER, diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index b199957..9b95725 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,11 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition. + (c_parser_oacc_data): New function. + (c_parser_omp_construct): Handle PRAGMA_OACC_DATA. + * c-tree.h (c_finish_oacc_data): New prototype. + * c-typeck.c (c_finish_oacc_data): New function. + 2014-02-17 Thomas Schwinge <thomas@codesourcery.com> * c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin, diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 7850eab..4643722 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -4776,10 +4776,14 @@ c_parser_label (c_parser *parser) openacc-construct: parallel-construct + data-construct parallel-construct: parallel-directive structured-block + data-construct: + data-directive structured-block + OpenMP: statement: @@ -11362,6 +11366,41 @@ c_parser_omp_structured_block (c_parser *parser) } /* OpenACC 2.0: + # pragma acc data oacc-data-clause[optseq] new-line + structured-block + + LOC is the location of the #pragma token. +*/ + +#define OACC_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) ) + +static tree +c_parser_oacc_data (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK, + "#pragma acc data"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + + stmt = c_finish_oacc_data (loc, clauses, block); + + return stmt; +} + +/* OpenACC 2.0: # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block @@ -13675,6 +13714,9 @@ c_parser_omp_construct (c_parser *parser) switch (p_kind) { + case PRAGMA_OACC_DATA: + stmt = c_parser_oacc_data (loc, parser); + break; case PRAGMA_OACC_PARALLEL: stmt = c_parser_oacc_parallel (loc, parser); break; diff --git gcc/c/c-tree.h gcc/c/c-tree.h index c174c7a..c84d3d7 100644 --- gcc/c/c-tree.h +++ gcc/c/c-tree.h @@ -634,6 +634,7 @@ extern tree c_finish_goto_label (location_t, tree); extern tree c_finish_goto_ptr (location_t, tree); extern tree c_expr_to_decl (tree, bool *, bool *); extern tree c_finish_oacc_parallel (location_t, tree, tree); +extern tree c_finish_oacc_data (location_t, tree, tree); extern tree c_begin_omp_parallel (void); extern tree c_finish_omp_parallel (location_t, tree, tree); extern tree c_begin_omp_task (void); diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 76d655b..8c4445b 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -11122,6 +11122,25 @@ c_finish_oacc_parallel (location_t loc, tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_DATA. */ + +tree +c_finish_oacc_data (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_DATA_CLAUSES (stmt) = clauses; + OACC_DATA_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Like c_begin_compound_stmt, except force the retention of the BLOCK. */ tree diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index fbccfa3..41d73b6 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,13 @@ +2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + + * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC + data construct. + * c-c++-common/goacc/nesting-fail-1.c: Likewise. + * c-c++-common/goacc/parallel-fail-1.c: Rename to... + * c-c++-common/goacc/clauses-fail.c: ... this new file. Extend + for OpenACC data construct. + * c-c++-common/goacc/data-1.c: New file. + 2014-02-18 Thomas Schwinge <thomas@codesourcery.com> * gcc.dg/goacc/parallel-sb-1.c: New file. diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 875ec66..78fb45b 100644 --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -1,7 +1,7 @@ /* TODO: Some of these should either be allowed or fail with a more sensible error message. */ void -f1 (void) +f_omp (void) { int i; @@ -9,6 +9,8 @@ f1 (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp for @@ -16,49 +18,68 @@ f1 (void) { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp sections { + { #pragma acc parallel /* { dg-error "may not be nested" } */ - ; + ; + } +#pragma omp section + { +#pragma acc data /* { dg-error "may not be nested" } */ + ; + } } #pragma omp single { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp task { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp master { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp critical { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } #pragma omp ordered { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } } /* TODO: Some of these should either be allowed or fail with a more sensible error message. */ void -f2 (void) +f_acc_parallel (void) { #pragma acc parallel { @@ -119,3 +140,68 @@ f2 (void) ; } } + +/* TODO: Some of these should either be allowed or fail with a more sensible + error message. */ +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma omp parallel /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { + int i; +#pragma omp for /* { dg-error "may not be nested" } */ + for (i = 0; i < 3; i++) + ; + } + +#pragma acc data + { +#pragma omp sections /* { dg-error "may not be nested" } */ + { + ; + } + } + +#pragma acc data + { +#pragma omp single /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { +#pragma omp task /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { +#pragma omp master /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { +#pragma omp critical /* { dg-error "may not be nested" } */ + ; + } + +#pragma acc data + { + int i; +#pragma omp atomic write + i = 0; /* { dg-error "may not be nested" } */ + } + +#pragma acc data + { +#pragma omp ordered /* { dg-error "may not be nested" } */ + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c new file mode 100644 index 0000000..b0dd042 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c @@ -0,0 +1,9 @@ +void +f (void) +{ +#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */ + ; + +#pragma acc data two /* { dg-error "expected clause before 'two'" } */ + ; +} diff --git gcc/testsuite/c-c++-common/goacc/data-1.c gcc/testsuite/c-c++-common/goacc/data-1.c new file mode 100644 index 0000000..8094575 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/data-1.c @@ -0,0 +1,6 @@ +void +foo (void) +{ +#pragma acc data + ; +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 6501397..24a4c11 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -1,11 +1,27 @@ /* TODO: While the OpenACC specification does allow for certain kinds of nesting, we don't support that yet. */ void -f1 (void) +f_acc_parallel (void) { #pragma acc parallel { #pragma acc parallel /* { dg-error "may not be nested" } */ ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; + } +} + +/* TODO: While the OpenACC specification does allow for certain kinds of + nesting, we don't support that yet. */ +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma acc parallel /* { dg-error "may not be nested" } */ + ; +#pragma acc data /* { dg-error "may not be nested" } */ + ; } } diff --git gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c deleted file mode 100644 index efc6f14..0000000 --- gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c +++ /dev/null @@ -1,6 +0,0 @@ -void -foo (void) -{ -#pragma acc parallel foo /* { dg-error "expected clause before 'foo'" } */ - foo (); -} diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 5c15656..b90b09b 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,7 @@ 2014-02-21 Thomas Schwinge <thomas@codesourcery.com> + * testsuite/libgomp.oacc-c/data-1.c: New file. + * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New diff --git libgomp/testsuite/libgomp.oacc-c/data-1.c libgomp/testsuite/libgomp.oacc-c/data-1.c new file mode 100644 index 0000000..8f9a17a --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c/data-1.c @@ -0,0 +1,170 @@ +/* { dg-do run } */ + +extern void abort (); + +int i; + +int main(void) +{ + int j; + +#if 0 + i = -1; + j = -2; +#pragma acc data copyin (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != 2 || j != 1) + abort (); + + i = -1; + j = -2; +#pragma acc data copyout (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); + + i = -1; + j = -2; +#pragma acc data copy (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); + + i = -1; + j = -2; +#pragma acc data create (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + + i = -1; + j = -2; +#pragma acc data present_or_copyin (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != 2 || j != 1) + abort (); + +#if 0 + i = -1; + j = -2; +#pragma acc data present_or_copyout (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + + i = -1; + j = -2; +#pragma acc data present_or_copy (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); + +#if 0 + i = -1; + j = -2; +#pragma acc data present_or_create (i, j) + { + // TODO: check that variables have been mapped. + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + +#if 0 + i = -1; + j = -2; +#pragma acc data present (i, j) + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + +#if 0 + i = -1; + j = -2; +#pragma acc data + { + // TODO: check that variables have been mapped. + if (i != -1 || j != -2) + abort (); + i = 2; + j = 1; + if (i != 2 || j != 1) + abort (); + } + if (i != -1 || j != -2) + abort (); +#endif + + return 0; +} -- 1.8.1.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA. 2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge 2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge @ 2014-03-12 13:48 ` Thomas Schwinge 2014-03-20 14:39 ` [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) Thomas Schwinge 2 siblings, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-03-12 13:48 UTC (permalink / raw) To: gcc-patches; +Cc: jakub [-- Attachment #1: Type: text/plain, Size: 4687 bytes --] Hi! On Fri, 21 Feb 2014 21:32:14 +0100, I wrote: > --- gcc/omp-low.c > +++ gcc/omp-low.c > @@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) > { > tree c, decl; > bool scan_array_reductions = false; > + bool offloaded; > + switch (gimple_code (ctx->stmt)) > + { > + case GIMPLE_OACC_PARALLEL: > + offloaded = true; > + break; > + case GIMPLE_OMP_TARGET: > + switch (gimple_omp_target_kind (ctx->stmt)) > + { > + case GF_OMP_TARGET_KIND_REGION: > + offloaded = true; > + break; > + case GF_OMP_TARGET_KIND_DATA: > + case GF_OMP_TARGET_KIND_UPDATE: > + case GF_OMP_TARGET_KIND_OACC_DATA: > + offloaded = false; > + break; > + default: > + gcc_unreachable (); > + } > + break; > + default: > + offloaded = false; > + } I now have a need for this information elsewhere; in gomp-4_0-branch r208513 changed as follows: commit 326592ef8fe7501f9ba7e67157d68c6c541e5601 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Wed Mar 12 13:40:07 2014 +0000 is_gimple_omp_offloaded. gcc/ * omp-low.c (scan_sharing_clauses): Move offloaded logic into... * gimple.h (is_gimple_omp_offloaded): ... this new static inline function. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208513 138bc75d-0d04-0410-961f-82ee72b054a4 diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 79030d6..4ee843f 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,9 @@ +2014-03-12 Thomas Schwinge <thomas@codesourcery.com> + + * omp-low.c (scan_sharing_clauses): Move offloaded logic into... + * gimple.h (is_gimple_omp_offloaded): ... this new static inline + function. + 2014-02-28 Thomas Schwinge <thomas@codesourcery.com> * gimple.def (GIMPLE_OACC_KERNELS): New code. diff --git gcc/gimple.h gcc/gimple.h index 514af32..910072d 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -5823,6 +5823,31 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) } +/* Return true if OMP_* STMT is offloaded. */ + +static inline bool +is_gimple_omp_offloaded (const_gimple stmt) +{ + gcc_assert (is_gimple_omp (stmt)); + switch (gimple_code (stmt)) + { + case GIMPLE_OACC_KERNELS: + case GIMPLE_OACC_PARALLEL: + return true; + case GIMPLE_OMP_TARGET: + switch (gimple_omp_target_kind (stmt)) + { + case GF_OMP_TARGET_KIND_REGION: + return true; + default: + return false; + } + default: + return false; + } +} + + /* Returns TRUE if statement G is a GIMPLE_NOP. */ static inline bool diff --git gcc/omp-low.c gcc/omp-low.c index 2f13fb4..6b676e5 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1499,31 +1499,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { tree c, decl; bool scan_array_reductions = false; - bool offloaded; - switch (gimple_code (ctx->stmt)) - { - case GIMPLE_OACC_KERNELS: - case GIMPLE_OACC_PARALLEL: - offloaded = true; - break; - case GIMPLE_OMP_TARGET: - switch (gimple_omp_target_kind (ctx->stmt)) - { - case GF_OMP_TARGET_KIND_REGION: - offloaded = true; - break; - case GF_OMP_TARGET_KIND_DATA: - case GF_OMP_TARGET_KIND_UPDATE: - case GF_OMP_TARGET_KIND_OACC_DATA: - offloaded = false; - break; - default: - gcc_unreachable (); - } - break; - default: - offloaded = false; - } for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { @@ -1696,7 +1671,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in target regions that are not offloaded; there is nothing to map for those. */ - if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl))) + if (!is_gimple_omp_offloaded (ctx->stmt) + && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } if (DECL_P (decl)) @@ -1721,7 +1697,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_field (decl, true, 7, ctx); else install_var_field (decl, true, 3, ctx); - if (offloaded) + if (is_gimple_omp_offloaded (ctx->stmt)) install_var_local (decl, ctx); } } @@ -1845,7 +1821,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET || (gimple_omp_target_kind (ctx->stmt) != GF_OMP_TARGET_KIND_UPDATE)); - if (!offloaded) + if (!is_gimple_omp_offloaded (ctx->stmt)) break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) 2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge 2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge 2014-03-12 13:48 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge @ 2014-03-20 14:39 ` Thomas Schwinge 2 siblings, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-03-20 14:39 UTC (permalink / raw) To: gcc-patches [-- Attachment #1: Type: text/plain, Size: 6560 bytes --] Hi! Applied in r208701 to gomp-4_0-branch: commit 22dd36a31c433dcd8bcc890d245a9e4ac6ed9c7f Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Thu Mar 20 14:33:28 2014 +0000 Nesting of OpenACC constructs inside of OpenACC data constructs. gcc/ * omp-low.c (check_omp_nesting_restrictions): Allow nesting of OpenACC constructs inside of OpenACC data constructs. gcc/testsuite/ * c-c++-common/goacc/nesting-1.c: New file. * c-c++-common/goacc/nesting-data-1.c: Likewise. * c-c++-common/goacc/nesting-fail-1.c: Update. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208701 138bc75d-0d04-0410-961f-82ee72b054a4 diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 1aebc4d..f43452c 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-03-20 Thomas Schwinge <thomas@codesourcery.com> + + * omp-low.c (check_omp_nesting_restrictions): Allow nesting of + OpenACC constructs inside of OpenACC data constructs. + 2014-03-18 Ilmir Usmanov <i.usmanov@samsung.com> * tree.def (OACC_LOOP): New tree code. diff --git gcc/omp-low.c gcc/omp-low.c index f1b0fa5..23a0dda 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -2416,26 +2416,31 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx) static bool check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) { - omp_context *ctx_; - /* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ - /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin) - inside any OpenACC CTX. */ - for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) - if (is_gimple_omp (ctx_->stmt) - && is_gimple_omp_oacc_specifically (ctx_->stmt)) - { - error_at (gimple_location (stmt), - "may not be nested"); - return false; - } - /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX. */ + nesting, we don't support many of these yet. */ if (is_gimple_omp (stmt) && is_gimple_omp_oacc_specifically (stmt)) { - for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) - if (is_gimple_omp (ctx_->stmt)) + /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different + from an OpenACC data construct. */ + for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) + if (is_gimple_omp (ctx_->stmt) + && !(gimple_code (ctx_->stmt) == GIMPLE_OMP_TARGET + && (gimple_omp_target_kind (ctx_->stmt) + == GF_OMP_TARGET_KIND_OACC_DATA))) + { + error_at (gimple_location (stmt), + "may not be nested"); + return false; + } + } + else + { + /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP + builtin) inside any OpenACC CTX. */ + for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) + if (is_gimple_omp (ctx_->stmt) + && is_gimple_omp_oacc_specifically (ctx_->stmt)) { error_at (gimple_location (stmt), "may not be nested"); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index fd38d80..13e99d5 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,9 @@ 2014-03-20 Thomas Schwinge <thomas@codesourcery.com> + * c-c++-common/goacc/nesting-1.c: New file. + * c-c++-common/goacc/nesting-data-1.c: Likewise. + * c-c++-common/goacc/nesting-fail-1.c: Update. + * c-c++-common/goacc/nesting-fail-1.c (f_acc_kernels): Replace OpenACC parallel with kernels directive. diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c new file mode 100644 index 0000000..3a22292 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -0,0 +1,13 @@ +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma acc parallel + ; +#pragma acc kernels + ; +#pragma acc data + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-data-1.c gcc/testsuite/c-c++-common/goacc/nesting-data-1.c new file mode 100644 index 0000000..fefe6cd --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/nesting-data-1.c @@ -0,0 +1,61 @@ +void +f (void) +{ + unsigned char c, ca[15], caa[20][30]; + +#pragma acc data copyin(c) + { + c = 5; + ca[3] = c; + caa[3][12] = ca[3] + caa[3][12]; + +#pragma acc data copyin(ca[2:4]) + { + c = 6; + ca[4] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc parallel copyout(ca[3:4]) + { + c = 7; + ca[5] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc kernels copy(ca[4:4]) + { + c = 8; + ca[6] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc data pcopy(ca[5:7]) + { + c = 15; + ca[7] = c; + caa[3][12] = ca[3] + caa[3][12]; + +#pragma acc data pcopyin(caa[3:7][0:30]) + { + c = 16; + ca[8] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc parallel pcopyout(caa[3:7][0:30]) + { + c = 17; + ca[9] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc kernels pcopy(caa[3:7][0:30]) + { + c = 18; + ca[10] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + } + } +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index ca8921f..00dc602 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -1,5 +1,5 @@ /* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ + nesting, we don't support many of these yet. */ void f_acc_parallel (void) { @@ -15,7 +15,7 @@ f_acc_parallel (void) } /* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ + nesting, we don't support many of these yet. */ void f_acc_kernels (void) { @@ -29,19 +29,3 @@ f_acc_kernels (void) ; } } - -/* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ -void -f_acc_data (void) -{ -#pragma acc data - { -#pragma acc parallel /* { dg-error "may not be nested" } */ - ; -#pragma acc kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; - } -} Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
[parent not found: <538DF785.3050206@mentor.com>]
[parent not found: <87egz645j4.fsf@schwinge.name>]
* Re: [gomp4 3/6] Initial support for OpenACC memory mapping semantics. [not found] ` <87egz645j4.fsf@schwinge.name> @ 2014-11-13 12:21 ` Thomas Schwinge 2014-11-13 13:13 ` Jakub Jelinek 0 siblings, 1 reply; 22+ messages in thread From: Thomas Schwinge @ 2014-11-13 12:21 UTC (permalink / raw) To: gcc-patches [-- Attachment #1: Type: text/plain, Size: 25086 bytes --] Hi! On Tue, 14 Jan 2014 16:10:05 +0100, I wrote: > --- gcc/gimplify.c > +++ gcc/gimplify.c > @@ -69,7 +69,13 @@ enum gimplify_omp_var_data > + /* Force a specific behavior (or else, a run-time error). */ > + GOVD_MAP_FORCE = 16384, > @@ -86,7 +92,11 @@ enum omp_region_type > + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */ > + ORT_TARGET_MAP_FORCE = 64 > }; > @@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) > OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; > else if (code == OMP_CLAUSE_MAP) > { > - OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY > - ? OMP_CLAUSE_MAP_TO > - : OMP_CLAUSE_MAP_TOFROM; > + unsigned map_kind; > + map_kind = (flags & GOVD_MAP_TO_ONLY > + ? OMP_CLAUSE_MAP_TO > + : OMP_CLAUSE_MAP_TOFROM); > + if (flags & GOVD_MAP_FORCE) > + map_kind |= OMP_CLAUSE_MAP_FORCE; > + OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind; > + > if (DECL_SIZE (decl) > && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) > { > @@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p) > tree expr = *expr_p; > gimple g; > gimple_seq body = NULL; > + enum omp_region_type ort = > + (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE); > > - gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, > - ORT_TARGET); > + gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort); > > push_gimplify_context (); I don't remember what I have been thinking when implementing this -- per the OpenACC specification's rules for implicitly determined data attributes, it should be present_or_copy (that is, OpenMP's tofrom, without "force" semantics), and firstprivate/copy for scalar variables for the parallel/kernels constructs, respectively (which is still to be implemented, for now not considering scalar variables different from non-scalar ones). Committed to gomp-4_0-branch in r217482: commit 7058203891bd6e1696763603673090f161e172b8 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Thu Nov 13 12:18:34 2014 +0000 Middle end: Don't use mapping kinds with "force" semantics for OpenACC. ..., which is the wrong thing to do. Also extend libgomp to actually distinguish between "non-force"/"force" semantics. gcc/ * gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS, OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE. (enum gimplify_omp_var_data, enum omp_region_type): Remove GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update all users. include/ * gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and _GOMP_MAP_FLAG_FORCE. libgomp/ * target.c (gomp_map_vars_existing): Error out if "force" semantics. (gomp_map_vars): Actually pass kinds to gomp_map_vars_existing. Remove FIXMEs. * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise. * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217482 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 6 ++ gcc/gimplify.c | 65 +++------------------- include/ChangeLog.gomp | 4 ++ include/gomp-constants.h | 3 + libgomp/ChangeLog.gomp | 23 ++++++++ libgomp/target.c | 21 ++++--- .../libgomp.oacc-c-c++-common/data-already-1.c | 19 +++++++ .../libgomp.oacc-c-c++-common/data-already-2.c | 16 ++++++ .../libgomp.oacc-c-c++-common/data-already-3.c | 17 ++++++ .../libgomp.oacc-c-c++-common/data-already-4.c | 17 ++++++ .../libgomp.oacc-c-c++-common/data-already-5.c | 17 ++++++ .../libgomp.oacc-c-c++-common/data-already-6.c | 17 ++++++ .../libgomp.oacc-c-c++-common/data-already-7.c | 17 ++++++ .../libgomp.oacc-c-c++-common/data-already-8.c | 16 ++++++ .../libgomp.oacc-fortran/data-already-1.f | 17 ++++++ .../libgomp.oacc-fortran/data-already-2.f | 16 ++++++ .../libgomp.oacc-fortran/data-already-3.f | 15 +++++ .../libgomp.oacc-fortran/data-already-4.f | 14 +++++ .../libgomp.oacc-fortran/data-already-5.f | 14 +++++ .../libgomp.oacc-fortran/data-already-6.f | 14 +++++ .../libgomp.oacc-fortran/data-already-7.f | 14 +++++ .../libgomp.oacc-fortran/data-already-8.f | 16 ++++++ 22 files changed, 311 insertions(+), 67 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 174235d..a499755 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,11 @@ 2014-11-13 Thomas Schwinge <thomas@codesourcery.com> + * gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS, + OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE. + (enum gimplify_omp_var_data, enum omp_region_type): Remove + GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update + all users. + * omp-low.c (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Revert earlier change. diff --git gcc/gimplify.c gcc/gimplify.c index 233ac56..2c8c666 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -94,8 +94,6 @@ enum gimplify_omp_var_data /* Flags for GOVD_MAP. */ /* Don't copy back. */ GOVD_MAP_TO_ONLY = 8192, - /* Force a specific behavior (or else, a run-time error). */ - GOVD_MAP_FORCE = 16384, GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR @@ -116,9 +114,7 @@ enum omp_region_type /* Flags for ORT_TARGET. */ /* Prepare this region for offloading. */ - ORT_TARGET_OFFLOAD = 32, - /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */ - ORT_TARGET_MAP_FORCE = 64 + ORT_TARGET_OFFLOAD = 32 }; /* Gimplify hashtable helper. */ @@ -5585,15 +5581,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) if (!(flags & GOVD_LOCAL)) { if (flags & GOVD_MAP) - { - nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT; -#if 0 - /* Not sure if this is actually needed; haven't found a case - where this would change anything; TODO. */ - if (flags & GOVD_MAP_FORCE) - nflags |= OMP_CLAUSE_MAP_FORCE; -#endif - } + nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT; else if (flags & GOVD_PRIVATE) nflags = GOVD_PRIVATE; else @@ -5667,8 +5655,6 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, if ((octx->region_type & ORT_TARGET) && (octx->region_type & ORT_TARGET_OFFLOAD)) { - gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE)); - n = splay_tree_lookup (octx->variables, (splay_tree_key)decl); if (n == NULL) { @@ -5731,11 +5717,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) if ((ctx->region_type & ORT_TARGET) && (ctx->region_type & ORT_TARGET_OFFLOAD)) { - unsigned map_force; - if (ctx->region_type & ORT_TARGET_MAP_FORCE) - map_force = GOVD_MAP_FORCE; - else - map_force = 0; ret = lang_hooks.decls.omp_disregard_value_expr (decl, true); if (n == NULL) { @@ -5743,32 +5724,13 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) { error ("%qD referenced in target region does not have " "a mappable type", decl); - omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT | flags); + omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags); } else - omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags); + omp_add_variable (ctx, decl, GOVD_MAP | flags); } else { -#if 0 - /* The following fails for: - - int l = 10; - float c[l]; - #pragma acc parallel copy(c[2:4]) - { - #pragma acc parallel - { - int t = sizeof c; - } - } - - ..., which we currently don't have to care about (nesting - disabled), but eventually will have to; TODO. */ - if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT)) - gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force); -#endif - /* If nothing changed, there's nothing left to do. */ if ((n->value & flags) == flags) return ret; @@ -6423,13 +6385,11 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; else if (code == OMP_CLAUSE_MAP) { - unsigned map_kind; + enum omp_clause_map_kind map_kind; map_kind = (flags & GOVD_MAP_TO_ONLY ? OMP_CLAUSE_MAP_TO : OMP_CLAUSE_MAP_TOFROM); - if (flags & GOVD_MAP_FORCE) - map_kind |= OMP_CLAUSE_MAP_FORCE; - OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind; + OMP_CLAUSE_MAP_KIND (clause) = map_kind; if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) @@ -7258,23 +7218,16 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) switch (TREE_CODE (expr)) { - case OACC_DATA: - ort = (enum omp_region_type) (ORT_TARGET - | ORT_TARGET_MAP_FORCE); - break; - case OACC_KERNELS: - case OACC_PARALLEL: - ort = (enum omp_region_type) (ORT_TARGET - | ORT_TARGET_OFFLOAD - | ORT_TARGET_MAP_FORCE); - break; case OMP_SECTIONS: case OMP_SINGLE: ort = ORT_WORKSHARE; break; + case OACC_KERNELS: + case OACC_PARALLEL: case OMP_TARGET: ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD); break; + case OACC_DATA: case OMP_TARGET_DATA: ort = ORT_TARGET; break; diff --git include/ChangeLog.gomp include/ChangeLog.gomp new file mode 100644 index 0000000..9172c26 --- /dev/null +++ include/ChangeLog.gomp @@ -0,0 +1,4 @@ +2014-11-13 Thomas Schwinge <thomas@codesourcery.com> + + * gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and + _GOMP_MAP_FLAG_FORCE. diff --git include/gomp-constants.h include/gomp-constants.h index e600766..15b658f 100644 --- include/gomp-constants.h +++ include/gomp-constants.h @@ -28,6 +28,9 @@ /* Enumerated variable mapping types used to communicate between GCC and libgomp. These values are used for both OpenMP and OpenACC. */ +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2) +#define _GOMP_MAP_FLAG_FORCE (1 << 3) + #define GOMP_MAP_ALLOC 0x00 #define GOMP_MAP_ALLOC_TO 0x01 #define GOMP_MAP_ALLOC_FROM 0x02 diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 0528531..254846f 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,26 @@ +2014-11-13 Thomas Schwinge <thomas@codesourcery.com> + + * target.c (gomp_map_vars_existing): Error out if "force" + semantics. + (gomp_map_vars): Actually pass kinds to gomp_map_vars_existing. + Remove FIXMEs. + * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: New file. + * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise. + * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise. + 2014-11-12 Thomas Schwinge <thomas@codesourcery.com> * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file. diff --git libgomp/target.c libgomp/target.c index 052c59d..2b9f08f 100644 --- libgomp/target.c +++ libgomp/target.c @@ -117,9 +117,11 @@ static inline void gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn, unsigned char kind) { - if (oldn->host_start > newn->host_start + if ((!(kind & _GOMP_MAP_FLAG_SPECIAL) + && (kind & _GOMP_MAP_FLAG_FORCE)) + || oldn->host_start > newn->host_start || oldn->host_end < newn->host_end) - gomp_fatal ("Trying to map into device [%p..%p) object when" + gomp_fatal ("Trying to map into device [%p..%p) object when " "[%p..%p) is already mapped", (void *) newn->host_start, (void *) newn->host_end, (void *) oldn->host_start, (void *) oldn->host_end); @@ -200,7 +202,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, if (n) { tgt->list[i] = n; - gomp_map_vars_existing (n, &cur_node, kind); + gomp_map_vars_existing (n, &cur_node, kind & typemask); } else { @@ -323,7 +325,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, if (n) { tgt->list[i] = n; - gomp_map_vars_existing (n, k, kind); + gomp_map_vars_existing (n, k, kind & typemask); } else { @@ -345,18 +347,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum, switch (kind & typemask) { - case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_FROM: - /* FIXME: No special handling (see comment in - oacc-parallel.c). */ case GOMP_MAP_ALLOC: case GOMP_MAP_ALLOC_FROM: + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_FROM: break; - case GOMP_MAP_FORCE_TO: - case GOMP_MAP_FORCE_TOFROM: - /* FIXME: No special handling, as above. */ case GOMP_MAP_ALLOC_TO: case GOMP_MAP_ALLOC_TOFROM: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_TOFROM: /* Copy from host to device memory. */ /* FIXME: Perhaps add some smarts, like if copying several adjacent fields from host to target, use some diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c new file mode 100644 index 0000000..83c0a42 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c @@ -0,0 +1,19 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include <openacc.h> + +int +main (int argc, char *argv[]) +{ + int i; + + acc_copyin (&i, sizeof i); + +#pragma acc data copy (i) + ++i; + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "Trying to map into device .* object when .* is already mapped" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c new file mode 100644 index 0000000..137d8ce --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c @@ -0,0 +1,16 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +int +main (int argc, char *argv[]) +{ + int i; + +#pragma acc data present_or_copy (i) +#pragma acc data copyout (i) + ++i; + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "Trying to map into device .* object when .* is already mapped" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c new file mode 100644 index 0000000..b993b78 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c @@ -0,0 +1,17 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include <openacc.h> + +int +main (int argc, char *argv[]) +{ + int i; + +#pragma acc data present_or_copy (i) + acc_copyin (&i, sizeof i); + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "already mapped to" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c new file mode 100644 index 0000000..82523f4 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c @@ -0,0 +1,17 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include <openacc.h> + +int +main (int argc, char *argv[]) +{ + int i; + + acc_present_or_copyin (&i, sizeof i); + acc_copyin (&i, sizeof i); + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "already mapped to" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c new file mode 100644 index 0000000..4961fe5 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c @@ -0,0 +1,17 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include <openacc.h> + +int +main (int argc, char *argv[]) +{ + int i; + +#pragma acc enter data create (i) + acc_copyin (&i, sizeof i); + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "already mapped to" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c new file mode 100644 index 0000000..77b56a9 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c @@ -0,0 +1,17 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include <openacc.h> + +int +main (int argc, char *argv[]) +{ + int i; + + acc_present_or_copyin (&i, sizeof i); +#pragma acc enter data create (i) + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "already mapped to" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c new file mode 100644 index 0000000..b08417b --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c @@ -0,0 +1,17 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +#include <openacc.h> + +int +main (int argc, char *argv[]) +{ + int i; + +#pragma acc enter data create (i) + acc_create (&i, sizeof i); + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "already mapped to" } */ diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c new file mode 100644 index 0000000..a50f7de --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c @@ -0,0 +1,16 @@ +/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ + +int +main (int argc, char *argv[]) +{ + int i; + +#pragma acc data create (i) +#pragma acc parallel copyin (i) + ++i; + + return 0; +} + +/* { dg-shouldfail "" } + { dg-output "Trying to map into device .* object when .* is already mapped" } */ diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f new file mode 100644 index 0000000..ac220ab --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f @@ -0,0 +1,17 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + + INTEGER I + + CALL ACC_COPYIN (I) + +!$ACC DATA COPY (I) + I = 0 +!$ACC END DATA + + END + +! { dg-shouldfail "" } +! { dg-output "Trying to map into device .* object when .* is already mapped" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f new file mode 100644 index 0000000..2c5254b --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f @@ -0,0 +1,16 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + + INTEGER I + +!$ACC DATA PRESENT_OR_COPY (I) +!$ACC DATA COPYOUT (I) + I = 0 +!$ACC END DATA +!$ACC END DATA + + END + +! { dg-shouldfail "" } +! { dg-output "Trying to map into device .* object when .* is already mapped" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f new file mode 100644 index 0000000..c41de28 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f @@ -0,0 +1,15 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + + INTEGER I + +!$ACC DATA PRESENT_OR_COPY (I) + CALL ACC_COPYIN (I) +!$ACC END DATA + + END + +! { dg-shouldfail "" } +! { dg-output "already mapped to" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f new file mode 100644 index 0000000..f54bf58 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f @@ -0,0 +1,14 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + + INTEGER I + + CALL ACC_PRESENT_OR_COPYIN (I) + CALL ACC_COPYIN (I) + + END + +! { dg-shouldfail "" } +! { dg-output "already mapped to" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f new file mode 100644 index 0000000..9a3e94f --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f @@ -0,0 +1,14 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + + INTEGER I + +!$ACC ENTER DATA CREATE (I) + CALL ACC_COPYIN (I) + + END + +! { dg-shouldfail "" } +! { dg-output "already mapped to" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f new file mode 100644 index 0000000..eaf5d98 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f @@ -0,0 +1,14 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + + INTEGER I + + CALL ACC_PRESENT_OR_COPYIN (I) +!$ACC ENTER DATA CREATE (I) + + END + +! { dg-shouldfail "" } +! { dg-output "already mapped to" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f new file mode 100644 index 0000000..d96bf0b --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f @@ -0,0 +1,14 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + INCLUDE "openacc_lib.h" + + INTEGER I + +!$ACC ENTER DATA CREATE (I) + CALL ACC_CREATE (I) + + END + +! { dg-shouldfail "" } +! { dg-output "already mapped to" } diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f new file mode 100644 index 0000000..16da048 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f @@ -0,0 +1,16 @@ +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + + IMPLICIT NONE + + INTEGER I + +!$ACC DATA CREATE (I) +!$ACC PARALLEL COPYIN (I) + I = 0 +!$ACC END PARALLEL +!$ACC END DATA + + END + +! { dg-shouldfail "" } +! { dg-output "Trying to map into device .* object when .* is already mapped" } Grüße, Thomas [-- Attachment #2: signature.asc --] [-- Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 3/6] Initial support for OpenACC memory mapping semantics. 2014-11-13 12:21 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge @ 2014-11-13 13:13 ` Jakub Jelinek 2014-11-13 13:39 ` gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) Thomas Schwinge 0 siblings, 1 reply; 22+ messages in thread From: Jakub Jelinek @ 2014-11-13 13:13 UTC (permalink / raw) To: Thomas Schwinge; +Cc: gcc-patches On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote: > --- include/gomp-constants.h > +++ include/gomp-constants.h > @@ -28,6 +28,9 @@ > /* Enumerated variable mapping types used to communicate between GCC and > libgomp. These values are used for both OpenMP and OpenACC. */ > > +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2) > +#define _GOMP_MAP_FLAG_FORCE (1 << 3) I'm worried about reserved namespace issues if you use _ followed by capital letter. Can't it be just GOMP_MAP_FLAG_* ? Jakub ^ permalink raw reply [flat|nested] 22+ messages in thread
* gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) 2014-11-13 13:13 ` Jakub Jelinek @ 2014-11-13 13:39 ` Thomas Schwinge 2014-11-13 13:57 ` Jakub Jelinek 0 siblings, 1 reply; 22+ messages in thread From: Thomas Schwinge @ 2014-11-13 13:39 UTC (permalink / raw) To: Jakub Jelinek; +Cc: gcc-patches [-- Attachment #1: Type: text/plain, Size: 1186 bytes --] Hi Jakub! On Thu, 13 Nov 2014 14:10:10 +0100, Jakub Jelinek <jakub@redhat.com> wrote: > On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote: > > --- include/gomp-constants.h > > +++ include/gomp-constants.h > > @@ -28,6 +28,9 @@ > > /* Enumerated variable mapping types used to communicate between GCC and > > libgomp. These values are used for both OpenMP and OpenACC. */ > > > > +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2) > > +#define _GOMP_MAP_FLAG_FORCE (1 << 3) > > I'm worried about reserved namespace issues if you use _ followed by > capital letter. Please remind me what those are reserved for? > Can't it be just GOMP_MAP_FLAG_* ? My worry is the other way round: gomp-constants.h is also #included from <openacc.h> (to grab some of its constants), and using plain GOMP_* would pollute the user's namespace? (I'm working on a patch to clean that up, and also use gomp-constants.h more often, also for OpenMP code.) (Such a shared (GCC/libgomp) header files had been discussed before, and now introduced in <http://news.gmane.org/find-root.php?message_id=%3C20140923191931.2177e60f%40octopus%3E>.) Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) 2014-11-13 13:39 ` gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) Thomas Schwinge @ 2014-11-13 13:57 ` Jakub Jelinek 0 siblings, 0 replies; 22+ messages in thread From: Jakub Jelinek @ 2014-11-13 13:57 UTC (permalink / raw) To: Thomas Schwinge; +Cc: gcc-patches On Thu, Nov 13, 2014 at 02:38:06PM +0100, Thomas Schwinge wrote: > Hi Jakub! > > On Thu, 13 Nov 2014 14:10:10 +0100, Jakub Jelinek <jakub@redhat.com> wrote: > > On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote: > > > --- include/gomp-constants.h > > > +++ include/gomp-constants.h > > > @@ -28,6 +28,9 @@ > > > /* Enumerated variable mapping types used to communicate between GCC and > > > libgomp. These values are used for both OpenMP and OpenACC. */ > > > > > > +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2) > > > +#define _GOMP_MAP_FLAG_FORCE (1 << 3) > > > > I'm worried about reserved namespace issues if you use _ followed by > > capital letter. > > Please remind me what those are reserved for? See e.g. http://www.gnu.org/software/libc/manual/html_node/Reserved-Names.html http://pubs.opengroup.org/onlinepubs/007904975/functions/xsh_chap02_02.html and remember that if you use gomp-constants.h in the compiler, it can be built by the system compiler, which can be a very different implementation. > > Can't it be just GOMP_MAP_FLAG_* ? > > My worry is the other way round: gomp-constants.h is also #included from > <openacc.h> (to grab some of its constants), and using plain GOMP_* would > pollute the user's namespace? (I'm working on a patch to clean that up, > and also use gomp-constants.h more often, also for OpenMP code.) (Such a > shared (GCC/libgomp) header files had been discussed before, and now > introduced in > <http://news.gmane.org/find-root.php?message_id=%3C20140923191931.2177e60f%40octopus%3E>.) I think including gomp-constants.h in openacc.h, if that is a publicly installed header, is a bad idea, you'll pollute namespace of that header. Just duplicate the values in there under the right standard required names, and you want, either add a testcase or some static assertions (e.g. of the kind extern char typedef1[condition ? 1 : -1]; in some macros) to verify that the openacc.h constants match the gomp-constants.h where required. Jakub ^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4] Initial support for OpenACC data clauses 2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge 2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas @ 2014-01-28 9:44 ` Thomas Schwinge 2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC deviceptr clause Thomas Schwinge 2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC present data clause Thomas Schwinge 3 siblings, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-01-28 9:44 UTC (permalink / raw) To: gcc-patches; +Cc: jakub [-- Attachment #1: Type: text/plain, Size: 643 bytes --] Hi! On Tue, 14 Jan 2014 07:09:33 -0800, I wrote: > Here is a patch series that adds initial support for OpenACC data > clauses. It is not yet complete, but I thought I might as well already > now strive to get this integrated upstream instead of "hoarding" the > patches locally. Committed to gomp-4_0-branch in r207173..8. > Would it be a good idea to also commit to trunk the (portions of the) > patches that don't directly relate with OpenACC stuff? That way, trunk > and gomp-4_0-branch would diverge a little less? Or, would you first > like to see all of this stabilitize on gomp-4_0-branch? Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* [GOMP4, COMMITTED] OpenACC deviceptr clause. 2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge 2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas 2014-01-28 9:44 ` [gomp4] Initial support for OpenACC data clauses Thomas Schwinge @ 2014-06-05 14:00 ` Thomas Schwinge [not found] ` <5460F49F.3040904@mentor.com> 2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC present data clause Thomas Schwinge 3 siblings, 1 reply; 22+ messages in thread From: Thomas Schwinge @ 2014-06-05 14:00 UTC (permalink / raw) To: gcc-patches From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/c/ * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses): Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR. gcc/ * gimplify.c (gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR. * omp-low.c (scan_sharing_clauses, lower_oacc_offload) (lower_omp_target): Likewise. * tree-core.h (enum omp_clause_map_kind) <OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC deviceptr clause is now supported. * c-c++-common/goacc/deviceptr-1.c: Extend. * c-c++-common/goacc/deviceptr-2.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211278 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 8 +++ gcc/c/ChangeLog.gomp | 5 ++ gcc/c/c-typeck.c | 5 +- gcc/gimplify.c | 7 ++- gcc/omp-low.c | 60 +++++++++++++++++++--- gcc/testsuite/ChangeLog.gomp | 5 ++ .../c-c++-common/goacc/data-clause-duplicate-1.c | 4 +- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 22 +++++++- gcc/testsuite/c-c++-common/goacc/deviceptr-2.c | 23 +++++++++ gcc/tree-core.h | 3 +- 10 files changed, 127 insertions(+), 15 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-2.c diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 7371aa5..88f09b3 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,13 @@ 2014-06-05 Thomas Schwinge <thomas@codesourcery.com> + * gimplify.c (gimplify_scan_omp_clauses) + (gimplify_adjust_omp_clauses): Handle + OMP_CLAUSE_MAP_FORCE_DEVICEPTR. + * omp-low.c (scan_sharing_clauses, lower_oacc_offload) + (lower_omp_target): Likewise. + * tree-core.h (enum omp_clause_map_kind) + <OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment. + * gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>: Don't block OMP_CLAUSE_MAP_FORCE_PRESENT. diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp index 91978db..1e80031 100644 --- gcc/c/ChangeLog.gomp +++ gcc/c/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-06-05 Thomas Schwinge <thomas@codesourcery.com> + + * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses): + Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR. + 2014-03-20 Thomas Schwinge <thomas@codesourcery.com> * c-parser.c: Update comments. diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index c4ba531..839cdf7 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -11747,6 +11747,7 @@ handle_omp_array_sections (tree c) OMP_CLAUSE_SIZE (c) = size; if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) return false; + gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER; if (!c_mark_addressable (t)) @@ -12168,7 +12169,9 @@ c_finish_omp_clauses (tree clauses) else if (!c_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == OMP_CLAUSE_MAP_FORCE_DEVICEPTR))) && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) { error_at (OMP_CLAUSE_LOCATION (c), diff --git gcc/gimplify.c gcc/gimplify.c index 6eaf6fd..a1b6be6 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -6015,7 +6015,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, switch (OMP_CLAUSE_MAP_KIND (c)) { case OMP_CLAUSE_MAP_FORCE_DEALLOC: - case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: input_location = OMP_CLAUSE_LOCATION (c); /* TODO. */ sorry ("data clause not yet implemented"); @@ -6533,6 +6532,12 @@ gimplify_adjust_omp_clauses (tree *list_p) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER) { + /* For OMP_CLAUSE_MAP_FORCE_DEVICEPTR, we'll never enter here, + because for these, TREE_CODE (DECL_SIZE (decl)) will always be + INTEGER_CST. */ + gcc_assert (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); + tree decl2 = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); diff --git gcc/omp-low.c gcc/omp-low.c index 3e282c0..39f0598 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1708,6 +1708,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } +#if 0 + /* In target regions that are not offloaded, libgomp won't pay + attention to OMP_CLAUSE_MAP_FORCE_DEVICEPTR -- but I think we need + to handle it here anyway, in order to create a visible copy of the + variable. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + { + if (!is_gimple_omp_offloaded (ctx->stmt)) + break; + } +#endif if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -1723,6 +1735,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } else { + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -1738,6 +1754,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { tree base = get_base_address (decl); tree nc = OMP_CLAUSE_CHAIN (c); + gcc_assert (nc == NULL_TREE + || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (nc) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)); if (DECL_P (base) && nc != NULL_TREE && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP @@ -1867,6 +1887,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } if (DECL_P (decl)) { + gcc_assert ((OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE); if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE && !COMPLETE_TYPE_P (TREE_TYPE (decl))) @@ -1878,6 +1901,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { + gcc_assert (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR); + tree decl2 = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); @@ -9100,6 +9126,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) { x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -9199,6 +9229,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree var = lookup_decl_in_outer_ctx (ovar, ctx); tree x = build_sender_ref (ovar, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -9219,12 +9253,14 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx) = OMP_CLAUSE_MAP_KIND (c); if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) && (map_kind & OMP_CLAUSE_MAP_TO)) - || map_kind == OMP_CLAUSE_MAP_POINTER) + || map_kind == OMP_CLAUSE_MAP_POINTER + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) gimplify_assign (avar, var, &ilist); avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); - if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) - && (map_kind & OMP_CLAUSE_MAP_FROM)) + if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_FROM)) + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) && !TYPE_READONLY (TREE_TYPE (var))) { x = build_sender_ref (ovar, ctx); @@ -10606,6 +10642,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) @@ -10732,12 +10772,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree var = lookup_decl_in_outer_ctx (ovar, ctx); tree x = build_sender_ref (ovar, ctx); + gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || (OMP_CLAUSE_MAP_KIND (c) + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR) + || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) { - gcc_assert (kind == GF_OMP_TARGET_KIND_REGION); tree avar = create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL); mark_addressable (avar); @@ -10747,19 +10790,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } else if (is_gimple_reg (var)) { - gcc_assert (kind == GF_OMP_TARGET_KIND_REGION); tree avar = create_tmp_var (TREE_TYPE (var), NULL); mark_addressable (avar); enum omp_clause_map_kind map_kind = OMP_CLAUSE_MAP_KIND (c); if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) && (map_kind & OMP_CLAUSE_MAP_TO)) - || map_kind == OMP_CLAUSE_MAP_POINTER) + || map_kind == OMP_CLAUSE_MAP_POINTER + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) gimplify_assign (avar, var, &ilist); avar = build_fold_addr_expr (avar); gimplify_assign (x, avar, &ilist); - if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) - && (map_kind & OMP_CLAUSE_MAP_FROM)) + if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL) + && (map_kind & OMP_CLAUSE_MAP_FROM)) + || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR) && !TYPE_READONLY (TREE_TYPE (var))) { x = build_sender_ref (ovar, ctx); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 4e0ee28..08ec907 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,10 @@ 2014-06-05 Thomas Schwinge <thomas@codesourcery.com> + * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC + deviceptr clause is now supported. + * c-c++-common/goacc/deviceptr-1.c: Extend. + * c-c++-common/goacc/deviceptr-2.c: New file. + * c-c++-common/goacc/data-clause-duplicate-1.c: Extend. * c-c++-common/goacc/present-1.c: New file. diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c index 5c5ab02..7a1cf68 100644 --- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -6,9 +6,7 @@ fun (void) ; #pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */ ; -#pragma acc data create(fp[:10]) deviceptr(fp) - /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */ - /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */ +#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ ; #pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ ; diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c index 1ac63bd..cf2d809 100644 --- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -61,4 +61,24 @@ fun3 (void) ; } -/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */ +extern struct s s1; +extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */ + +void +fun4 (void) +{ + struct s *s1_p = &s1; + struct s *s2_p = &s2; + +#pragma acc parallel deviceptr(s1) /* { dg-error "'s1' is not a pointer variable" } */ + ; + +#pragma acc parallel deviceptr(s2) + ; + +#pragma acc parallel deviceptr(s1_p) + s1_p = 0; + +#pragma acc parallel deviceptr(s2_p) + s2_p = 0; +} diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-2.c gcc/testsuite/c-c++-common/goacc/deviceptr-2.c new file mode 100644 index 0000000..ac162b4 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/deviceptr-2.c @@ -0,0 +1,23 @@ +void +fun1 (void) +{ + char *a = 0; + +#pragma acc data deviceptr(a) + ++a; + +#pragma acc data deviceptr(a) +#pragma acc parallel + ++a; + +#pragma acc data deviceptr(a) +#pragma acc parallel deviceptr(a) + ++a; + +#pragma acc data +#pragma acc parallel deviceptr(a) + ++a; + +#pragma acc parallel deviceptr(a) + ++a; +} diff --git gcc/tree-core.h gcc/tree-core.h index 8603553..8b70c5b 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -1225,7 +1225,8 @@ enum omp_clause_map_kind OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL, /* Deallocate a mapping, without copying from device. */ OMP_CLAUSE_MAP_FORCE_DEALLOC, - /* Is a device pointer. */ + /* Is a device pointer. OMP_CLAUSE_SIZE for these is unused; is implicitly + POINTER_SIZE / BITS_PER_UNIT. */ OMP_CLAUSE_MAP_FORCE_DEVICEPTR, /* End marker. */ -- 1.9.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
[parent not found: <5460F49F.3040904@mentor.com>]
* Re: [GOMP4, COMMITTED] OpenACC deviceptr clause. [not found] ` <5460F49F.3040904@mentor.com> @ 2014-11-11 21:30 ` Thomas Schwinge 0 siblings, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-11-11 21:30 UTC (permalink / raw) To: gcc-patches; +Cc: James Norris, Cesar Philippidis [-- Attachment #1: Type: text/plain, Size: 3292 bytes --] Hi! On Thu, 5 Jun 2014 16:00:16 +0200, I wrote: > --- gcc/omp-low.c > +++ gcc/omp-low.c > @@ -1738,6 +1754,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) > { > tree base = get_base_address (decl); > tree nc = OMP_CLAUSE_CHAIN (c); > + gcc_assert (nc == NULL_TREE > + || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP > + || (OMP_CLAUSE_MAP_KIND (nc) > + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)); > if (DECL_P (base) > && nc != NULL_TREE > && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP That's a bogus assertion; removed in r217372 on gomp-4_0-branch: commit 7ae51786d4a2aad4c82045dda780ae3e7904afa8 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Tue Nov 11 21:22:26 2014 +0000 OpenACC deviceptr clause: Remove bogus assertion. gcc/ * omp-low.c (scan_sharing_clauses): Remove bogus assertion. gcc/testsuite/ * c-c++-common/goacc/deviceptr-3.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217372 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 2 ++ gcc/omp-low.c | 4 ---- gcc/testsuite/ChangeLog.gomp | 5 +++++ gcc/testsuite/c-c++-common/goacc/deviceptr-3.c | 11 +++++++++++ 4 files changed, 18 insertions(+), 4 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 94a7f8c..4ea28e2 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,7 @@ 2014-11-11 Thomas Schwinge <thomas@codesourcery.com> + * omp-low.c (scan_sharing_clauses): Remove bogus assertion. + * omp-low.c (delete_omp_context): Dispose of reduction_map. * omp-low.c (maybe_lookup_reduction): Don't require an OpenACC diff --git gcc/omp-low.c gcc/omp-low.c index 5695ec3..1263409 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1920,10 +1920,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { tree base = get_base_address (decl); tree nc = OMP_CLAUSE_CHAIN (c); - gcc_assert (nc == NULL_TREE - || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP - || (OMP_CLAUSE_MAP_KIND (nc) - != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)); if (DECL_P (base) && nc != NULL_TREE && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index a02f58a..f8bacc3 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-11-07 Thomas Schwinge <thomas@codesourcery.com> + James Norris <jnorris@codesourcery.com> + + * c-c++-common/goacc/deviceptr-3.c: New file. + 2014-11-05 Thomas Schwinge <thomas@codesourcery.com> * c-c++-common/goacc/update-1.c: Extend. diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-3.c gcc/testsuite/c-c++-common/goacc/deviceptr-3.c new file mode 100644 index 0000000..bab56c3 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/deviceptr-3.c @@ -0,0 +1,11 @@ +float *d_a; + +void +f (float *a) +{ +#pragma acc parallel copyout (a[3:10]) deviceptr (d_a) + d_a[2] += 1.0; + +#pragma acc parallel deviceptr (d_a) copyout (a[3:10]) + d_a[2] += 1.0; +} Grüße, Thomas [-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --] ^ permalink raw reply [flat|nested] 22+ messages in thread
* [GOMP4, COMMITTED] OpenACC present data clause. 2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge ` (2 preceding siblings ...) 2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC deviceptr clause Thomas Schwinge @ 2014-06-05 14:00 ` Thomas Schwinge 3 siblings, 0 replies; 22+ messages in thread From: Thomas Schwinge @ 2014-06-05 14:00 UTC (permalink / raw) To: gcc-patches From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> gcc/ * gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>: Don't block OMP_CLAUSE_MAP_FORCE_PRESENT. gcc/testsuite/ * c-c++-common/goacc/data-clause-duplicate-1.c: Extend. * c-c++-common/goacc/present-1.c: New file. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211277 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 5 +++++ gcc/gimplify.c | 1 - gcc/testsuite/ChangeLog.gomp | 5 +++++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c | 2 ++ gcc/testsuite/c-c++-common/goacc/present-1.c | 11 +++++++++++ 5 files changed, 23 insertions(+), 1 deletion(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/present-1.c diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 011fe77..7371aa5 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-06-05 Thomas Schwinge <thomas@codesourcery.com> + + * gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>: + Don't block OMP_CLAUSE_MAP_FORCE_PRESENT. + 2014-06-04 Thomas Schwinge <thomas@codesourcery.com> * cgraphunit.c (ipa_passes, compile): Handle flag_openacc next to diff --git gcc/gimplify.c gcc/gimplify.c index e98e6e5..6eaf6fd 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -6014,7 +6014,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_MAP: switch (OMP_CLAUSE_MAP_KIND (c)) { - case OMP_CLAUSE_MAP_FORCE_PRESENT: case OMP_CLAUSE_MAP_FORCE_DEALLOC: case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: input_location = OMP_CLAUSE_LOCATION (c); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index 78882c0..4e0ee28 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-06-05 Thomas Schwinge <thomas@codesourcery.com> + + * c-c++-common/goacc/data-clause-duplicate-1.c: Extend. + * c-c++-common/goacc/present-1.c: New file. + 2014-03-20 Thomas Schwinge <thomas@codesourcery.com> * c-c++-common/goacc-gomp/nesting-1.c: New file. diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c index 4cb3cc2..5c5ab02 100644 --- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c +++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c @@ -10,4 +10,6 @@ fun (void) /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */ /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */ ; +#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */ + ; } diff --git gcc/testsuite/c-c++-common/goacc/present-1.c gcc/testsuite/c-c++-common/goacc/present-1.c new file mode 100644 index 0000000..03ee592 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/present-1.c @@ -0,0 +1,11 @@ +/* { dg-additional-options "-fdump-tree-original" } */ + +void +f (char *cp) +{ +#pragma acc parallel present(cp[7:9]) + ; +} + +/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */ +/* { dg-final { cleanup-tree-dump "original" } } */ -- 1.9.1 ^ permalink raw reply [flat|nested] 22+ messages in thread
end of thread, other threads:[~2014-11-13 13:54 UTC | newest]
Thread overview: 22+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
2014-01-14 15:10 ` [gomp4 2/6] Prepare for extending omp_clause_map_kind thomas
2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas
2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas
2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas
2014-01-14 15:10 ` [gomp4 6/6] Enable initial " thomas
2014-02-12 11:17 ` [gomp4 5/6] Initial " Thomas Schwinge
2014-02-21 19:48 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
2014-02-21 20:32 ` [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE Thomas Schwinge
2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge
2014-03-12 13:48 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-03-20 14:39 ` [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) Thomas Schwinge
[not found] ` <538DF785.3050206@mentor.com>
[not found] ` <87egz645j4.fsf@schwinge.name>
2014-11-13 12:21 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
2014-11-13 13:13 ` Jakub Jelinek
2014-11-13 13:39 ` gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) Thomas Schwinge
2014-11-13 13:57 ` Jakub Jelinek
2014-01-28 9:44 ` [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC deviceptr clause Thomas Schwinge
[not found] ` <5460F49F.3040904@mentor.com>
2014-11-11 21:30 ` Thomas Schwinge
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC present data clause 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).