From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 93103 invoked by alias); 29 Jul 2015 17:19:33 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 93092 invoked by uid 89); 29 Jul 2015 17:19:32 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.7 required=5.0 tests=AWL,BAYES_50,KAM_LAZY_DOMAIN_SECURITY,RP_MATCHES_RCVD,SPF_HELO_PASS autolearn=ham version=3.3.2 X-HELO: mx1.redhat.com Received: from mx1.redhat.com (HELO mx1.redhat.com) (209.132.183.28) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-GCM-SHA384 encrypted) ESMTPS; Wed, 29 Jul 2015 17:19:15 +0000 Received: from int-mx13.intmail.prod.int.phx2.redhat.com (int-mx13.intmail.prod.int.phx2.redhat.com [10.5.11.26]) by mx1.redhat.com (Postfix) with ESMTPS id CFE55461D9; Wed, 29 Jul 2015 17:19:13 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-30.ams2.redhat.com [10.36.116.30]) by int-mx13.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t6THJBuW027142 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Wed, 29 Jul 2015 13:19:12 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.14.9/8.14.9) with ESMTP id t6THJ9eE020974; Wed, 29 Jul 2015 19:19:09 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.14.9/8.14.9/Submit) id t6THJ7CV020973; Wed, 29 Jul 2015 19:19:07 +0200 Date: Wed, 29 Jul 2015 17:30:00 -0000 From: Jakub Jelinek To: Ilya Verbin Cc: Thomas Schwinge , gcc-patches@gcc.gnu.org, Kirill Yukhin Subject: [gomp4.1] Various accelerator updates from OpenMP 4.1 Message-ID: <20150729171907.GI1780@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20150609183608.GA47936@msticlxl57.ims.intel.com> <20150609202426.GG10247@tucnak.redhat.com> <20150625194529.GB33078@msticlxl57.ims.intel.com> <20150625201058.GK10247@tucnak.redhat.com> <20150717163136.GB15252@msticlxl57.ims.intel.com> <20150717164306.GT1780@tucnak.redhat.com> <20150720161422.GC1780@tucnak.redhat.com> <20150720181041.GE1780@tucnak.redhat.com> <20150722211348.GA1750@tucnak.redhat.com> <20150724200457.GB1750@tucnak.redhat.com> MIME-Version: 1.0 Content-Type: text/plain; charset=us-ascii Content-Disposition: inline In-Reply-To: <20150724200457.GB1750@tucnak.redhat.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes X-SW-Source: 2015-07/txt/msg02488.txt.bz2 On Fri, Jul 24, 2015 at 10:04:57PM +0200, Jakub Jelinek wrote: > Another version. > What to do with zero-length array sections vs. objects is still under heated > debates, so target8.f90 keeps failing intermittently. Here is a new version of the patch, with various additions (implemented GOMP_MAP_FIRSTPRIVATE_INT I've talked about, it now handles use_device_ptr and handles is_device_ptr with array decls (silly, but seems the accel folks want it for some strange reason), etc.) and it special cases zero length array sections rather than all zero length mappings. The heated debates continue, so perhaps that part - GOMP_MAP_ZERO_LEN_ARRAY_SECTION - will need reversion and replacement with something else, we'll see. This let's the testsuite pass for now except for the two LTO ICEs, both without offloading (host fallback only) and with Intel MIC offloading. Committed to gomp-4_1-branch. Ilya, I think now is the time to update your enter data/exit data patch. 2015-07-29 Jakub Jelinek gcc/ * tree.h (OMP_TARGET_COMBINED): Define. (OMP_CLAUSE_SET_MAP_KIND): Cast to unsigned int rather than unsigned char. (OMP_CLAUSE_MAP_PRIVATE, OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION): Define. * tree-core.h (struct tree_omp_clause): Change type of map_kind from unsigned char to unsigned int. * gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_0LEN_ARRAY. (enum omp_region_type): Add ORT_COMBINED_TARGET. (struct gimplify_omp_ctx): Add target_map_scalars_firstprivate, target_map_pointers_as_0len_arrays and target_firstprivatize_array_bases fields. (maybe_fold_stmt): Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. (omp_notice_threadprivate_variable): Likewise. (omp_firstprivatize_variable): Likewise. If ctx->target_map_scalars_firstprivate is set, firstprivatize as GOVD_FIRSTPRIVATE. (omp_add_variable): Allow map clause together with data sharing clauses. For data sharing clause with VLA decl on omp target/target data don't add firstprivate for the pointer. (omp_notice_variable): Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. Handle implicit mapping of pointers as zero length array sections and ctx->target_map_scalars_firstprivate mapping of scalars as firstprivate data sharing. (gimplify_scan_omp_clauses): Initialize ctx->target_map_scalars_firstprivate, ctx->target_firstprivatize_array_bases and ctx->target_map_pointers_as_0len_arrays. Add firstprivate for linear clause even to target region if combined. Remove map clauses with GOMP_MAP_FIRSTPRIVATE_POINTER kind from OMP_TARGET_{,ENTER_,EXIT_}DATA. For GOMP_MAP_FIRSTPRIVATE_POINTER map kind with non-INTEGER_CST OMP_CLAUSE_SIZE firstprivatize the bias. (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_0LEN_ARRAY. If gimplify_omp_ctxp->target_firstprivatize_array_bases, use GOMP_MAP_FIRSTPRIVATE_POINTER map kind instead of GOMP_MAP_POINTER. (gimplify_adjust_omp_clauses): Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. Use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER if ctx->target_firstprivatize_array_bases for VLAs. Set OMP_CLAUSE_MAP_PRIVATE if both data sharing and map clause appear together. (gimplify_omp_workshare): Adjust check for ORT_TARGET for the addition of ORT_COMBINED_TARGET. Use ORT_COMBINED_TARGET if OMP_TARGET_COMBINED. * omp-low.c (lookup_sfield): Change first argument to splay_tree_key, add overload with tree first argument. (maybe_lookup_field): Likewise. (build_sender_ref): Likewise. (scan_sharing_clauses): Handle VLAs in target firstprivate and is_device_ptr clauses. Fix up variable shadowing. Handle OMP_CLAUSE_USE_DEVICE_PTR. Handle OMP_CLAUSE_MAP_PRIVATE. Handle GOMP_MAP_FIRSTPRIVATE_POINTER map kind. (handle_simd_reference): Use get_name. (lower_rec_input_clauses): Likewise. Use BUILT_IN_ALLOCA_WITH_ALIGN instead of BUILT_IN_ALLOCA. (lower_send_clauses): Use new lookup_sfield overload. (lower_omp_target): Handle GOMP_MAP_FIRSTPRIVATE_POINTER map kind. Handle OMP_CLAUSE_PRIVATE VLAs. Handle OMP_CLAUSE_USE_DEVICE_PTR, handle arrays and references to arrays in OMP_CLAUSE_IS_DEVICE_PTR clause. Handle OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION. * tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_FIRSTPRIVATE_POINTER. gcc/c/ * c-tree.h (c_finish_omp_clauses): Add is_omp argument. * c-parser.c (c_parser_oacc_all_clauses, c_parser_omp_all_clauses, c_parser_oacc_cache, omp_split_clauses, c_parser_cilk_for): Adjust c_finish_omp_clauses callers. (c_parser_omp_target_data, c_parser_omp_target_enter_data, c_parser_omp_target_exit_data): Disallow GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER but don't set map_seen for it. (c_parser_omp_target): Set OMP_TARGET_COMBINED if combined. Disallow GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER. * c-typeck.c (handle_omp_array_sections): Add is_omp argument. Set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if needed. Use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER if is_omp. (c_finish_omp_clauses): Add is_omp argument, pass it down to handle_omp_array_sections. Handle GOMP_MAP_FIRSTPRIVATE_POINTER. For is_device_ptr/use_device_ptr clauses allow ARRAY_TYPE. gcc/cp/ * parser.c (cp_parser_omp_target_data, cp_parser_omp_target_enter_data, cp_parser_omp_target_exit_data): Formatting fixes. Disallow GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER but don't set map_seen for it. (cp_parser_omp_target): Set OMP_TARGET_COMBINED if combined. Disallow GOMP_MAP_POINTER, allow GOMP_MAP_FIRSTPRIVATE_POINTER. * semantics.c (handle_omp_array_sections): Add is_omp argument. Set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION if needed. Use GOMP_MAP_FIRSTPRIVATE_POINTER instead of GOMP_MAP_POINTER if is_omp. (finish_omp_clauses): Handle GOMP_MAP_FIRSTPRIVATE_POINTER. For is_device_ptr/use_device_ptr clauses allow ARRAY_TYPE and REFERENCE_TYPE to ARRAY_TYPE. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_USE_DEVICE_PTR, GOMP_MAP_ZERO_LEN_ARRAY_SECTION and GOMP_MAP_FIRSTPRIVATE_POINTER. libgomp/ * libgomp.h (struct target_var_desc): Fix up comments about offset and length fields. * target.c (gomp_map_lookup): New function. (gomp_map_pointer): Use it. (gomp_map_vars): Handle GOMP_MAP_FIRSTPRIVATE_INT, GOMP_MAP_USE_DEVICE_PTR and GOMP_MAP_ZERO_LEN_ARRAY_SECTION. Add tgt->list[i].offset for mappings with non-NULL tgt->list[i].key. (GOMP_target_41): Handle GOMP_MAP_FIRSTPRIVATE even for host fallback. (omp_target_is_present): Use gomp_map_lookup. (omp_target_associate_ptr): Likewise. (omp_target_disassociate_ptr): Likewise. * testsuite/libgomp.c++/target-2.C (fn2): Add map(tofrom: s). * testsuite/libgomp.c++/target-7.C: New test. * testsuite/libgomp.c++/target-8.C: New test. * testsuite/libgomp.c++/target-9.C: New test. * testsuite/libgomp.c/target-1.c (fn2, fn3, fn4): Add map(tofrom:s). * testsuite/libgomp.c/target-2.c (fn2, fn3, fn4): Likewise. * testsuite/libgomp.c/target-7.c (foo): Add map(h) where needed. * testsuite/libgomp.c/target-15.c: New test. * testsuite/libgomp.c/target-16.c: New test. * testsuite/libgomp.c/target-17.c: New test. * testsuite/libgomp.c/target-18.c: New test. * testsuite/libgomp.c/target-19.c: New test. * testsuite/libgomp.c/examples-4/e.51.3.c (gramSchmidt): Add map(tofrom:tmp). * testsuite/libgomp.c/examples-4/e.53.1.c (fib_wrapper): Add map(from:x). * testsuite/libgomp.c/examples-4/e.53.4.c (accum): Add map(tofrom:tmp). * testsuite/libgomp.c/examples-4/e.53.5.c (accum): Likewise. * testsuite/libgomp.c/examples-4/e.54.2.c (dotprod): Add map(tofrom: sum). * testsuite/libgomp.c/examples-4/e.54.3.c (dotprod): Likewise. * testsuite/libgomp.c/examples-4/e.54.4.c (dotprod): Likewise. * testsuite/libgomp.c/examples-4/e.57.1.c (main): Add map(from: c) and map(from: b, d) where needed. * testsuite/libgomp.c/examples-4/e.57.3.c (main): Add map(from: res). --- gcc/tree.h.jj 2015-07-16 17:56:41.000000000 +0200 +++ gcc/tree.h 2015-07-29 14:13:26.336307751 +0200 @@ -1341,6 +1341,11 @@ extern void protected_set_expr_location #define OMP_TEAMS_COMBINED(NODE) \ (OMP_TEAMS_CHECK (NODE)->base.private_flag) +/* True on an OMP_TARGET statement if it represents explicit + combined target teams, target parallel or target simd constructs. */ +#define OMP_TARGET_COMBINED(NODE) \ + (OMP_TARGET_CHECK (NODE)->base.private_flag) + /* True if OMP_ATOMIC* is supposed to be sequentially consistent as opposed to relaxed. */ #define OMP_ATOMIC_SEQ_CST(NODE) \ @@ -1445,13 +1450,21 @@ extern void protected_set_expr_location ((enum gomp_map_kind) OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind) #define OMP_CLAUSE_SET_MAP_KIND(NODE, MAP_KIND) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \ - = (unsigned char) (MAP_KIND)) + = (unsigned int) (MAP_KIND)) /* Nonzero if this map clause is for array (rather than pointer) based array section with zero bias. Both the non-decl OMP_CLAUSE_MAP and corresponding OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag. */ #define OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.public_flag) +/* Nonzero if the same decl appears both in OMP_CLAUSE_MAP and either + OMP_CLAUSE_PRIVATE or OMP_CLAUSE_FIRSTPRIVATE. */ +#define OMP_CLAUSE_MAP_PRIVATE(NODE) \ + TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) +/* Nonzero if this is a mapped array section, that might need special + treatment if OMP_CLAUSE_SIZE is zero. */ +#define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \ + TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) #define OMP_CLAUSE_PROC_BIND_KIND(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind) --- gcc/tree-core.h.jj 2015-07-17 09:30:44.000000000 +0200 +++ gcc/tree-core.h 2015-07-21 16:28:48.524156167 +0200 @@ -1354,7 +1354,7 @@ struct GTY(()) tree_omp_clause { enum omp_clause_schedule_kind schedule_kind; enum omp_clause_depend_kind depend_kind; /* See include/gomp-constants.h for enum gomp_map_kind's values. */ - unsigned char map_kind; + unsigned int map_kind; enum omp_clause_proc_bind_kind proc_bind_kind; enum tree_code reduction_code; enum omp_clause_linear_kind linear_kind; --- gcc/gimplify.c.jj 2015-07-16 17:56:41.000000000 +0200 +++ gcc/gimplify.c 2015-07-29 16:43:57.056823518 +0200 @@ -90,6 +90,8 @@ enum gimplify_omp_var_data /* Flag for GOVD_LINEAR or GOVD_LASTPRIVATE: no outer reference. */ GOVD_LINEAR_LASTPRIVATE_NO_OUTER = 16384, + GOVD_MAP_0LEN_ARRAY = 32768, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -110,6 +112,7 @@ enum omp_region_type ORT_TARGET_DATA = 16, /* Data region with offloading. */ ORT_TARGET = 32, + ORT_COMBINED_TARGET = 33, /* Dummy OpenMP region, used to disable expansion of DECL_VALUE_EXPRs in taskloop pre body. */ ORT_NONE = 64 @@ -156,6 +159,9 @@ struct gimplify_omp_ctx enum omp_region_type region_type; bool combined_loop; bool distribute; + bool target_map_scalars_firstprivate; + bool target_map_pointers_as_0len_arrays; + bool target_firstprivatize_array_bases; }; static struct gimplify_ctx *gimplify_ctxp; @@ -2260,7 +2266,7 @@ maybe_fold_stmt (gimple_stmt_iterator *g { 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) != 0) return false; return fold_stmt (gsi); } @@ -5561,8 +5567,13 @@ omp_firstprivatize_variable (struct gimp else return; } - else if (ctx->region_type == ORT_TARGET) - omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY); + else if ((ctx->region_type & ORT_TARGET) != 0) + { + if (ctx->target_map_scalars_firstprivate) + omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE); + else + 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) @@ -5648,7 +5659,7 @@ omp_add_variable (struct gimplify_omp_ct flags |= GOVD_SEEN; n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); - if (n != NULL && n->value != GOVD_ALIGNED) + if (n != NULL && (n->value & GOVD_DATA_SHARE_CLASS) != 0) { /* We shouldn't be re-adding the decl with the same data sharing class. */ @@ -5678,6 +5689,9 @@ omp_add_variable (struct gimplify_omp_ct nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT; else if (flags & GOVD_PRIVATE) nflags = GOVD_PRIVATE; + else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0 + && (flags & GOVD_FIRSTPRIVATE)) + nflags = GOVD_PRIVATE | GOVD_EXPLICIT; else nflags = GOVD_FIRSTPRIVATE; nflags |= flags & GOVD_SEEN; @@ -5746,7 +5760,7 @@ omp_notice_threadprivate_variable (struc 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) != 0) { n = splay_tree_lookup (octx->variables, (splay_tree_key)decl); if (n == NULL) @@ -5810,19 +5824,66 @@ omp_notice_variable (struct gimplify_omp } n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); - if (ctx->region_type == ORT_TARGET) + if ((ctx->region_type & ORT_TARGET) != 0) { ret = lang_hooks.decls.omp_disregard_value_expr (decl, true); if (n == NULL) { - if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl))) + unsigned nflags = flags; + if (ctx->target_map_pointers_as_0len_arrays + || ctx->target_map_scalars_firstprivate) + { + bool is_declare_target = false; + bool is_scalar = false; + if (is_global_var (decl) + && varpool_node::get_create (decl)->offloadable) + { + struct gimplify_omp_ctx *octx; + for (octx = ctx->outer_context; + octx; octx = octx->outer_context) + { + n = splay_tree_lookup (octx->variables, + (splay_tree_key)decl); + if (n + && (n->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED + && (n->value & GOVD_DATA_SHARE_CLASS) != 0) + break; + } + is_declare_target = octx == NULL; + } + if (!is_declare_target && ctx->target_map_scalars_firstprivate) + { + tree type = TREE_TYPE (decl); + if (TREE_CODE (type) == REFERENCE_TYPE) + type = TREE_TYPE (type); + if (TREE_CODE (type) == COMPLEX_TYPE) + type = TREE_TYPE (type); + if (INTEGRAL_TYPE_P (type) + || SCALAR_FLOAT_TYPE_P (type) + || TREE_CODE (type) == POINTER_TYPE) + is_scalar = true; + } + if (is_declare_target) + ; + else if (ctx->target_map_pointers_as_0len_arrays + && (TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE + || (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) + == POINTER_TYPE))) + nflags |= GOVD_MAP | GOVD_MAP_0LEN_ARRAY; + else if (is_scalar) + nflags |= GOVD_FIRSTPRIVATE; + } + if (nflags == flags + && !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); + nflags |= GOVD_MAP | GOVD_EXPLICIT; } - else - omp_add_variable (ctx, decl, GOVD_MAP | flags); + else if (nflags == flags) + nflags |= GOVD_MAP; + omp_add_variable (ctx, decl, nflags); } else { @@ -6144,6 +6205,24 @@ gimplify_scan_omp_clauses (tree *list_p, ctx = new_omp_context (region_type); outer_ctx = ctx->outer_context; + if (code == OMP_TARGET && !lang_GNU_Fortran ()) + { + ctx->target_map_pointers_as_0len_arrays = true; + /* FIXME: For Fortran we want to set this too, when + the Fortran FE is updated to OpenMP 4.1. */ + ctx->target_map_scalars_firstprivate = true; + } + if (!lang_GNU_Fortran ()) + switch (code) + { + case OMP_TARGET: + case OMP_TARGET_DATA: + case OMP_TARGET_ENTER_DATA: + case OMP_TARGET_EXIT_DATA: + ctx->target_firstprivatize_array_bases = true; + default: + break; + } while ((c = *list_p) != NULL) { @@ -6290,11 +6369,18 @@ gimplify_scan_omp_clauses (tree *list_p, && ctx->region_type == ORT_WORKSHARE && octx == outer_ctx) flags = GOVD_SEEN | GOVD_SHARED; + else if (octx + && octx->region_type == ORT_COMBINED_TARGET) + flags &= ~GOVD_LASTPRIVATE; else break; - gcc_checking_assert (splay_tree_lookup (octx->variables, - (splay_tree_key) - decl) == NULL); + splay_tree_node on + = splay_tree_lookup (octx->variables, + (splay_tree_key) decl); + gcc_assert (on == NULL + || (octx->region_type == ORT_COMBINED_TARGET + && (on->value + & GOVD_DATA_SHARE_CLASS) == 0)); omp_add_variable (octx, decl, flags); if (octx->outer_context == NULL) break; @@ -6319,10 +6405,24 @@ gimplify_scan_omp_clauses (tree *list_p, case OMP_CLAUSE_MAP: decl = OMP_CLAUSE_DECL (c); if (error_operand_p (decl)) + remove = true; + switch (code) { - remove = true; + case OMP_TARGET: + break; + case OMP_TARGET_DATA: + case OMP_TARGET_ENTER_DATA: + case OMP_TARGET_EXIT_DATA: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + /* For target {,enter ,exit }data only the array slice is + mapped, but not the pointer to it. */ + remove = true; + break; + default: break; } + if (remove) + break; if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); @@ -6332,6 +6432,14 @@ gimplify_scan_omp_clauses (tree *list_p, remove = true; break; } + else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER + && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) + { + OMP_CLAUSE_SIZE (c) + = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL); + omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), + GOVD_FIRSTPRIVATE | GOVD_SEEN); + } if (!DECL_P (decl)) { if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, @@ -6643,7 +6751,10 @@ gimplify_scan_omp_clauses (tree *list_p, case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_THREADS: case OMP_CLAUSE_SIMD: + break; + case OMP_CLAUSE_DEFAULTMAP: + ctx->target_map_scalars_firstprivate = false; break; case OMP_CLAUSE_ALIGNED: @@ -6759,6 +6870,30 @@ gimplify_adjust_omp_clauses_1 (splay_tre OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1; else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; + else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0) + { + tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = decl; + if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE + && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == POINTER_TYPE) + OMP_CLAUSE_DECL (clause) + = build_simple_mem_ref_loc (input_location, decl); + OMP_CLAUSE_DECL (clause) + = build2 (MEM_REF, char_type_node, OMP_CLAUSE_DECL (clause), + build_int_cst (build_pointer_type (char_type_node), 0)); + OMP_CLAUSE_SIZE (clause) = size_zero_node; + OMP_CLAUSE_SIZE (nc) = size_zero_node; + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC); + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1; + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER); + OMP_CLAUSE_CHAIN (nc) = *list_p; + OMP_CLAUSE_CHAIN (clause) = nc; + struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; + gimplify_omp_ctxp = ctx->outer_context; + gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0), + pre_p, NULL, is_gimple_val, fb_rvalue); + gimplify_omp_ctxp = ctx; + } else if (code == OMP_CLAUSE_MAP) { OMP_CLAUSE_SET_MAP_KIND (clause, @@ -6785,7 +6920,10 @@ gimplify_adjust_omp_clauses_1 (splay_tre OMP_CLAUSE_MAP); OMP_CLAUSE_DECL (nc) = decl; OMP_CLAUSE_SIZE (nc) = size_zero_node; - OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER); + if (gimplify_omp_ctxp->target_firstprivatize_array_bases) + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER); + else + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER); OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause); OMP_CLAUSE_CHAIN (clause) = nc; } @@ -6910,12 +7048,14 @@ gimplify_adjust_omp_clauses (gimple_seq 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) != 0 + && !(n->value & GOVD_SEEN) && !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)) remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER) { /* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because for these, TREE_CODE (DECL_SIZE (decl)) will always be @@ -6935,17 +7075,33 @@ gimplify_adjust_omp_clauses (gimple_seq omp_notice_variable (ctx->outer_context, OMP_CLAUSE_SIZE (c), true); } - tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), - OMP_CLAUSE_MAP); - OMP_CLAUSE_DECL (nc) = decl; - OMP_CLAUSE_SIZE (nc) = size_zero_node; - OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER); - OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); - OMP_CLAUSE_CHAIN (c) = nc; - c = nc; + if (((ctx->region_type & ORT_TARGET) != 0 + || !ctx->target_firstprivatize_array_bases) + && ((n->value & GOVD_SEEN) == 0 + || (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) == 0)) + { + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (c), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (nc) = decl; + OMP_CLAUSE_SIZE (nc) = size_zero_node; + if (ctx->target_firstprivatize_array_bases) + OMP_CLAUSE_SET_MAP_KIND (nc, + GOMP_MAP_FIRSTPRIVATE_POINTER); + else + OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER); + OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); + OMP_CLAUSE_CHAIN (c) = nc; + c = nc; + } + } + else + { + if (OMP_CLAUSE_SIZE (c) == NULL_TREE) + OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); + if ((n->value & GOVD_SEEN) + && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))) + OMP_CLAUSE_MAP_PRIVATE (c) = 1; } - else if (OMP_CLAUSE_SIZE (c) == NULL_TREE) - OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); break; case OMP_CLAUSE_TO: @@ -7888,9 +8044,11 @@ gimplify_omp_workshare (tree *expr_p, gi case OMP_SINGLE: ort = ORT_WORKSHARE; break; + case OMP_TARGET: + ort = OMP_TARGET_COMBINED (expr) ? ORT_COMBINED_TARGET : ORT_TARGET; + break; case OACC_KERNELS: case OACC_PARALLEL: - case OMP_TARGET: ort = ORT_TARGET; break; case OACC_DATA: @@ -7905,7 +8063,7 @@ gimplify_omp_workshare (tree *expr_p, gi } gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort, TREE_CODE (expr)); - if (ort == ORT_TARGET || ort == ORT_TARGET_DATA) + if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0) { push_gimplify_context (); gimple g = gimplify_and_return_first (OMP_BODY (expr), &body); --- gcc/omp-low.c.jj 2015-07-21 09:07:23.000000000 +0200 +++ gcc/omp-low.c 2015-07-29 16:13:33.209580272 +0200 @@ -1071,24 +1071,35 @@ lookup_field (tree var, omp_context *ctx } static inline tree -lookup_sfield (tree var, omp_context *ctx) +lookup_sfield (splay_tree_key key, omp_context *ctx) { splay_tree_node n; n = splay_tree_lookup (ctx->sfield_map - ? ctx->sfield_map : ctx->field_map, - (splay_tree_key) var); + ? ctx->sfield_map : ctx->field_map, key); return (tree) n->value; } static inline tree -maybe_lookup_field (tree var, omp_context *ctx) +lookup_sfield (tree var, omp_context *ctx) +{ + return lookup_sfield ((splay_tree_key) var, ctx); +} + +static inline tree +maybe_lookup_field (splay_tree_key key, omp_context *ctx) { splay_tree_node n; - n = splay_tree_lookup (ctx->field_map, (splay_tree_key) var); + n = splay_tree_lookup (ctx->field_map, key); return n ? (tree) n->value : NULL_TREE; } static inline tree +maybe_lookup_field (tree var, omp_context *ctx) +{ + return maybe_lookup_field ((splay_tree_key) var, ctx); +} + +static inline tree lookup_oacc_reduction (const char *id, omp_context *ctx) { splay_tree_node n; @@ -1359,12 +1370,18 @@ build_outer_var_ref (tree var, omp_conte /* Build tree nodes to access the field for VAR on the sender side. */ static tree -build_sender_ref (tree var, omp_context *ctx) +build_sender_ref (splay_tree_key key, omp_context *ctx) { - tree field = lookup_sfield (var, ctx); + tree field = lookup_sfield (key, ctx); return omp_build_component_ref (ctx->sender_decl, field); } +static tree +build_sender_ref (tree var, omp_context *ctx) +{ + return build_sender_ref ((splay_tree_key) var, ctx); +} + /* Add a new field for VAR inside the structure CTX->SENDER_DECL. */ static void @@ -1908,6 +1925,17 @@ scan_sharing_clauses (tree clauses, omp_ case OMP_CLAUSE_LINEAR: decl = OMP_CLAUSE_DECL (c); do_private: + if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + && is_gimple_omp_offloaded (ctx->stmt)) + { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) + install_var_field (decl, !is_reference (decl), 3, ctx); + else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + install_var_field (decl, true, 3, ctx); + else + install_var_field (decl, false, 3, ctx); + } if (is_variable_sized (decl)) { if (is_task_ctx (ctx)) @@ -1930,10 +1958,6 @@ scan_sharing_clauses (tree clauses, omp_ else if (!global) install_var_field (decl, by_ref, 3, ctx); } - else if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) - && is_gimple_omp_offloaded (ctx->stmt)) - install_var_field (decl, !is_reference (decl), 3, ctx); install_var_local (decl, ctx); if (is_gimple_omp_oacc (ctx->stmt) && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION) @@ -1944,9 +1968,9 @@ scan_sharing_clauses (tree clauses, omp_ tree ptype = build_pointer_type (type); tree array = create_tmp_var (ptype, oacc_get_reduction_array_id (var)); - omp_context *c = (ctx->field_map ? ctx : ctx->outer); - install_var_field (array, true, 3, c); - install_var_local (array, c); + omp_context *octx = (ctx->field_map ? ctx : ctx->outer); + install_var_field (array, true, 3, octx); + install_var_local (array, octx); /* Insert it into the current context. */ splay_tree_insert (ctx->reduction_map, (splay_tree_key) @@ -1959,6 +1983,23 @@ scan_sharing_clauses (tree clauses, omp_ break; case OMP_CLAUSE_USE_DEVICE_PTR: + decl = OMP_CLAUSE_DECL (c); + if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + install_var_field (decl, true, 3, ctx); + else + install_var_field (decl, false, 3, ctx); + if (DECL_SIZE (decl) + && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) + { + tree decl2 = DECL_VALUE_EXPR (decl); + gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + install_var_local (decl2, ctx); + } + install_var_local (decl, ctx); + break; + case OMP_CLAUSE_IS_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); goto do_private; @@ -2025,6 +2066,21 @@ scan_sharing_clauses (tree clauses, omp_ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)) break; } + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + { + if (DECL_SIZE (decl) + && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) + { + tree decl2 = DECL_VALUE_EXPR (decl); + gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + install_var_local (decl2, ctx); + } + install_var_local (decl, ctx); + break; + } if (DECL_P (decl)) { if (DECL_SIZE (decl) @@ -2034,7 +2090,11 @@ scan_sharing_clauses (tree clauses, omp_ gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); decl2 = TREE_OPERAND (decl2, 0); gcc_assert (DECL_P (decl2)); - install_var_field (decl2, true, 3, ctx); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_PRIVATE (c)) + install_var_field (decl2, true, 11, ctx); + else + install_var_field (decl2, true, 3, ctx); install_var_local (decl2, ctx); install_var_local (decl, ctx); } @@ -2045,6 +2105,9 @@ scan_sharing_clauses (tree clauses, omp_ && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) install_var_field (decl, true, 7, ctx); + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_PRIVATE (c)) + install_var_field (decl, true, 11, ctx); else install_var_field (decl, true, 3, ctx); if (is_gimple_omp_offloaded (ctx->stmt)) @@ -2147,11 +2210,23 @@ scan_sharing_clauses (tree clauses, omp_ /* FALLTHRU */ case OMP_CLAUSE_PRIVATE: case OMP_CLAUSE_LINEAR: - case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); if (is_variable_sized (decl)) - install_var_local (decl, ctx); + { + if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR) + && is_gimple_omp_offloaded (ctx->stmt)) + { + tree decl2 = DECL_VALUE_EXPR (decl); + gcc_assert (TREE_CODE (decl2) == INDIRECT_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + install_var_local (decl2, ctx); + fixup_remapped_decl (decl2, ctx, false); + } + install_var_local (decl, ctx); + } fixup_remapped_decl (decl, ctx, OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE && OMP_CLAUSE_PRIVATE_DEBUG (c)); @@ -2201,7 +2276,8 @@ scan_sharing_clauses (tree clauses, omp_ break; if (DECL_P (decl)) { - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE && !COMPLETE_TYPE_P (TREE_TYPE (decl))) { @@ -2255,6 +2331,7 @@ scan_sharing_clauses (tree clauses, omp_ case OMP_CLAUSE_SIMD: case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_DEFAULTMAP: + case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: @@ -3924,11 +4001,8 @@ handle_simd_reference (location_t loc, t tree z = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_vard))); if (TREE_CONSTANT (z)) { - const char *name = NULL; - if (DECL_NAME (new_vard)) - name = IDENTIFIER_POINTER (DECL_NAME (new_vard)); - - z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)), name); + z = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_vard)), + get_name (new_vard)); gimple_add_tmp_var (z); TREE_ADDRESSABLE (z) = 1; z = build_fold_addr_expr_loc (loc, z); @@ -4127,9 +4201,7 @@ lower_rec_input_clauses (tree clauses, g tree type = TREE_TYPE (d); gcc_assert (TREE_CODE (type) == ARRAY_TYPE); tree v = TYPE_MAX_VALUE (TYPE_DOMAIN (type)); - const char *name = NULL; - if (DECL_NAME (orig_var)) - name = IDENTIFIER_POINTER (DECL_NAME (orig_var)); + const char *name = get_name (orig_var); if (TREE_CONSTANT (v)) { x = create_tmp_var_raw (type, name); @@ -4139,7 +4211,8 @@ lower_rec_input_clauses (tree clauses, g } else { - tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA); + tree atmp + = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); tree t = maybe_lookup_decl (v, ctx); if (t) v = t; @@ -4152,7 +4225,8 @@ lower_rec_input_clauses (tree clauses, g t = fold_build2_loc (clause_loc, MULT_EXPR, TREE_TYPE (v), t, TYPE_SIZE_UNIT (TREE_TYPE (type))); - x = build_call_expr_loc (clause_loc, atmp, 1, t); + tree al = size_int (TYPE_ALIGN (TREE_TYPE (type))); + x = build_call_expr_loc (clause_loc, atmp, 2, t, al); } tree ptype = build_pointer_type (TREE_TYPE (type)); @@ -4362,8 +4436,9 @@ lower_rec_input_clauses (tree clauses, g x = TYPE_SIZE_UNIT (TREE_TYPE (new_var)); /* void *tmp = __builtin_alloca */ - atmp = builtin_decl_explicit (BUILT_IN_ALLOCA); - stmt = gimple_build_call (atmp, 1, x); + atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); + stmt = gimple_build_call (atmp, 2, x, + size_int (DECL_ALIGN (var))); tmp = create_tmp_var_raw (ptr_type_node); gimple_add_tmp_var (tmp); gimple_call_set_lhs (stmt, tmp); @@ -4400,12 +4475,8 @@ lower_rec_input_clauses (tree clauses, g x = NULL_TREE; else { - const char *name = NULL; - if (DECL_NAME (var)) - name = IDENTIFIER_POINTER (DECL_NAME (new_var)); - x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)), - name); + get_name (var)); gimple_add_tmp_var (x); TREE_ADDRESSABLE (x) = 1; x = build_fold_addr_expr_loc (clause_loc, x); @@ -4413,8 +4484,11 @@ lower_rec_input_clauses (tree clauses, g } else { - tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA); - x = build_call_expr_loc (clause_loc, atmp, 1, x); + tree atmp + = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); + tree rtype = TREE_TYPE (TREE_TYPE (new_var)); + tree al = size_int (TYPE_ALIGN (rtype)); + x = build_call_expr_loc (clause_loc, atmp, 2, x, al); } if (x) @@ -5489,11 +5563,7 @@ lower_send_clauses (tree clauses, gimple /* Handle taskloop firstprivate/lastprivate, where the lastprivate on GIMPLE_OMP_TASK is represented as OMP_CLAUSE_SHARED_FIRSTPRIVATE. */ - tree f - = (tree) - splay_tree_lookup (ctx->sfield_map - ? ctx->sfield_map : ctx->field_map, - (splay_tree_key) &DECL_UID (val))->value; + tree f = lookup_sfield ((splay_tree_key) &DECL_UID (val), ctx); x = omp_build_component_ref (ctx->sender_decl, f); if (use_pointer_for_field (val, ctx)) var = build_fold_addr_expr (var); @@ -12883,6 +12953,7 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_FIRSTPRIVATE_POINTER: break; case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_TO: @@ -12918,6 +12989,28 @@ lower_omp_target (gimple_stmt_iterator * var = var2; } + if (offloaded + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + { + if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + { + tree type = build_pointer_type (TREE_TYPE (var)); + tree new_var = lookup_decl (var, ctx); + x = create_tmp_var_raw (type, get_name (new_var)); + gimple_add_tmp_var (x); + x = build_simple_mem_ref (x); + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + continue; + } + + if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c)) + { + map_cnt++; + continue; + } + if (!maybe_lookup_field (var, ctx)) continue; @@ -12925,6 +13018,7 @@ lower_omp_target (gimple_stmt_iterator * { x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) && TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) @@ -12936,14 +13030,70 @@ lower_omp_target (gimple_stmt_iterator * break; case OMP_CLAUSE_FIRSTPRIVATE: - case OMP_CLAUSE_IS_DEVICE_PTR: map_cnt++; var = OMP_CLAUSE_DECL (c); if (!is_reference (var) && !is_gimple_reg_type (TREE_TYPE (var))) { - x = build_receiver_ref (var, true, ctx); tree new_var = lookup_decl (var, ctx); + if (is_variable_sized (var)) + { + tree pvar = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); + pvar = TREE_OPERAND (pvar, 0); + gcc_assert (DECL_P (pvar)); + tree new_pvar = lookup_decl (pvar, ctx); + x = build_fold_indirect_ref (new_pvar); + TREE_THIS_NOTRAP (x) = 1; + } + else + x = build_receiver_ref (var, true, ctx); + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + break; + + case OMP_CLAUSE_PRIVATE: + var = OMP_CLAUSE_DECL (c); + if (is_variable_sized (var)) + { + tree new_var = lookup_decl (var, ctx); + tree pvar = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); + pvar = TREE_OPERAND (pvar, 0); + gcc_assert (DECL_P (pvar)); + tree new_pvar = lookup_decl (pvar, ctx); + x = build_fold_indirect_ref (new_pvar); + TREE_THIS_NOTRAP (x) = 1; + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + break; + + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: + var = OMP_CLAUSE_DECL (c); + map_cnt++; + if (is_variable_sized (var)) + { + tree new_var = lookup_decl (var, ctx); + tree pvar = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); + pvar = TREE_OPERAND (pvar, 0); + gcc_assert (DECL_P (pvar)); + tree new_pvar = lookup_decl (pvar, ctx); + x = build_fold_indirect_ref (new_pvar); + TREE_THIS_NOTRAP (x) = 1; + SET_DECL_VALUE_EXPR (new_var, x); + DECL_HAS_VALUE_EXPR_P (new_var) = 1; + } + else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + { + tree new_var = lookup_decl (var, ctx); + tree type = build_pointer_type (TREE_TYPE (var)); + x = create_tmp_var_raw (type, get_name (new_var)); + gimple_add_tmp_var (x); + x = build_simple_mem_ref (x); SET_DECL_VALUE_EXPR (new_var, x); DECL_HAS_VALUE_EXPR_P (new_var) = 1; } @@ -13013,7 +13163,7 @@ lower_omp_target (gimple_stmt_iterator * for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { - tree ovar, nc, s, purpose, var, x; + tree ovar, nc, s, purpose, var, x, type; unsigned int talign; default: @@ -13044,6 +13194,10 @@ lower_omp_target (gimple_stmt_iterator * } else { + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + break; if (DECL_SIZE (ovar) && TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST) { @@ -13053,7 +13207,14 @@ lower_omp_target (gimple_stmt_iterator * gcc_assert (DECL_P (ovar2)); ovar = ovar2; } - if (!maybe_lookup_field (ovar, ctx)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_PRIVATE (c)) + { + if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar), + ctx)) + continue; + } + else if (!maybe_lookup_field (ovar, ctx)) continue; } @@ -13063,7 +13224,12 @@ lower_omp_target (gimple_stmt_iterator * if (nc) { var = lookup_decl_in_outer_ctx (ovar, ctx); - x = build_sender_ref (ovar, ctx); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_PRIVATE (c)) + x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), + ctx); + else + x = build_sender_ref (ovar, ctx); if (maybe_lookup_oacc_reduction (var, ctx)) { gcc_checking_assert (offloaded @@ -13101,7 +13267,7 @@ lower_omp_target (gimple_stmt_iterator * || map_kind == GOMP_MAP_FORCE_DEVICEPTR) && !TYPE_READONLY (TREE_TYPE (var))) { - x = build_sender_ref (ovar, ctx); + x = unshare_expr (x); x = build_simple_mem_ref (x); gimplify_assign (var, x, &olist); } @@ -13121,35 +13287,74 @@ lower_omp_target (gimple_stmt_iterator * if (TREE_CODE (s) != INTEGER_CST) TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0; - unsigned HOST_WIDE_INT tkind; + unsigned HOST_WIDE_INT tkind, tkind_zero; switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_MAP: tkind = OMP_CLAUSE_MAP_KIND (c); + tkind_zero = tkind; + if (OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c)) + switch (tkind) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION; + break; + default: + break; + } + if (tkind_zero != tkind) + { + if (integer_zerop (s)) + tkind = tkind_zero; + else if (integer_nonzerop (s)) + tkind_zero = tkind; + } break; case OMP_CLAUSE_TO: tkind = GOMP_MAP_TO; + tkind_zero = tkind; break; case OMP_CLAUSE_FROM: tkind = GOMP_MAP_FROM; + tkind_zero = tkind; break; default: gcc_unreachable (); } gcc_checking_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift)); + gcc_checking_assert (tkind_zero + < (HOST_WIDE_INT_C (1U) << talign_shift)); talign = ceil_log2 (talign); tkind |= talign << talign_shift; + tkind_zero |= talign << talign_shift; gcc_checking_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); - CONSTRUCTOR_APPEND_ELT (vkind, purpose, - build_int_cstu (tkind_type, tkind)); + gcc_checking_assert (tkind_zero + <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); + if (tkind == tkind_zero) + x = build_int_cstu (tkind_type, tkind); + else + { + TREE_STATIC (TREE_VEC_ELT (t, 2)) = 0; + x = build3 (COND_EXPR, tkind_type, + fold_build2 (EQ_EXPR, boolean_type_node, + unshare_expr (s), size_zero_node), + build_int_cstu (tkind_type, tkind_zero), + build_int_cstu (tkind_type, tkind)); + } + CONSTRUCTOR_APPEND_ELT (vkind, purpose, x); if (nc && nc != c) c = nc; break; case OMP_CLAUSE_FIRSTPRIVATE: - case OMP_CLAUSE_IS_DEVICE_PTR: ovar = OMP_CLAUSE_DECL (c); if (is_reference (ovar)) talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (ovar))); @@ -13157,7 +13362,24 @@ lower_omp_target (gimple_stmt_iterator * talign = DECL_ALIGN_UNIT (ovar); var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); - if (is_reference (var)) + tkind = GOMP_MAP_FIRSTPRIVATE; + type = TREE_TYPE (ovar); + if (is_reference (ovar)) + type = TREE_TYPE (type); + if ((INTEGRAL_TYPE_P (type) + && TYPE_PRECISION (type) <= POINTER_SIZE) + || TREE_CODE (type) == POINTER_TYPE) + { + tkind = GOMP_MAP_FIRSTPRIVATE_INT; + tree t = var; + if (is_reference (var)) + t = build_simple_mem_ref (var); + if (TREE_CODE (type) != POINTER_TYPE) + t = fold_convert (pointer_sized_int_node, t); + t = fold_convert (TREE_TYPE (x), t); + gimplify_assign (x, t, &ilist); + } + else if (is_reference (var)) gimplify_assign (x, var, &ilist); else if (is_gimple_reg (var)) { @@ -13172,7 +13394,9 @@ lower_omp_target (gimple_stmt_iterator * var = build_fold_addr_expr (var); gimplify_assign (x, var, &ilist); } - if (is_reference (var)) + if (tkind == GOMP_MAP_FIRSTPRIVATE_INT) + s = size_int (0); + else if (is_reference (var)) s = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ovar))); else s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); @@ -13182,7 +13406,6 @@ lower_omp_target (gimple_stmt_iterator * if (TREE_CODE (s) != INTEGER_CST) TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0; - tkind = GOMP_MAP_FIRSTPRIVATE; gcc_checking_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift)); talign = ceil_log2 (talign); @@ -13192,6 +13415,40 @@ lower_omp_target (gimple_stmt_iterator * CONSTRUCTOR_APPEND_ELT (vkind, purpose, build_int_cstu (tkind_type, tkind)); break; + + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: + ovar = OMP_CLAUSE_DECL (c); + var = lookup_decl_in_outer_ctx (ovar, ctx); + x = build_sender_ref (ovar, ctx); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) + tkind = GOMP_MAP_USE_DEVICE_PTR; + else + tkind = GOMP_MAP_FIRSTPRIVATE_INT; + type = TREE_TYPE (ovar); + if (TREE_CODE (type) == ARRAY_TYPE) + var = build_fold_addr_expr (var); + else + { + if (is_reference (ovar)) + { + type = TREE_TYPE (type); + if (TREE_CODE (type) != ARRAY_TYPE) + var = build_simple_mem_ref (var); + var = fold_convert (TREE_TYPE (x), var); + } + } + gimplify_assign (x, var, &ilist); + s = size_int (0); + purpose = size_int (map_idx++); + CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); + gcc_checking_assert (tkind + < (HOST_WIDE_INT_C (1U) << talign_shift)); + gcc_checking_assert (tkind + <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); + CONSTRUCTOR_APPEND_ELT (vkind, purpose, + build_int_cstu (tkind_type, tkind)); + break; } gcc_assert (map_idx == map_cnt); @@ -13200,21 +13457,22 @@ lower_omp_target (gimple_stmt_iterator * = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), vsize); DECL_INITIAL (TREE_VEC_ELT (t, 2)) = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind); - if (!TREE_STATIC (TREE_VEC_ELT (t, 1))) - { - gimple_seq initlist = NULL; - force_gimple_operand (build1 (DECL_EXPR, void_type_node, - TREE_VEC_ELT (t, 1)), - &initlist, true, NULL_TREE); - gimple_seq_add_seq (&ilist, initlist); - - tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 1)), - NULL); - TREE_THIS_VOLATILE (clobber) = 1; - gimple_seq_add_stmt (&olist, - gimple_build_assign (TREE_VEC_ELT (t, 1), - clobber)); - } + for (int i = 1; i <= 2; i++) + if (!TREE_STATIC (TREE_VEC_ELT (t, i))) + { + gimple_seq initlist = NULL; + force_gimple_operand (build1 (DECL_EXPR, void_type_node, + TREE_VEC_ELT (t, i)), + &initlist, true, NULL_TREE); + gimple_seq_add_seq (&ilist, initlist); + + tree clobber = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, i)), + NULL); + TREE_THIS_VOLATILE (clobber) = 1; + gimple_seq_add_stmt (&olist, + gimple_build_assign (TREE_VEC_ELT (t, i), + clobber)); + } tree clobber = build_constructor (ctx->record_type, NULL); TREE_THIS_VOLATILE (clobber) = 1; @@ -13237,22 +13495,64 @@ lower_omp_target (gimple_stmt_iterator * gimple_build_assign (ctx->receiver_decl, t)); } - if (offloaded) + if (offloaded || data_region) { + tree prev = NULL_TREE; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) switch (OMP_CLAUSE_CODE (c)) { - tree var; + tree var, x; default: break; case OMP_CLAUSE_FIRSTPRIVATE: - case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); if (is_reference (var) || is_gimple_reg_type (TREE_TYPE (var))) { tree new_var = lookup_decl (var, ctx); - tree x = build_receiver_ref (var, !is_reference (var), ctx); + tree type; + type = TREE_TYPE (var); + if (is_reference (var)) + type = TREE_TYPE (type); + if ((INTEGRAL_TYPE_P (type) + && TYPE_PRECISION (type) <= POINTER_SIZE) + || TREE_CODE (type) == POINTER_TYPE) + { + x = build_receiver_ref (var, false, ctx); + if (TREE_CODE (type) != POINTER_TYPE) + x = fold_convert (pointer_sized_int_node, x); + x = fold_convert (type, x); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, + fb_rvalue); + if (is_reference (var)) + { + tree v = create_tmp_var_raw (type, get_name (var)); + gimple_add_tmp_var (v); + TREE_ADDRESSABLE (v) = 1; + gimple_seq_add_stmt (&new_body, + gimple_build_assign (v, x)); + x = build_fold_addr_expr (v); + } + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + } + else + { + x = build_receiver_ref (var, !is_reference (var), ctx); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, + fb_rvalue); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + } + } + else if (is_variable_sized (var)) + { + tree pvar = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); + pvar = TREE_OPERAND (pvar, 0); + gcc_assert (DECL_P (pvar)); + tree new_var = lookup_decl (pvar, ctx); + x = build_receiver_ref (var, false, ctx); gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); gimple_seq_add_stmt (&new_body, gimple_build_assign (new_var, x)); @@ -13264,23 +13564,22 @@ lower_omp_target (gimple_stmt_iterator * { location_t clause_loc = OMP_CLAUSE_LOCATION (c); tree new_var = lookup_decl (var, ctx); - tree x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var))); + x = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (new_var))); if (TREE_CONSTANT (x)) { - const char *name = NULL; - if (DECL_NAME (var)) - name = IDENTIFIER_POINTER (DECL_NAME (new_var)); - x = create_tmp_var_raw (TREE_TYPE (TREE_TYPE (new_var)), - name); + get_name (var)); gimple_add_tmp_var (x); TREE_ADDRESSABLE (x) = 1; x = build_fold_addr_expr_loc (clause_loc, x); } else { - tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA); - x = build_call_expr_loc (clause_loc, atmp, 1, x); + tree atmp + = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); + tree rtype = TREE_TYPE (TREE_TYPE (new_var)); + tree al = size_int (TYPE_ALIGN (rtype)); + x = build_call_expr_loc (clause_loc, atmp, 2, x, al); } x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x); @@ -13289,9 +13588,169 @@ lower_omp_target (gimple_stmt_iterator * gimple_build_assign (new_var, x)); } break; + case OMP_CLAUSE_USE_DEVICE_PTR: + case OMP_CLAUSE_IS_DEVICE_PTR: + var = OMP_CLAUSE_DECL (c); + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) + x = build_sender_ref (var, ctx); + else + x = build_receiver_ref (var, false, ctx); + if (is_variable_sized (var)) + { + tree pvar = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); + pvar = TREE_OPERAND (pvar, 0); + gcc_assert (DECL_P (pvar)); + tree new_var = lookup_decl (pvar, ctx); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + } + else if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE) + { + tree new_var = lookup_decl (var, ctx); + new_var = DECL_VALUE_EXPR (new_var); + gcc_assert (TREE_CODE (new_var) == MEM_REF); + new_var = TREE_OPERAND (new_var, 0); + gcc_assert (DECL_P (new_var)); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + } + else + { + tree type = TREE_TYPE (var); + tree new_var = lookup_decl (var, ctx); + if (is_reference (var)) + { + type = TREE_TYPE (type); + if (TREE_CODE (type) != ARRAY_TYPE) + { + tree v = create_tmp_var_raw (type, get_name (var)); + gimple_add_tmp_var (v); + TREE_ADDRESSABLE (v) = 1; + x = fold_convert (type, x); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, + fb_rvalue); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (v, x)); + x = build_fold_addr_expr (v); + } + } + x = fold_convert (TREE_TYPE (new_var), x); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + } + break; + } + /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass, + so that firstprivate vars holding OMP_CLAUSE_SIZE if needed + are already handled. */ + for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) + switch (OMP_CLAUSE_CODE (c)) + { + tree var; + default: + break; + case OMP_CLAUSE_MAP: + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER) + { + location_t clause_loc = OMP_CLAUSE_LOCATION (c); + gcc_assert (prev); + var = OMP_CLAUSE_DECL (c); + if (DECL_SIZE (var) + && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST) + { + tree var2 = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (var2) == INDIRECT_REF); + var2 = TREE_OPERAND (var2, 0); + gcc_assert (DECL_P (var2)); + var = var2; + } + tree new_var = lookup_decl (var, ctx), x; + tree type = TREE_TYPE (new_var); + bool is_ref = is_reference (var); + bool ref_to_array = false; + if (is_ref) + { + type = TREE_TYPE (type); + if (TREE_CODE (type) == ARRAY_TYPE) + { + type = build_pointer_type (type); + ref_to_array = true; + } + } + else if (TREE_CODE (type) == ARRAY_TYPE) + { + tree decl2 = DECL_VALUE_EXPR (new_var); + gcc_assert (TREE_CODE (decl2) == MEM_REF); + decl2 = TREE_OPERAND (decl2, 0); + gcc_assert (DECL_P (decl2)); + new_var = decl2; + type = TREE_TYPE (new_var); + } + x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx); + x = fold_convert_loc (clause_loc, type, x); + if (!integer_zerop (OMP_CLAUSE_SIZE (c))) + { + tree bias = OMP_CLAUSE_SIZE (c); + if (DECL_P (bias)) + bias = lookup_decl (bias, ctx); + bias = fold_convert_loc (clause_loc, sizetype, bias); + bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype, + bias); + x = fold_build2_loc (clause_loc, POINTER_PLUS_EXPR, + TREE_TYPE (x), x, bias); + } + if (ref_to_array) + x = fold_convert_loc (clause_loc, TREE_TYPE (new_var), x); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + if (is_ref && !ref_to_array) + { + tree t = create_tmp_var_raw (type, get_name (var)); + gimple_add_tmp_var (t); + TREE_ADDRESSABLE (t) = 1; + gimple_seq_add_stmt (&new_body, + gimple_build_assign (t, x)); + x = build_fold_addr_expr_loc (clause_loc, t); + } + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_var, x)); + prev = NULL_TREE; + } + else if (OMP_CLAUSE_CHAIN (c) + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) + == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + prev = c; + break; + case OMP_CLAUSE_PRIVATE: + var = OMP_CLAUSE_DECL (c); + if (is_variable_sized (var)) + { + location_t clause_loc = OMP_CLAUSE_LOCATION (c); + tree new_var = lookup_decl (var, ctx); + tree pvar = DECL_VALUE_EXPR (var); + gcc_assert (TREE_CODE (pvar) == INDIRECT_REF); + pvar = TREE_OPERAND (pvar, 0); + gcc_assert (DECL_P (pvar)); + tree new_pvar = lookup_decl (pvar, ctx); + tree atmp = builtin_decl_explicit (BUILT_IN_ALLOCA_WITH_ALIGN); + tree al = size_int (DECL_ALIGN (var)); + tree x = TYPE_SIZE_UNIT (TREE_TYPE (new_var)); + x = build_call_expr_loc (clause_loc, atmp, 2, x, al); + x = fold_convert_loc (clause_loc, TREE_TYPE (new_pvar), x); + gimplify_expr (&x, &new_body, NULL, is_gimple_val, fb_rvalue); + gimple_seq_add_stmt (&new_body, + gimple_build_assign (new_pvar, x)); + } + break; } gimple_seq_add_seq (&new_body, tgt_body); - new_body = maybe_catch_exception (new_body); + if (offloaded) + new_body = maybe_catch_exception (new_body); } else if (data_region) new_body = tgt_body; --- gcc/tree-pretty-print.c.jj 2015-07-21 09:06:42.000000000 +0200 +++ gcc/tree-pretty-print.c 2015-07-22 13:53:51.406065024 +0200 @@ -639,6 +639,9 @@ dump_omp_clause (pretty_printer *pp, tre case GOMP_MAP_RELEASE: pp_string (pp, "release"); break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + pp_string (pp, "firstprivate"); + break; default: gcc_unreachable (); } @@ -649,7 +652,9 @@ dump_omp_clause (pretty_printer *pp, tre if (OMP_CLAUSE_SIZE (clause)) { if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER + || OMP_CLAUSE_MAP_KIND (clause) + == GOMP_MAP_FIRSTPRIVATE_POINTER)) pp_string (pp, " [pointer assign, bias: "); else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET) --- gcc/c/c-tree.h.jj 2015-07-01 12:50:49.000000000 +0200 +++ gcc/c/c-tree.h 2015-07-22 12:47:49.185826677 +0200 @@ -649,7 +649,7 @@ extern tree c_begin_omp_task (void); extern tree c_finish_omp_task (location_t, tree, tree); extern void c_finish_omp_cancel (location_t, tree); extern void c_finish_omp_cancellation_point (location_t, tree); -extern tree c_finish_omp_clauses (tree, bool = false); +extern tree c_finish_omp_clauses (tree, bool, bool = false); extern tree c_build_va_arg (location_t, tree, tree); extern tree c_finish_transaction (location_t, tree, int); extern bool c_tree_equal (tree, tree); --- gcc/c/c-parser.c.jj 2015-07-21 09:06:42.000000000 +0200 +++ gcc/c/c-parser.c 2015-07-23 12:51:02.000000000 +0200 @@ -12435,7 +12435,7 @@ c_parser_oacc_all_clauses (c_parser *par c_parser_skip_to_pragma_eol (parser); if (finish_p) - return c_finish_omp_clauses (clauses); + return c_finish_omp_clauses (clauses, false); return clauses; } @@ -12720,8 +12720,8 @@ c_parser_omp_all_clauses (c_parser *pars if (finish_p) { if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_UNIFORM)) != 0) - return c_finish_omp_clauses (clauses, true); - return c_finish_omp_clauses (clauses); + return c_finish_omp_clauses (clauses, true, true); + return c_finish_omp_clauses (clauses, true); } return clauses; @@ -12755,7 +12755,7 @@ c_parser_oacc_cache (location_t loc, c_p tree stmt, clauses; clauses = c_parser_omp_var_list_parens (parser, OMP_CLAUSE__CACHE_, NULL); - clauses = c_finish_omp_clauses (clauses); + clauses = c_finish_omp_clauses (clauses, false); c_parser_skip_to_pragma_eol (parser); @@ -13902,7 +13902,7 @@ omp_split_clauses (location_t loc, enum c_omp_split_clauses (loc, code, mask, clauses, cclauses); for (i = 0; i < C_OMP_CLAUSE_SPLIT_COUNT; i++) if (cclauses[i]) - cclauses[i] = c_finish_omp_clauses (cclauses[i]); + cclauses[i] = c_finish_omp_clauses (cclauses[i], true); } /* OpenMP 4.0: @@ -14668,9 +14668,10 @@ c_parser_omp_target_data (location_t loc case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: map_seen = 3; break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), @@ -14800,9 +14801,10 @@ c_parser_omp_target_enter_data (location case GOMP_MAP_TO: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: map_seen = 3; break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), @@ -14885,9 +14887,10 @@ c_parser_omp_target_exit_data (location_ case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - case GOMP_MAP_POINTER: map_seen = 3; break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; default: map_seen |= 1; error_at (OMP_CLAUSE_LOCATION (*pc), @@ -15016,6 +15019,7 @@ c_parser_omp_target (c_parser *parser, e TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; OMP_TARGET_BODY (stmt) = block; + OMP_TARGET_COMBINED (stmt) = 1; add_stmt (stmt); pc = &OMP_TARGET_CLAUSES (stmt); goto check_clauses; @@ -15078,7 +15082,7 @@ check_clauses: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), @@ -16379,7 +16383,7 @@ c_parser_cilk_for (c_parser *parser, tre tree clauses = build_omp_clause (EXPR_LOCATION (grain), OMP_CLAUSE_SCHEDULE); OMP_CLAUSE_SCHEDULE_KIND (clauses) = OMP_CLAUSE_SCHEDULE_CILKFOR; OMP_CLAUSE_SCHEDULE_CHUNK_EXPR (clauses) = grain; - clauses = c_finish_omp_clauses (clauses); + clauses = c_finish_omp_clauses (clauses, false); tree block = c_begin_compound_stmt (true); tree sb = push_stmt_list (); @@ -16444,7 +16448,7 @@ c_parser_cilk_for (c_parser *parser, tre OMP_CLAUSE_OPERAND (c, 0) = cilk_for_number_of_iterations (omp_for); OMP_CLAUSE_CHAIN (c) = clauses; - OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c); + OMP_PARALLEL_CLAUSES (omp_par) = c_finish_omp_clauses (c, true); add_stmt (omp_par); } --- gcc/c/c-typeck.c.jj 2015-07-17 13:06:58.000000000 +0200 +++ gcc/c/c-typeck.c 2015-07-29 16:14:08.276065810 +0200 @@ -11850,7 +11850,7 @@ handle_omp_array_sections_1 (tree c, tre /* Handle array sections for clause C. */ static bool -handle_omp_array_sections (tree c) +handle_omp_array_sections (tree c, bool is_omp) { bool maybe_zero_len = false; unsigned int first_non_one = 0; @@ -12030,9 +12030,26 @@ handle_omp_array_sections (tree c) if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) return false; gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR); + if (is_omp) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + break; + default: + break; + } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - if (!c_mark_addressable (t)) + OMP_CLAUSE_SET_MAP_KIND (c2, is_omp + ? GOMP_MAP_FIRSTPRIVATE_POINTER + : GOMP_MAP_POINTER); + if (!is_omp && !c_mark_addressable (t)) return false; OMP_CLAUSE_DECL (c2) = t; t = build_fold_addr_expr (first); @@ -12097,7 +12114,7 @@ c_find_omp_placeholder_r (tree *tp, int Remove any elements from the list that are invalid. */ tree -c_finish_omp_clauses (tree clauses, bool declare_simd) +c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) { bitmap_head generic_head, firstprivate_head, lastprivate_head; bitmap_head aligned_head, map_head; @@ -12136,7 +12153,7 @@ c_finish_omp_clauses (tree clauses, bool t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, is_omp)) { remove = true; break; @@ -12496,7 +12513,7 @@ c_finish_omp_clauses (tree clauses, bool } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, is_omp)) remove = true; break; } @@ -12519,7 +12536,7 @@ c_finish_omp_clauses (tree clauses, bool t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, is_omp)) remove = true; else { @@ -12556,6 +12573,8 @@ c_finish_omp_clauses (tree clauses, bool else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_POINTER) + || (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR))) && !lang_hooks.types.omp_mappable_type (TREE_TYPE (t))) { @@ -12624,10 +12643,11 @@ c_finish_omp_clauses (tree clauses, bool case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE) + if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE + && TREE_CODE (TREE_TYPE (t)) != ARRAY_TYPE) { error_at (OMP_CLAUSE_LOCATION (c), - "%qs variable is not a pointer", + "%qs variable is neither a pointer nor an array", omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } --- gcc/cp/parser.c.jj 2015-07-21 09:06:42.000000000 +0200 +++ gcc/cp/parser.c 2015-07-23 12:46:22.172652420 +0200 @@ -32276,27 +32276,28 @@ cp_parser_omp_target_data (cp_parser *pa for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_TOFROM: - case GOMP_MAP_ALWAYS_TOFROM: - case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: - map_seen = 3; - break; - default: - map_seen |= 1; - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target data%> with map-type other " - "than %, %, % or % " - "on % clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALLOC: + map_seen = 3; + break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target data%> with map-type other " + "than %, %, % or % " + "on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -32370,22 +32371,23 @@ cp_parser_omp_target_enter_data (cp_pars for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_TO: - case GOMP_MAP_ALWAYS_TO: - case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: - map_seen = 3; - break; - default: - map_seen |= 1; - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target enter data%> with map-type other " - "than % or % on % clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALLOC: + map_seen = 3; + break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target enter data%> with map-type other " + "than % or % on % clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -32455,24 +32457,25 @@ cp_parser_omp_target_exit_data (cp_parse for (tree *pc = &clauses; *pc;) { if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP) - switch (OMP_CLAUSE_MAP_KIND (*pc)) - { - case GOMP_MAP_FROM: - case GOMP_MAP_ALWAYS_FROM: - case GOMP_MAP_RELEASE: - case GOMP_MAP_DELETE: - case GOMP_MAP_POINTER: - map_seen = 3; - break; - default: - map_seen |= 1; - error_at (OMP_CLAUSE_LOCATION (*pc), - "%<#pragma omp target exit data%> with map-type other " - "than %, % or % on %" - " clause"); - *pc = OMP_CLAUSE_CHAIN (*pc); - continue; - } + switch (OMP_CLAUSE_MAP_KIND (*pc)) + { + case GOMP_MAP_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: + map_seen = 3; + break; + case GOMP_MAP_FIRSTPRIVATE_POINTER: + break; + default: + map_seen |= 1; + error_at (OMP_CLAUSE_LOCATION (*pc), + "%<#pragma omp target exit data%> with map-type other " + "than %, % or % on %" + " clause"); + *pc = OMP_CLAUSE_CHAIN (*pc); + continue; + } pc = &OMP_CLAUSE_CHAIN (*pc); } @@ -32637,6 +32640,7 @@ cp_parser_omp_target (cp_parser *parser, TREE_TYPE (stmt) = void_type_node; OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET]; OMP_TARGET_BODY (stmt) = body; + OMP_TARGET_COMBINED (stmt) = 1; add_stmt (stmt); pc = &OMP_TARGET_CLAUSES (stmt); goto check_clauses; @@ -32697,7 +32701,7 @@ check_clauses: case GOMP_MAP_TOFROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_ALLOC: - case GOMP_MAP_POINTER: + case GOMP_MAP_FIRSTPRIVATE_POINTER: break; default: error_at (OMP_CLAUSE_LOCATION (*pc), --- gcc/cp/semantics.c.jj 2015-07-17 13:59:27.000000000 +0200 +++ gcc/cp/semantics.c 2015-07-29 16:14:49.040467753 +0200 @@ -4650,7 +4650,7 @@ handle_omp_array_sections_1 (tree c, tre /* Handle array sections for clause C. */ static bool -handle_omp_array_sections (tree c) +handle_omp_array_sections (tree c, bool is_omp) { bool maybe_zero_len = false; unsigned int first_non_one = 0; @@ -4826,10 +4826,26 @@ handle_omp_array_sections (tree c) OMP_CLAUSE_SIZE (c) = size; if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) return false; + if (is_omp) + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_TOFROM: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; + break; + default: + break; + } tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER); - if (!cxx_mark_addressable (t)) + OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER + : GOMP_MAP_POINTER); + if (!is_omp && !cxx_mark_addressable (t)) return false; OMP_CLAUSE_DECL (c2) = t; t = build_fold_addr_expr (first); @@ -4847,7 +4863,8 @@ handle_omp_array_sections (tree c) OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = c2; ptr = OMP_CLAUSE_DECL (c2); - if (TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE + if (!is_omp + && TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE && POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr)))) { tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), @@ -5569,7 +5586,7 @@ finish_omp_clauses (tree clauses, bool a t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, allow_fields)) { remove = true; break; @@ -6155,7 +6172,7 @@ finish_omp_clauses (tree clauses, bool a } if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, allow_fields)) remove = true; break; } @@ -6189,7 +6206,7 @@ finish_omp_clauses (tree clauses, bool a t = OMP_CLAUSE_DECL (c); if (TREE_CODE (t) == TREE_LIST) { - if (handle_omp_array_sections (c)) + if (handle_omp_array_sections (c, allow_fields)) remove = true; else { @@ -6242,7 +6259,9 @@ finish_omp_clauses (tree clauses, bool a && !cxx_mark_addressable (t)) remove = true; else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER) + && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER + || (OMP_CLAUSE_MAP_KIND (c) + == GOMP_MAP_FIRSTPRIVATE_POINTER))) && !type_dependent_expression_p (t) && !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) @@ -6428,12 +6447,14 @@ finish_omp_clauses (tree clauses, bool a { tree type = TREE_TYPE (t); if (TREE_CODE (type) != POINTER_TYPE + && TREE_CODE (type) != ARRAY_TYPE && (TREE_CODE (type) != REFERENCE_TYPE - || TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE)) + || (TREE_CODE (TREE_TYPE (type)) != POINTER_TYPE + && TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE))) { error_at (OMP_CLAUSE_LOCATION (c), - "%qs variable is not a pointer or reference " - "to pointer", + "%qs variable is neither a pointer, nor an array" + "nor reference to pointer or array", omp_clause_code_name[OMP_CLAUSE_CODE (c)]); remove = true; } --- include/gomp-constants.h.jj 2015-07-21 09:07:23.689851239 +0200 +++ include/gomp-constants.h 2015-07-29 16:15:20.101012063 +0200 @@ -74,6 +74,17 @@ enum gomp_map_kind GOMP_MAP_FORCE_DEVICEPTR = (GOMP_MAP_FLAG_SPECIAL_1 | 0), /* Do not map, copy bits for firstprivate instead. */ GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0), + /* Similarly, but store the value in the pointer rather than + pointed by the pointer. */ + GOMP_MAP_FIRSTPRIVATE_INT = (GOMP_MAP_FLAG_SPECIAL | 1), + /* Pointer translate host address into device address and copy that + back to host. */ + GOMP_MAP_USE_DEVICE_PTR = (GOMP_MAP_FLAG_SPECIAL | 2), + /* Allocate a zero length array section. Prefer next non-zero length + mapping over previous non-zero length mapping over zero length mapping + at the address. If not already mapped, do nothing (and pointer translate + to NULL). */ + GOMP_MAP_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL | 3), /* Allocate. */ GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC), /* ..., and copy to device. */ @@ -95,7 +106,11 @@ enum gomp_map_kind GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC, /* Decrement usage count and deallocate if zero. */ GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_ALWAYS - | GOMP_MAP_FORCE_DEALLOC) + | GOMP_MAP_FORCE_DEALLOC), + + /* Internal to GCC, not used in libgomp. */ + /* Do not map, but pointer assign a pointer instead. */ + GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1) }; #define GOMP_MAP_COPY_TO_P(X) \ --- libgomp/libgomp.h.jj 2015-07-15 13:00:32.000000000 +0200 +++ libgomp/libgomp.h 2015-07-22 21:09:39.023307107 +0200 @@ -647,11 +647,9 @@ struct target_var_desc { bool copy_from; /* True if data always should be copied from device to host at the end. */ bool always_copy_from; - /* Used for unmapping of array sections, can be nonzero only when - always_copy_from is true. */ + /* Relative offset against key host_start. */ uintptr_t offset; - /* Used for unmapping of array sections, can be less than the size of the - whole object only when always_copy_from is true. */ + /* Actual length. */ uintptr_t length; }; --- libgomp/target.c.jj 2015-07-21 09:07:23.690851224 +0200 +++ libgomp/target.c 2015-07-29 17:12:06.377060519 +0200 @@ -142,7 +142,26 @@ resolve_device (int device_id) } -/* Handle the case where splay_tree_lookup found oldn for newn. +static inline splay_tree_key +gomp_map_lookup (splay_tree mem_map, splay_tree_key key) +{ + if (key->host_start != key->host_end) + return splay_tree_lookup (mem_map, key); + + key->host_end++; + splay_tree_key n = splay_tree_lookup (mem_map, key); + key->host_end--; + if (n) + return n; + key->host_start--; + n = splay_tree_lookup (mem_map, key); + key->host_start++; + if (n) + return n; + return splay_tree_lookup (mem_map, key); +} + +/* Handle the case where gomp_map_lookup found oldn for newn. Helper function of gomp_map_vars. */ static inline void @@ -204,20 +223,8 @@ gomp_map_pointer (struct target_mem_desc } /* Add bias to the pointer value. */ cur_node.host_start += bias; - cur_node.host_end = cur_node.host_start + 1; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - /* Could be possibly zero size array section. */ - cur_node.host_end--; - n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - cur_node.host_start--; - n = splay_tree_lookup (mem_map, &cur_node); - cur_node.host_start++; - } - } + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { gomp_mutex_unlock (&devicep->lock); @@ -271,9 +278,29 @@ gomp_map_vars (struct gomp_device_descr for (i = 0; i < mapnum; i++) { int kind = get_kind (short_mapkind, kinds, i); - if (hostaddrs[i] == NULL) + if (hostaddrs[i] == NULL + || (kind & typemask) == GOMP_MAP_FIRSTPRIVATE_INT) { tgt->list[i].key = NULL; + tgt->list[i].offset = ~(uintptr_t) 0; + continue; + } + else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR) + { + cur_node.host_start = (uintptr_t) hostaddrs[i]; + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); + if (n == NULL) + { + gomp_mutex_unlock (&devicep->lock); + gomp_fatal ("use_device_ptr pointer wasn't mapped"); + } + cur_node.host_start -= n->host_start; + hostaddrs[i] + = (void *) (n->tgt->tgt_start + n->tgt_offset + + cur_node.host_start); + tgt->list[i].key = NULL; + tgt->list[i].offset = ~(uintptr_t) 0; continue; } cur_node.host_start = (uintptr_t) hostaddrs[i]; @@ -293,7 +320,19 @@ gomp_map_vars (struct gomp_device_descr has_firstprivate = true; continue; } - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + splay_tree_key n; + if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) + { + n = gomp_map_lookup (mem_map, &cur_node); + if (!n) + { + tgt->list[i].key = NULL; + tgt->list[i].offset = ~(uintptr_t) 1; + continue; + } + } + else + n = splay_tree_lookup (mem_map, &cur_node); if (n) gomp_map_vars_existing (devicep, n, &cur_node, &tgt->list[i], kind & typemask); @@ -386,6 +425,15 @@ gomp_map_vars (struct gomp_device_descr tgt_size += len; continue; } + switch (kind & typemask) + { + case GOMP_MAP_FIRSTPRIVATE_INT: + case GOMP_MAP_USE_DEVICE_PTR: + case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: + continue; + default: + break; + } splay_tree_key k = &array->key; k->host_start = (uintptr_t) hostaddrs[i]; if (!GOMP_MAP_POINTER_P (kind & typemask)) @@ -518,15 +566,18 @@ gomp_map_vars (struct gomp_device_descr { if (tgt->list[i].key == NULL) { - if (hostaddrs[i] == NULL) - cur_node.tgt_offset = (uintptr_t) NULL; + if (tgt->list[i].offset == ~(uintptr_t) 0) + cur_node.tgt_offset = (uintptr_t) hostaddrs[i]; + else if (tgt->list[i].offset == ~(uintptr_t) 1) + cur_node.tgt_offset = 0; else cur_node.tgt_offset = tgt->tgt_start + tgt->list[i].offset; } else cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start - + tgt->list[i].key->tgt_offset; + + tgt->list[i].key->tgt_offset + + tgt->list[i].offset; /* FIXME: see above FIXME comment. */ devicep->host2dev_func (devicep->target_id, (void *) (tgt->tgt_start @@ -1052,7 +1103,38 @@ GOMP_target_41 (int device, void (*fn) ( if (devicep == NULL || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) - return gomp_target_fallback (fn, hostaddrs); + { + size_t i, tgt_align = 0, tgt_size = 0; + char *tgt = NULL; + for (i = 0; i < mapnum; i++) + if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE) + { + size_t align = (size_t) 1 << (kinds[i] >> 8); + if (tgt_align < align) + tgt_align = align; + tgt_size = (tgt_size + align - 1) & ~(align - 1); + tgt_size += sizes[i]; + } + if (tgt_align) + { + tgt = gomp_alloca (tgt_size + tgt_align - 1); + uintptr_t al = (uintptr_t) tgt & (tgt_align - 1); + if (al) + tgt += tgt_align - al; + tgt_size = 0; + for (i = 0; i < mapnum; i++) + if ((kinds[i] & 0xff) == GOMP_MAP_FIRSTPRIVATE) + { + size_t align = (size_t) 1 << (kinds[i] >> 8); + tgt_size = (tgt_size + align - 1) & ~(align - 1); + memcpy (tgt + tgt_size, hostaddrs[i], sizes[i]); + hostaddrs[i] = tgt + tgt_size; + tgt_size = tgt_size + sizes[i]; + } + } + gomp_target_fallback (fn, hostaddrs); + return; + } void *fn_addr = gomp_get_target_fn_addr (devicep, fn); @@ -1289,20 +1371,8 @@ omp_target_is_present (void *ptr, size_t struct splay_tree_key_s cur_node; cur_node.host_start = (uintptr_t) ptr + offset; - cur_node.host_end = cur_node.host_start + 1; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - /* Could be possibly zero size array section. */ - cur_node.host_end--; - n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - cur_node.host_start--; - n = splay_tree_lookup (mem_map, &cur_node); - cur_node.host_start++; - } - } + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); int ret = n != NULL; gomp_mutex_unlock (&devicep->lock); return ret; @@ -1524,7 +1594,7 @@ omp_target_associate_ptr (void *host_ptr cur_node.host_start = (uintptr_t) host_ptr; cur_node.host_end = cur_node.host_start + size; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n) { if (n->tgt->tgt_start + n->tgt_offset @@ -1584,13 +1654,8 @@ omp_target_disassociate_ptr (void *ptr, int ret = EINVAL; cur_node.host_start = (uintptr_t) ptr; - cur_node.host_end = cur_node.host_start + 1; - splay_tree_key n = splay_tree_lookup (mem_map, &cur_node); - if (n == NULL) - { - cur_node.host_end--; - n = splay_tree_lookup (mem_map, &cur_node); - } + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n && n->host_start == cur_node.host_start && n->refcount == REFCOUNT_INFINITY --- libgomp/testsuite/libgomp.c++/target-2.C.jj 2015-06-30 14:24:03.000000000 +0200 +++ libgomp/testsuite/libgomp.c++/target-2.C 2015-07-23 17:48:08.978674497 +0200 @@ -33,7 +33,8 @@ fn2 (int x, double (&dr) [1024], double int j; fn1 (hr + 2 * x, ir + 2 * x, x); #pragma omp target map(to: br[:x], cr[0:x], dr[x:x], er[x:x]) \ - map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x]) + map(to: fr[0:x], gr[0:x], hr[2 * x:x], ir[2 * x:x]) \ + map(tofrom: s) #pragma omp parallel for reduction(+:s) for (j = 0; j < x; j++) s += br[j] * cr[j] + dr[x + j] + er[x + j] --- libgomp/testsuite/libgomp.c++/target-7.C.jj 2015-07-22 11:36:53.042867520 +0200 +++ libgomp/testsuite/libgomp.c++/target-7.C 2015-07-22 11:32:00.000000000 +0200 @@ -0,0 +1,90 @@ +extern "C" void abort (); + +void +foo (int *x, int *&y, int (&z)[15]) +{ + int a[10], b[15], err, i; + for (i = 0; i < 10; i++) + a[i] = 7 * i; + for (i = 0; i < 15; i++) + b[i] = 8 * i; + #pragma omp target map(to:x[5:10], y[5:10], z[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if (x[5 + i] != 20 + 4 * i + || y[5 + i] != 25 + 5 * i + || z[5 + i] != 30 + 6 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n, int v) +{ + int a[n], b[n], c[n], d[n], e[n], err, i; + int (*x)[n] = &c; + int (*y2)[n] = &d; + int (*&y)[n] = y2; + int (&z)[n] = e; + for (i = 0; i < n; i++) + { + (*x)[i] = 4 * i; + (*y)[i] = 5 * i; + z[i] = 6 * i; + a[i] = 7 * i; + b[i] = 8 * i; + } + #pragma omp target map(to:x[0][5:10], y[0][5:10], z[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 20 + 4 * i + || (*y)[5 + i] != 25 + 5 * i + || z[5 + i] != 30 + 6 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + { + (*x)[i] = 9 * i; + (*y)[i] = 10 * i; + z[i] = 11 * i; + a[i] = 12 * i; + b[i] = 13 * i; + } + #pragma omp target map(to:x[0][v:v+5], y[0][v:v+5], z[v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 45 + 9 * i + || (*y)[5 + i] != 50 + 10 * i + || z[5 + i] != 55 + 11 * i + || a[i] != 12 * i + || b[5 + i] != 65 + 13 * i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + int x[15], y2[15], z[15], *y = y2, i; + for (i = 0; i < 15; i++) + { + x[i] = 4 * i; + y[i] = 5 * i; + z[i] = 6 * i; + } + foo (x, y, z); + bar (15, 5); +} --- libgomp/testsuite/libgomp.c++/target-8.C.jj 2015-07-27 13:39:49.446401028 +0200 +++ libgomp/testsuite/libgomp.c++/target-8.C 2015-07-27 13:39:27.000000000 +0200 @@ -0,0 +1,58 @@ +extern "C" void abort (); +struct S { int a; }; +#ifdef __SIZEOF_INT128__ +typedef __int128 T; +#else +typedef long long int T; +#endif + +void +foo (T a, int b, struct S c) +{ + int err; + #pragma omp target firstprivate (a, b, c) map(from:err) + { + err = 0; + if (a != 131 || b != 276 || c.a != 59) + err = 1; + a = 936; + b = 27; + c.a = 98; + if (a != 936 || b != 27 || c.a != 98) + err = 1; + } + if (err || a != 131 || b != 276 || c.a != 59) + abort (); +} + +void +bar (T &a, int &b, struct S &c) +{ + int err; + #pragma omp target firstprivate (a, b, c) map(from:err) + { + err = 0; + if (a != 131 || b != 276 || c.a != 59) + err = 1; + a = 936; + b = 27; + c.a = 98; + if (a != 936 || b != 27 || c.a != 98) + err = 1; + } + if (err || a != 131 || b != 276 || c.a != 59) + abort (); +} + +int +main () +{ + T a = 131; + int b = 276; + struct S c; + c.a = 59; + foo (a, b, c); + bar (a, b, c); + if (a != 131 || b != 276 || c.a != 59) + abort (); +} --- libgomp/testsuite/libgomp.c++/target-9.C.jj 2015-07-28 16:57:29.940191999 +0200 +++ libgomp/testsuite/libgomp.c++/target-9.C 2015-07-28 20:30:05.951617430 +0200 @@ -0,0 +1,73 @@ +extern "C" void abort (void); + +void +foo (int *&p, int (&s)[5], int n) +{ + int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 }; + int *r = a + 1, *q = p - 1, i, err; + for (i = 0; i < n; i++) + b[i] = 9 + i; + #pragma omp target data map(to:a) + #pragma omp target data use_device_ptr(r) map(from:err) + #pragma omp target is_device_ptr(r) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (r[i - 1] != 7 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:q[:4]) + #pragma omp target data use_device_ptr(p) map(from:err) + #pragma omp target is_device_ptr(p) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (p[i - 1] != i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:b) + #pragma omp target data use_device_ptr(b) map(from:err) + #pragma omp target is_device_ptr(b) private(i) map(from:err) + { + err = 0; + for (i = 0; i < n; i++) + if (b[i] != 9 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:c) + #pragma omp target data use_device_ptr(c) map(from:err) + #pragma omp target is_device_ptr(c) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 3; i++) + if (c[i] != 20 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:s[:5]) + #pragma omp target data use_device_ptr(s) map(from:err) + #pragma omp target is_device_ptr(s) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 5; i++) + if (s[i] != 17 + i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 }; + int *p = a + 1; + foo (p, b, 9); +} --- libgomp/testsuite/libgomp.c/target-1.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/target-1.c 2015-07-23 17:08:32.474133124 +0200 @@ -34,7 +34,7 @@ fn2 (int x, int y, int z) fn1 (b, c, x); #pragma omp target data map(to: b) { - #pragma omp target map(tofrom: c) + #pragma omp target map(tofrom: c, s) #pragma omp teams num_teams(y) thread_limit(z) reduction(+:s) firstprivate(x) #pragma omp distribute dist_schedule(static, 4) collapse(1) for (j=0; j < x; j += y) @@ -52,7 +52,7 @@ fn3 (int x) double b[1024], c[1024], s = 0; int i; fn1 (b, c, x); - #pragma omp target map(to: b, c) + #pragma omp target map(to: b, c) map(tofrom:s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) tgt (), s += b[i] * c[i]; @@ -66,7 +66,8 @@ fn4 (int x, double *p) int i; fn1 (b, c, x); fn1 (d + x, p + x, x); - #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) + #pragma omp target map(to: b, c[0:x], d[x:x]) map(to:p[x:64 + (x & 31)]) \ + map(tofrom: s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i] + d[x + i] + p[x + i]; --- libgomp/testsuite/libgomp.c/target-2.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/target-2.c 2015-07-23 17:09:27.987350372 +0200 @@ -23,7 +23,7 @@ fn2 (int x) int i; fn1 (b, c, x); fn1 (e, d + x, x); - #pragma omp target map(to: b, c[:x], d[x:x], e) + #pragma omp target map(to: b, c[:x], d[x:x], e) map(tofrom: s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i] + d[x + i] + sizeof (b) - sizeof (c); @@ -38,7 +38,7 @@ fn3 (int x) int i; fn1 (b, c, x); fn1 (e, d, x); - #pragma omp target + #pragma omp target map(tofrom: s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) s += b[i] * c[i] + d[i]; @@ -56,7 +56,7 @@ fn4 (int x) #pragma omp target data map(from: b, c[:x], d[x:x], e) { #pragma omp target update to(b, c[:x], d[x:x], e) - #pragma omp target map(c[:x], d[x:x]) + #pragma omp target map(c[:x], d[x:x], s) #pragma omp parallel for reduction(+:s) for (i = 0; i < x; i++) { --- libgomp/testsuite/libgomp.c/target-7.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/target-7.c 2015-07-23 17:12:33.159753962 +0200 @@ -37,63 +37,63 @@ foo (int f) abort (); #pragma omp target data device (d) map (to: h) { - #pragma omp target device (d) + #pragma omp target device (d) map (h) if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 5) abort (); #pragma omp target update device (d) from (h) } #pragma omp target data if (v > 1) map (to: h) { - #pragma omp target if (v > 1) + #pragma omp target if (v > 1) map(h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 6) abort (); #pragma omp target update if (v > 1) from (h) } #pragma omp target data device (d) if (v > 1) map (to: h) { - #pragma omp target device (d) if (v > 1) + #pragma omp target device (d) if (v > 1) map(h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 7) abort (); #pragma omp target update device (d) if (v > 1) from (h) } #pragma omp target data if (v <= 1) map (to: h) { - #pragma omp target if (v <= 1) + #pragma omp target if (v <= 1) map (tofrom: h) if (omp_get_level () != 0 || h++ != 8) abort (); #pragma omp target update if (v <= 1) from (h) } #pragma omp target data device (d) if (v <= 1) map (to: h) { - #pragma omp target device (d) if (v <= 1) + #pragma omp target device (d) if (v <= 1) map (h) if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 9) abort (); #pragma omp target update device (d) if (v <= 1) from (h) } #pragma omp target data if (0) map (to: h) { - #pragma omp target if (0) + #pragma omp target if (0) map (h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 10) abort (); #pragma omp target update if (0) from (h) } #pragma omp target data device (d) if (0) map (to: h) { - #pragma omp target device (d) if (0) + #pragma omp target device (d) if (0) map (h) if (omp_get_level () != 0 || !omp_is_initial_device () || h++ != 11) abort (); #pragma omp target update device (d) if (0) from (h) } #pragma omp target data if (1) map (to: h) { - #pragma omp target if (1) + #pragma omp target if (1) map (tofrom: h) if (omp_get_level () != 0 || h++ != 12) abort (); #pragma omp target update if (1) from (h) } #pragma omp target data device (d) if (1) map (to: h) { - #pragma omp target device (d) if (1) + #pragma omp target device (d) if (1) map (tofrom: h) if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 13) abort (); #pragma omp target update device (d) if (1) from (h) --- libgomp/testsuite/libgomp.c/target-15.c.jj 2015-07-22 11:37:11.655612690 +0200 +++ libgomp/testsuite/libgomp.c/target-15.c 2015-07-23 21:53:37.354632916 +0200 @@ -0,0 +1,74 @@ +extern void abort (void); + +void +foo (int *x) +{ + int a[10], b[15], err, i; + for (i = 0; i < 10; i++) + a[i] = 7 * i; + for (i = 0; i < 15; i++) + b[i] = 8 * i; + #pragma omp target map(to:x[5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if (x[5 + i] != 20 + 4 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n, int v) +{ + int a[n], b[n], c[n], d[n], e[n], err, i; + int (*x)[n] = &c; + for (i = 0; i < n; i++) + { + (*x)[i] = 4 * i; + a[i] = 7 * i; + b[i] = 8 * i; + } + #pragma omp target map(to:x[0][5:10], a[0:10], b[5:10]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 20 + 4 * i + || a[i] != 7 * i + || b[5 + i] != 40 + 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + { + (*x)[i] = 9 * i; + a[i] = 12 * i; + b[i] = 13 * i; + } + #pragma omp target map(to:x[0][v:v+5], a[v-5:v+5], b[v:v+5]) map(from:err) + { + err = 0; + for (i = 0; i < 10; i++) + if ((*x)[5 + i] != 45 + 9 * i + || a[i] != 12 * i + || b[5 + i] != 65 + 13 * i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + int x[15], i; + for (i = 0; i < 15; i++) + x[i] = 4 * i; + foo (x); + bar (15, 5); + return 0; +} --- libgomp/testsuite/libgomp.c/target-16.c.jj 2015-07-23 21:53:28.905753778 +0200 +++ libgomp/testsuite/libgomp.c/target-16.c 2015-07-24 12:20:32.048722516 +0200 @@ -0,0 +1,45 @@ +extern void abort (void); + +void +foo (int n) +{ + int a[n], i, err; + for (i = 0; i < n; i++) + a[i] = 7 * i; + #pragma omp target firstprivate (a) map(from:err) private (i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 7 * i) + err = 1; + } + if (err) + abort (); +} + +void +bar (int n) +{ + int a[n], i, err; + #pragma omp target private (a) map(from:err) + { + #pragma omp parallel for + for (i = 0; i < n; i++) + a[i] = 7 * i; + err = 0; + #pragma omp parallel for reduction(|:err) + for (i = 0; i < n; i++) + if (a[i] != 7 * i) + err |= 1; + } + if (err) + abort (); +} + +int +main () +{ + foo (7); + bar (7); + return 0; +} --- libgomp/testsuite/libgomp.c/target-17.c.jj 2015-07-24 19:50:14.275109272 +0200 +++ libgomp/testsuite/libgomp.c/target-17.c 2015-07-24 19:47:57.000000000 +0200 @@ -0,0 +1,99 @@ +extern void abort (void); + +void +foo (int n) +{ + int a[n], i, err; + for (i = 0; i < n; i++) + a[i] = 5 * i; + #pragma omp target map(to:a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 5 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 6 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target firstprivate (a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 7 * i) + err = 1; + } + if (err) + abort (); + int on = n; + #pragma omp target firstprivate (n) map(tofrom: n) + { + n++; + } + if (on != n) + abort (); + #pragma omp target map(tofrom: n) private (n) + { + n = 25; + } + if (on != n) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 8 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i) + { + err = 0; + for (i = 0; i < n; i++) + if (a[i] != 9 * i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + a[i] += i; + #pragma omp target map(tofrom:a) map(from:err) private(a, i) + { + err = 0; + for (i = 0; i < n; i++) + a[i] = 7; + #pragma omp parallel for reduction(|:err) + for (i = 0; i < n; i++) + if (a[i] != 7) + err |= 1; + } + if (err) + abort (); + for (i = 0; i < n; i++) + if (a[i] != 10 * i) + abort (); +} + +int +main () +{ + foo (9); + return 0; +} --- libgomp/testsuite/libgomp.c/target-18.c.jj 2015-07-28 16:50:12.139587099 +0200 +++ libgomp/testsuite/libgomp.c/target-18.c 2015-07-28 19:59:41.000000000 +0200 @@ -0,0 +1,52 @@ +extern void abort (void); + +void +foo (int n) +{ + int a[4] = { 0, 1, 2, 3 }, b[n]; + int *p = a + 1, i, err; + for (i = 0; i < n; i++) + b[i] = 9 + i; + #pragma omp target data map(to:a) + #pragma omp target data use_device_ptr(p) map(from:err) + #pragma omp target is_device_ptr(p) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (p[i - 1] != i) + err = 1; + } + if (err) + abort (); + for (i = 0; i < 4; i++) + a[i] = 23 + i; + #pragma omp target data map(to:a) + #pragma omp target data use_device_ptr(a) map(from:err) + #pragma omp target is_device_ptr(a) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (a[i] != 23 + i) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:b) + #pragma omp target data use_device_ptr(b) map(from:err) + #pragma omp target is_device_ptr(b) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 4; i++) + if (b[i] != 9 + i) + err = 1; + } + if (err) + abort (); +} + +int +main () +{ + foo (9); + return 0; +} --- libgomp/testsuite/libgomp.c/target-19.c.jj 2015-07-29 16:28:01.783837512 +0200 +++ libgomp/testsuite/libgomp.c/target-19.c 2015-07-29 16:32:42.800714833 +0200 @@ -0,0 +1,127 @@ +extern void abort (void); + +void +foo (int *p, int *q, int *r, int n, int m) +{ + int i, err, *s = r; + #pragma omp target data map(to:p[0:8]) + { + /* For zero length array sections, p points to the start of + already mapped range, q to the end of it, and r does not point + to an mapped range. */ + #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1 || q[i - 8] != i + 1) + err = 1; + if (p + 8 != q || (r != (int *) 0 && r != s)) + err = 1; + } + if (err) + abort (); + /* Implicit mapping of pointers behaves the same way. */ + #pragma omp target private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1 || q[i - 8] != i + 1) + err = 1; + if (p + 8 != q || (r != (int *) 0 && r != s)) + err = 1; + } + if (err) + abort (); + /* And zero-length array sections, though not known at compile + time, behave the same. */ + #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1 || q[i - 8] != i + 1) + err = 1; + if (p + 8 != q || (r != (int *) 0 && r != s)) + err = 1; + } + if (err) + abort (); + /* Non-zero length array sections, though not known at compile, + behave differently. */ + #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r[0] != 10) + err = 1; + } + if (err) + abort (); + #pragma omp target data map(to:q[0:1]) + { + /* For zero length array sections, p points to the start of + already mapped range, q points to the start of another one, + and r to the end of the second one. */ + #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r != q + 1) + err = 1; + } + if (err) + abort (); + /* Implicit mapping of pointers behaves the same way. */ + #pragma omp target private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r != q + 1) + err = 1; + } + if (err) + abort (); + /* And zero-length array sections, though not known at compile + time, behave the same. */ + #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r != q + 1) + err = 1; + } + if (err) + abort (); + /* Non-zero length array sections, though not known at compile, + behave differently. */ + #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err) + { + err = 0; + for (i = 0; i < 8; i++) + if (p[i] != i + 1) + err = 1; + if (q[0] != 9 || r[0] != 10) + err = 1; + } + if (err) + abort (); + } + } +} + +int +main () +{ + int a[32], i; + for (i = 0; i < 32; i++) + a[i] = i; + foo (a + 1, a + 9, a + 10, 0, 1); + return 0; +} --- libgomp/testsuite/libgomp.c/examples-4/e.51.3.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.51.3.c 2015-07-23 15:58:15.867779262 +0200 @@ -47,7 +47,7 @@ void gramSchmidt (int Q[][COLS], const i { int tmp = 0; - #pragma omp target + #pragma omp target map(tofrom:tmp) #pragma omp parallel for reduction(+:tmp) for (i = 0; i < rows; i++) tmp += (Q[i][k] * Q[i][k]); --- libgomp/testsuite/libgomp.c/examples-4/e.53.1.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.53.1.c 2015-07-23 15:59:44.430518114 +0200 @@ -20,7 +20,7 @@ int fib_wrapper (int n) { int x = 0; - #pragma omp target if(n > THRESHOLD) + #pragma omp target if(n > THRESHOLD) map(from:x) x = fib (n); return x; --- libgomp/testsuite/libgomp.c/examples-4/e.53.4.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.53.4.c 2015-07-23 16:00:22.468976440 +0200 @@ -41,7 +41,7 @@ float accum (int k) int i; float tmp = 0.0; - #pragma omp target + #pragma omp target map(tofrom:tmp) #pragma omp parallel for reduction(+:tmp) for (i = 0; i < N; i++) tmp += Pfun (i, k); --- libgomp/testsuite/libgomp.c/examples-4/e.53.5.c.jj 2015-06-17 21:00:36.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.53.5.c 2015-07-23 16:01:17.802188485 +0200 @@ -48,7 +48,7 @@ float accum () int i, k; float tmp = 0.0; - #pragma omp target + #pragma omp target map(tofrom:tmp) #pragma omp parallel for reduction(+:tmp) for (i = 0; i < N; i++) { --- libgomp/testsuite/libgomp.c/examples-4/e.54.2.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.54.2.c 2015-07-23 16:02:02.343554209 +0200 @@ -32,7 +32,7 @@ float dotprod (float B[], float C[], int int i, i0; float sum = 0; - #pragma omp target map(to: B[0:n], C[0:n]) + #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom: sum) #pragma omp teams num_teams(num_teams) thread_limit(block_threads) \ reduction(+:sum) #pragma omp distribute --- libgomp/testsuite/libgomp.c/examples-4/e.54.3.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.54.3.c 2015-07-23 16:02:28.060187999 +0200 @@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int int i; float sum = 0; - #pragma omp target teams map(to: B[0:n], C[0:n]) + #pragma omp target teams map(to: B[0:n], C[0:n]) map(tofrom: sum) #pragma omp distribute parallel for reduction(+:sum) for (i = 0; i < n; i++) sum += B[i] * C[i]; --- libgomp/testsuite/libgomp.c/examples-4/e.54.4.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.54.4.c 2015-07-23 16:03:21.446427770 +0200 @@ -31,7 +31,7 @@ float dotprod (float B[], float C[], int int i; float sum = 0; - #pragma omp target map(to: B[0:n], C[0:n]) + #pragma omp target map(to: B[0:n], C[0:n]) map(tofrom:sum) #pragma omp teams num_teams(8) thread_limit(16) #pragma omp distribute parallel for reduction(+:sum) \ dist_schedule(static, 1024) \ --- libgomp/testsuite/libgomp.c/examples-4/e.57.1.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.57.1.c 2015-07-23 17:37:01.880139916 +0200 @@ -10,11 +10,11 @@ int main () int b = 0; int c, d; - #pragma omp target if(a > 200 && a < 400) + #pragma omp target if(a > 200 && a < 400) map(from: c) c = omp_is_initial_device (); #pragma omp target data map(to: b) if(a > 200 && a < 400) - #pragma omp target + #pragma omp target map(from: b, d) { b = 100; d = omp_is_initial_device (); @@ -26,11 +26,11 @@ int main () a += 200; b = 0; - #pragma omp target if(a > 200 && a < 400) + #pragma omp target if(a > 200 && a < 400) map(from: c) c = omp_is_initial_device (); #pragma omp target data map(to: b) if(a > 200 && a < 400) - #pragma omp target + #pragma omp target map(from: b, d) { b = 100; d = omp_is_initial_device (); @@ -42,11 +42,11 @@ int main () a += 200; b = 0; - #pragma omp target if(a > 200 && a < 400) + #pragma omp target if(a > 200 && a < 400) map(from: c) c = omp_is_initial_device (); #pragma omp target data map(to: b) if(a > 200 && a < 400) - #pragma omp target + #pragma omp target map(from: b, d) { b = 100; d = omp_is_initial_device (); --- libgomp/testsuite/libgomp.c/examples-4/e.57.3.c.jj 2015-04-24 12:30:40.000000000 +0200 +++ libgomp/testsuite/libgomp.c/examples-4/e.57.3.c 2015-07-23 16:08:48.176775074 +0200 @@ -9,7 +9,7 @@ int main () int res; int default_device = omp_get_default_device (); - #pragma omp target + #pragma omp target map(from: res) res = omp_is_initial_device (); if (res) @@ -17,7 +17,7 @@ int main () omp_set_default_device (omp_get_num_devices ()); - #pragma omp target + #pragma omp target map(from: res) res = omp_is_initial_device (); if (!res) Jakub