From patchwork Wed Sep 2 15:58:54 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Jakub Jelinek X-Patchwork-Id: 513554 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 56BF61401E7 for ; Thu, 3 Sep 2015 01:59:12 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=isRSVa+b; dkim-atps=neutral DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:reply-to:references:mime-version :content-type:content-transfer-encoding:in-reply-to; q=dns; s= default; b=LFEHqBn3qvv+NsE0YV0/MjdUg5SZbRLjuzayCEy2vk/gNTrLhAq7a NW+CBxZPFTwLgXhmSV2aWkBRv49qDDIJRmxrZ63BAUyZdio7QC/4ghNpIkWKcGxg ND7oCWaViweBTfX6KX0jkpgYJN+pVXEF5Zv0CpMdjLw1gc9ZeJsfjE= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:cc:subject:message-id:reply-to:references:mime-version :content-type:content-transfer-encoding:in-reply-to; s=default; bh=dzGExA20sBYTHTWl2MnlIAX7ULU=; b=isRSVa+bkaElWCgCSFBZ7m9Qclqo F2Xl+bG29KiO6kx9J9VU1M6aH+AZp6IkUI5iqFypNCXxd0J796qQD4tcWSBYoje4 0h47brCIoZ2hhzaTHpwNDeVN3bIUYjk/3/YqdXiQSBu5SNxvbAL23HrJoHXUAtQZ 7F2MjZ6344c0pYM= Received: (qmail 52979 invoked by alias); 2 Sep 2015 15:59:05 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 52967 invoked by uid 89); 2 Sep 2015 15:59:04 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-0.0 required=5.0 tests=AWL, BAYES_50, KAM_LAZY_DOMAIN_SECURITY, SPF_HELO_PASS, T_RP_MATCHES_RCVD autolearn=no 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, 02 Sep 2015 15:59:01 +0000 Received: from int-mx11.intmail.prod.int.phx2.redhat.com (int-mx11.intmail.prod.int.phx2.redhat.com [10.5.11.24]) by mx1.redhat.com (Postfix) with ESMTPS id 6BDF28CF6F; Wed, 2 Sep 2015 15:58:59 +0000 (UTC) Received: from tucnak.zalov.cz (ovpn-116-44.ams2.redhat.com [10.36.116.44]) by int-mx11.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id t82FwvMm014483 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=NO); Wed, 2 Sep 2015 11:58:58 -0400 Received: from tucnak.zalov.cz (localhost [127.0.0.1]) by tucnak.zalov.cz (8.15.2/8.15.2) with ESMTP id t82Fwtgx008565; Wed, 2 Sep 2015 17:58:56 +0200 Received: (from jakub@localhost) by tucnak.zalov.cz (8.15.2/8.15.2/Submit) id t82Fwsru008564; Wed, 2 Sep 2015 17:58:54 +0200 Date: Wed, 2 Sep 2015 17:58:54 +0200 From: Jakub Jelinek To: Ilya Verbin Cc: gcc-patches@gcc.gnu.org Subject: [gomp4.1] Depend clause support for offloading Message-ID: <20150902155854.GD1847@tucnak.redhat.com> Reply-To: Jakub Jelinek References: <20150731161610.GF1780@tucnak.redhat.com> <20150828181335.GS9425@tucnak.redhat.com> <20150831150753.GC1847@tucnak.redhat.com> <20150902112114.GA19034@msticlxl57.ims.intel.com> MIME-Version: 1.0 Content-Disposition: inline In-Reply-To: <20150902112114.GA19034@msticlxl57.ims.intel.com> User-Agent: Mutt/1.5.23 (2014-03-12) X-IsSubscribed: yes Hi! On Wed, Sep 02, 2015 at 02:21:14PM +0300, Ilya Verbin wrote: > On Mon, Aug 31, 2015 at 17:07:53 +0200, Jakub Jelinek wrote: > > * gimplify.c (gimplify_scan_omp_clauses): Handle > > struct element GOMP_MAP_FIRSTPRIVATE_POINTER. > > Have you seen this? > > gcc/gimplify.c: In function ‘void gimplify_scan_omp_clauses(tree_node**, gimple_statement_base**, omp_region_type, tree_code)’: > gcc/gimplify.c:6578:12: error: ‘sc’ may be used uninitialized in this function [-Werror=maybe-uninitialized] > : *sc != c; > ^ I haven't, but I haven't bootstrapped it for a while, just keep doing make -C gcc -j16 -k check RUNTESTFLAGS=gomp.exp and make check-target-libgomp. That said, this looks like a false positive, but I've added a NULL initialization for it anyway. Here is the start of the async offloading support I've talked about, but nowait is not supported on the library side yet, only depend clause (and for that I haven't added a testcase yet). 2015-09-02 Jakub Jelinek * gimplify.c (gimplify_scan_omp_clauses): Initialize sc to NULL to avoid false positive warnings. * omp-low.c (check_omp_nesting_restrictions): Diagnose depend(source) or depend(sink:...) on #pragma omp target *. (expand_omp_target): Pass flags and depend arguments to GOMP_target_{41,update_41,enter_exit_data} libcalls. (lower_depend_clauses): Change first argument from gimple to tree * pointing to the stmt's clauses. (lower_omp_taskreg): Adjust caller. (lower_omp_target): Lower depend clauses. Always use 16-bit kinds and 8 as align shift. Use GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION for zero length array section in map clause with delete kind. * omp-builtins.def (BUILT_IN_GOMP_TARGET, BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA): Add flags and depend arguments. (BUILT_IN_GOMP_TARGET_UPDATE): Change library function name to GOMP_target_update_41. Add flags and depend arguments, remove unused argument. * builtin-types.def (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR): Remove. (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): New. gcc/c/ * c-typeck.c (handle_omp_array_sections): Set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION even for GOMP_MAP_DELETE kinds. gcc/cp/ * semantics.c (handle_omp_array_sections): Set OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION even for GOMP_MAP_DELETE kinds. gcc/fortran/ * types.def (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR): Remove. (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): New. include/ * gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION. (GOMP_TARGET_FLAG_NOWAIT, GOMP_TARGET_FLAG_EXIT_DATA): Define. libgomp/ * libgomp_g.h (GOMP_target_41, GOMP_target_enter_exit_data): Add flags and depend arguments. (GOMP_target_update_41): New prototype. * libgomp.h (gomp_task_maybe_wait_for_dependencies): New prototype. * libgomp.map (GOMP_4.1): Add GOMP_target_update_41. * task.c (gomp_task_maybe_wait_for_dependencies): Remove prototype. No longer static. * target.c (GOMP_target_41): Add flags and depend arguments. If depend is non-NULL, wait until all dependencies are satisfied. (GOMP_target_enter_exit_data): Likewise. Use flags & GOMP_TARGET_FLAG_EXIT_DATA to determine if it is enter or exit data construct, instead of analysing kinds. (gomp_exit_data): Handle GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION. (GOMP_target_update_41): New function. * testsuite/libgomp.c/target-24.c: New test. Jakub --- gcc/gimplify.c.jj 2015-08-31 16:57:23.000000000 +0200 +++ gcc/gimplify.c 2015-09-02 14:20:41.012253248 +0200 @@ -6557,8 +6557,8 @@ gimplify_scan_omp_clauses (tree *list_p, } else { - tree *osc = struct_map_to_clause->get (decl), *sc; - tree *pt = NULL; + tree *osc = struct_map_to_clause->get (decl); + tree *sc = NULL, *pt = NULL; if (!ptr && TREE_CODE (*osc) == TREE_LIST) osc = &TREE_PURPOSE (*osc); if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) --- gcc/omp-low.c.jj 2015-09-01 17:39:05.000000000 +0200 +++ gcc/omp-low.c 2015-09-02 15:13:13.726567918 +0200 @@ -3440,6 +3440,19 @@ check_omp_nesting_restrictions (gimple s } break; case GIMPLE_OMP_TARGET: + for (c = gimple_omp_target_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND + && (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SOURCE + || OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)) + { + enum omp_clause_depend_kind kind = OMP_CLAUSE_DEPEND_KIND (c); + gcc_assert (kind == OMP_CLAUSE_DEPEND_SOURCE + || kind == OMP_CLAUSE_DEPEND_SINK); + error_at (OMP_CLAUSE_LOCATION (c), + "% is only allowed in %", + kind == OMP_CLAUSE_DEPEND_SOURCE ? "source" : "sink"); + return false; + } for (; ctx != NULL; ctx = ctx->outer) { if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET) @@ -10639,9 +10652,10 @@ expand_omp_target (struct omp_region *re /* Emit a library call to launch the offloading region, or do data transfers. */ - tree t1, t2, t3, t4, device, cond, c, clauses; + tree t1, t2, t3, t4, device, cond, depend, c, clauses; enum built_in_function start_ix; location_t clause_loc; + unsigned int flags_i = 0; switch (gimple_omp_target_kind (entry_stmt)) { @@ -10655,8 +10669,11 @@ expand_omp_target (struct omp_region *re start_ix = BUILT_IN_GOMP_TARGET_UPDATE; break; case GF_OMP_TARGET_KIND_ENTER_DATA: + start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA; + break; case GF_OMP_TARGET_KIND_EXIT_DATA: start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA; + flags_i |= GOMP_TARGET_FLAG_EXIT_DATA; break; case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: @@ -10702,6 +10719,10 @@ expand_omp_target (struct omp_region *re else clause_loc = gimple_location (entry_stmt); + c = find_omp_clause (clauses, OMP_CLAUSE_NOWAIT); + if (c) + flags_i |= GOMP_TARGET_FLAG_NOWAIT; + /* Ensure 'device' is of the correct type. */ device = fold_convert_loc (clause_loc, integer_type_node, device); @@ -10781,10 +10802,6 @@ expand_omp_target (struct omp_region *re args.quick_push (device); if (offloaded) args.quick_push (build_fold_addr_expr (child_fn)); - /* This const void * is part of the current ABI, but we're not actually using - it. */ - if (start_ix == BUILT_IN_GOMP_TARGET_UPDATE) - args.quick_push (build_zero_cst (ptr_type_node)); args.quick_push (t1); args.quick_push (t2); args.quick_push (t3); @@ -10792,10 +10809,18 @@ expand_omp_target (struct omp_region *re switch (start_ix) { case BUILT_IN_GOACC_DATA_START: - case BUILT_IN_GOMP_TARGET: case BUILT_IN_GOMP_TARGET_DATA: + break; + case BUILT_IN_GOMP_TARGET: case BUILT_IN_GOMP_TARGET_UPDATE: case BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA: + args.quick_push (build_int_cst (unsigned_type_node, flags_i)); + c = find_omp_clause (clauses, OMP_CLAUSE_DEPEND); + if (c) + depend = OMP_CLAUSE_DECL (c); + else + depend = build_int_cst (ptr_type_node, 0); + args.quick_push (depend); break; case BUILT_IN_GOACC_PARALLEL: { @@ -10891,8 +10916,7 @@ expand_omp_target (struct omp_region *re gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET); gsi_remove (&gsi, true); } - if (data_region - && region->exit) + if (data_region && region->exit) { gsi = gsi_last_bb (region->exit); g = gsi_stmt (gsi); @@ -12923,14 +12947,13 @@ create_task_copyfn (gomp_task *task_stmt } static void -lower_depend_clauses (gimple stmt, gimple_seq *iseq, gimple_seq *oseq) +lower_depend_clauses (tree *pclauses, gimple_seq *iseq, gimple_seq *oseq) { tree c, clauses; gimple g; size_t n_in = 0, n_out = 0, idx = 2, i; - clauses = find_omp_clause (gimple_omp_task_clauses (stmt), - OMP_CLAUSE_DEPEND); + clauses = find_omp_clause (*pclauses, OMP_CLAUSE_DEPEND); gcc_assert (clauses); for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND) @@ -12977,11 +13000,10 @@ lower_depend_clauses (gimple stmt, gimpl gimple_seq_add_stmt (iseq, g); } } - tree *p = gimple_omp_task_clauses_ptr (stmt); c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEPEND); OMP_CLAUSE_DECL (c) = build_fold_addr_expr (array); - OMP_CLAUSE_CHAIN (c) = *p; - *p = c; + OMP_CLAUSE_CHAIN (c) = *pclauses; + *pclauses = c; tree clobber = build_constructor (type, NULL); TREE_THIS_VOLATILE (clobber) = 1; g = gimple_build_assign (array, clobber); @@ -13026,7 +13048,8 @@ lower_omp_taskreg (gimple_stmt_iterator { push_gimplify_context (); dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK)); - lower_depend_clauses (stmt, &dep_ilist, &dep_olist); + lower_depend_clauses (gimple_omp_task_clauses_ptr (stmt), + &dep_ilist, &dep_olist); } if (ctx->srecord_type) @@ -13124,7 +13147,7 @@ lower_omp_target (gimple_stmt_iterator * tree clauses; tree child_fn, t, c; gomp_target *stmt = as_a (gsi_stmt (*gsi_p)); - gbind *tgt_bind, *bind; + gbind *tgt_bind, *bind, *dep_bind = NULL; gimple_seq tgt_body, olist, ilist, orlist, irlist, new_body; location_t loc = gimple_location (stmt); bool offloaded, data_region; @@ -13153,6 +13176,16 @@ lower_omp_target (gimple_stmt_iterator * clauses = gimple_omp_target_clauses (stmt); + gimple_seq dep_ilist = NULL; + gimple_seq dep_olist = NULL; + if (find_omp_clause (clauses, OMP_CLAUSE_DEPEND)) + { + push_gimplify_context (); + dep_bind = gimple_build_bind (NULL, NULL, make_node (BLOCK)); + lower_depend_clauses (gimple_omp_task_clauses_ptr (stmt), + &dep_ilist, &dep_olist); + } + tgt_bind = NULL; tgt_body = NULL; if (offloaded) @@ -13378,19 +13411,8 @@ lower_omp_target (gimple_stmt_iterator * DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1; TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1; TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1; - tree tkind_type; - int talign_shift; - if (is_gimple_omp_oacc (stmt) - || gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_UPDATE) - { - tkind_type = short_unsigned_type_node; - talign_shift = 8; - } - else - { - tkind_type = unsigned_char_type_node; - talign_shift = 3; - } + tree tkind_type = short_unsigned_type_node; + int talign_shift = 8; TREE_VEC_ELT (t, 2) = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt), ".omp_data_kinds"); @@ -13550,6 +13572,8 @@ lower_omp_target (gimple_stmt_iterator * case GOMP_MAP_RELEASE: tkind_zero = GOMP_MAP_ZERO_LEN_ARRAY_SECTION; break; + case GOMP_MAP_DELETE: + tkind_zero = GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION; default: break; } @@ -14039,7 +14063,7 @@ lower_omp_target (gimple_stmt_iterator * bind = gimple_build_bind (NULL, NULL, tgt_bind ? gimple_bind_block (tgt_bind) : NULL_TREE); - gsi_replace (gsi_p, bind, true); + gsi_replace (gsi_p, dep_bind ? dep_bind : bind, true); gimple_bind_add_seq (bind, irlist); gimple_bind_add_seq (bind, ilist); gimple_bind_add_stmt (bind, stmt); @@ -14047,6 +14071,14 @@ lower_omp_target (gimple_stmt_iterator * gimple_bind_add_seq (bind, orlist); pop_gimplify_context (NULL); + + if (dep_bind) + { + gimple_bind_add_seq (dep_bind, dep_ilist); + gimple_bind_add_stmt (dep_bind, bind); + gimple_bind_add_seq (dep_bind, dep_olist); + pop_gimplify_context (dep_bind); + } } /* Expand code for an OpenMP teams directive. */ --- gcc/omp-builtins.def.jj 2015-06-18 15:24:31.000000000 +0200 +++ gcc/omp-builtins.def 2015-09-02 12:51:00.710561827 +0200 @@ -263,15 +263,17 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_C DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_END, "GOMP_single_copy_end", BT_FN_VOID_PTR, ATTR_NOTHROW_LEAF_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target_41", - BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) + BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR, + ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_DATA, "GOMP_target_data_41", BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_END_DATA, "GOMP_target_end_data", BT_FN_VOID, ATTR_NOTHROW_LIST) -DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update", - BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_UPDATE, "GOMP_target_update_41", + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, + ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA, "GOMP_target_enter_exit_data", - BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) + BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TEAMS, "GOMP_teams", BT_FN_VOID_UINT_UINT, ATTR_NOTHROW_LIST) --- gcc/builtin-types.def.jj 2015-06-18 15:24:31.000000000 +0200 +++ gcc/builtin-types.def 2015-09-02 12:51:51.201829660 +0200 @@ -524,11 +524,6 @@ DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_VPTR_PTR BT_INT) DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_SIZE_VPTR_PTR_PTR_INT_INT, BT_BOOL, BT_SIZE, BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT) -DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, - BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR) -DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR, - BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, - BT_PTR, BT_PTR) DEF_FUNCTION_TYPE_7 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, @@ -537,7 +532,13 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG, BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG, BT_PTR_ULONGLONG) +DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, + BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_UINT, + BT_PTR) +DEF_FUNCTION_TYPE_8 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR, + BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, + BT_PTR, BT_PTR, BT_UINT, BT_PTR) DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_UINT) --- gcc/c/c-typeck.c.jj 2015-08-31 16:57:23.000000000 +0200 +++ gcc/c/c-typeck.c 2015-09-02 13:53:39.487580457 +0200 @@ -12070,6 +12070,7 @@ handle_omp_array_sections (tree c, bool case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; break; default: --- gcc/cp/semantics.c.jj 2015-08-31 16:57:23.000000000 +0200 +++ gcc/cp/semantics.c 2015-09-02 13:54:11.019128248 +0200 @@ -4869,6 +4869,7 @@ handle_omp_array_sections (tree c, bool case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: case GOMP_MAP_RELEASE: + case GOMP_MAP_DELETE: OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1; break; default: --- gcc/fortran/types.def.jj 2015-06-18 15:24:31.000000000 +0200 +++ gcc/fortran/types.def 2015-09-02 12:52:20.089410765 +0200 @@ -189,11 +189,6 @@ DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_VPTR_PTR BT_INT) DEF_FUNCTION_TYPE_6 (BT_FN_BOOL_SIZE_VPTR_PTR_PTR_INT_INT, BT_BOOL, BT_SIZE, BT_VOLATILE_PTR, BT_PTR, BT_PTR, BT_INT, BT_INT) -DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR, - BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, - BT_PTR, BT_PTR) -DEF_FUNCTION_TYPE_6 (BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, - BT_VOID, BT_INT, BT_PTR, BT_SIZE, BT_PTR, BT_PTR, BT_PTR) DEF_FUNCTION_TYPE_7 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_UINT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, @@ -202,10 +197,16 @@ DEF_FUNCTION_TYPE_7 (BT_FN_BOOL_BOOL_ULL BT_BOOL, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG, BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG, BT_PTR_ULONGLONG) +DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_PTR, + BT_VOID, BT_INT, BT_SIZE, BT_PTR, BT_PTR, BT_PTR, BT_UINT, + BT_PTR) DEF_FUNCTION_TYPE_8 (BT_FN_VOID_OMPFN_PTR_UINT_LONG_LONG_LONG_LONG_UINT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_LONG, BT_LONG, BT_LONG, BT_LONG, BT_UINT) +DEF_FUNCTION_TYPE_8 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR, + BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, + BT_PTR, BT_PTR, BT_UINT, BT_PTR) DEF_FUNCTION_TYPE_9 (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, --- include/gomp-constants.h.jj 2015-07-31 16:55:38.000000000 +0200 +++ include/gomp-constants.h 2015-09-02 13:53:09.065016663 +0200 @@ -110,6 +110,10 @@ enum gomp_map_kind (address of the last adjacent entry plus its size). */ GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_ALWAYS | GOMP_MAP_FLAG_SPECIAL | 0), + /* Forced deallocation of zero length array section. */ + GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION + = (GOMP_MAP_FLAG_ALWAYS + | GOMP_MAP_FLAG_SPECIAL | 3), /* OpenMP 4.1 alias for forced deallocation. */ GOMP_MAP_DELETE = GOMP_MAP_FORCE_DEALLOC, /* Decrement usage count and deallocate if zero. */ @@ -171,4 +175,8 @@ enum gomp_map_kind #define GOMP_TASK_FLAG_IF (1 << 10) #define GOMP_TASK_FLAG_NOGROUP (1 << 11) +/* GOMP_target{_41,update_41,enter_exit_data} flags argument. */ +#define GOMP_TARGET_FLAG_NOWAIT (1 << 0) +#define GOMP_TARGET_FLAG_EXIT_DATA (1 << 1) + #endif --- libgomp/libgomp_g.h.jj 2015-06-18 15:24:32.000000000 +0200 +++ libgomp/libgomp_g.h 2015-09-02 12:50:21.794126150 +0200 @@ -217,7 +217,7 @@ extern void GOMP_single_copy_end (void * extern void GOMP_target (int, void (*) (void *), const void *, size_t, void **, size_t *, unsigned char *); extern void GOMP_target_41 (int, void (*) (void *), size_t, void **, size_t *, - unsigned short *); + unsigned short *, unsigned int, void **); extern void GOMP_target_data (int, const void *, size_t, void **, size_t *, unsigned char *); extern void GOMP_target_data_41 (int, size_t, void **, size_t *, @@ -225,8 +225,11 @@ extern void GOMP_target_data_41 (int, si extern void GOMP_target_end_data (void); extern void GOMP_target_update (int, const void *, size_t, void **, size_t *, unsigned char *); +extern void GOMP_target_update_41 (int, size_t, void **, size_t *, + unsigned short *, unsigned int, void **); extern void GOMP_target_enter_exit_data (int, size_t, void **, size_t *, - unsigned short *); + unsigned short *, unsigned int, + void **); extern void GOMP_teams (unsigned int, unsigned int); /* oacc-parallel.c */ --- libgomp/libgomp.h.jj 2015-08-31 16:54:12.000000000 +0200 +++ libgomp/libgomp.h 2015-09-02 15:21:44.722166933 +0200 @@ -650,6 +650,7 @@ extern void gomp_init_task (struct gomp_ struct gomp_task_icv *); extern void gomp_end_task (void); extern void gomp_barrier_handle_tasks (gomp_barrier_state_t); +extern void gomp_task_maybe_wait_for_dependencies (void **); static void inline gomp_finish_task (struct gomp_task *task) --- libgomp/libgomp.map.jj 2015-07-10 18:49:17.000000000 +0200 +++ libgomp/libgomp.map 2015-09-02 12:01:18.132047752 +0200 @@ -268,6 +268,7 @@ GOMP_4.1 { global: GOMP_target_41; GOMP_target_data_41; + GOMP_target_update_41; GOMP_target_enter_exit_data; GOMP_taskloop; GOMP_taskloop_ull; --- libgomp/task.c.jj 2015-08-31 16:54:12.000000000 +0200 +++ libgomp/task.c 2015-09-02 15:22:14.162740580 +0200 @@ -108,8 +108,6 @@ gomp_clear_parent (struct gomp_task *chi while (task != children); } -static void gomp_task_maybe_wait_for_dependencies (void **depend); - /* Called when encountering an explicit task directive. If IF_CLAUSE is false, then we must not delay in executing the task. If UNTIED is true, then the task may be executed by any member of the team. @@ -987,7 +985,7 @@ GOMP_taskwait (void) DEPEND is as in GOMP_task. */ -static void +void gomp_task_maybe_wait_for_dependencies (void **depend) { struct gomp_thread *thr = gomp_thread (); --- libgomp/target.c.jj 2015-08-31 16:57:23.000000000 +0200 +++ libgomp/target.c 2015-09-02 15:30:23.350656259 +0200 @@ -1247,10 +1247,22 @@ GOMP_target (int device, void (*fn) (voi void GOMP_target_41 (int device, void (*fn) (void *), size_t mapnum, - void **hostaddrs, size_t *sizes, unsigned short *kinds) + void **hostaddrs, size_t *sizes, unsigned short *kinds, + unsigned int flags, void **depend) { struct gomp_device_descr *devicep = resolve_device (device); + /* If there are depend clauses, but nowait is not present, + block the parent task until the dependencies are resolved + and then just continue with the rest of the function as if it + is a merged task. */ + if (depend != NULL) + { + struct gomp_thread *thr = gomp_thread (); + if (thr->task && thr->task->depend_hash) + gomp_task_maybe_wait_for_dependencies (depend); + } + if (devicep == NULL || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) { @@ -1386,6 +1398,31 @@ GOMP_target_update (int device, const vo gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, false); } +void +GOMP_target_update_41 (int device, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, + unsigned int flags, void **depend) +{ + struct gomp_device_descr *devicep = resolve_device (device); + + /* If there are depend clauses, but nowait is not present, + block the parent task until the dependencies are resolved + and then just continue with the rest of the function as if it + is a merged task. */ + if (depend != NULL) + { + struct gomp_thread *thr = gomp_thread (); + if (thr->task && thr->task->depend_hash) + gomp_task_maybe_wait_for_dependencies (depend); + } + + if (devicep == NULL + || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) + return; + + gomp_update (devicep, mapnum, hostaddrs, sizes, kinds, true); +} + static void gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned short *kinds) @@ -1404,9 +1441,11 @@ gomp_exit_data (struct gomp_device_descr case GOMP_MAP_DELETE: case GOMP_MAP_RELEASE: case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: + case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION: cur_node.host_start = (uintptr_t) hostaddrs[i]; cur_node.host_end = cur_node.host_start + sizes[i]; - splay_tree_key k = kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION + splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION + || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) ? gomp_map_lookup (&devicep->mem_map, &cur_node) : splay_tree_lookup (&devicep->mem_map, &cur_node); if (!k) @@ -1414,7 +1453,9 @@ gomp_exit_data (struct gomp_device_descr if (k->refcount > 0 && k->refcount != REFCOUNT_INFINITY) k->refcount--; - if (kind == GOMP_MAP_DELETE && k->refcount != REFCOUNT_INFINITY) + if ((kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION) + && k->refcount != REFCOUNT_INFINITY) k->refcount = 0; if ((kind == GOMP_MAP_FROM && k->refcount == 0) @@ -1447,42 +1488,28 @@ gomp_exit_data (struct gomp_device_descr void GOMP_target_enter_exit_data (int device, size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds) + size_t *sizes, unsigned short *kinds, + unsigned int flags, void **depend) { struct gomp_device_descr *devicep = resolve_device (device); + /* If there are depend clauses, but nowait is not present, + block the parent task until the dependencies are resolved + and then just continue with the rest of the function as if it + is a merged task. */ + if (depend != NULL) + { + struct gomp_thread *thr = gomp_thread (); + if (thr->task && thr->task->depend_hash) + gomp_task_maybe_wait_for_dependencies (depend); + } + if (devicep == NULL || !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)) return; - /* Determine if this is an "omp target enter data". */ - const int typemask = 0xff; - bool is_enter_data = false; size_t i; - for (i = 0; i < mapnum; i++) - { - unsigned char kind = kinds[i] & typemask; - - if (kind == GOMP_MAP_ALLOC - || kind == GOMP_MAP_TO - || kind == GOMP_MAP_ALWAYS_TO - || kind == GOMP_MAP_STRUCT) - { - is_enter_data = true; - break; - } - - if (kind == GOMP_MAP_FROM - || kind == GOMP_MAP_ALWAYS_FROM - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_RELEASE - || kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION) - break; - - gomp_fatal ("GOMP_target_enter_exit_data unhandled kind 0x%.2x", kind); - } - - if (is_enter_data) + if ((flags & GOMP_TARGET_FLAG_EXIT_DATA) == 0) for (i = 0; i < mapnum; i++) if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT) { --- libgomp/testsuite/libgomp.c/target-24.c.jj 2015-09-02 16:52:08.540815330 +0200 +++ libgomp/testsuite/libgomp.c/target-24.c 2015-09-02 16:54:13.176019999 +0200 @@ -0,0 +1,43 @@ +#include +#include + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + int a[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }; + int *b = a; + int shared_mem = 0; + #pragma omp target map (alloc: shared_mem) + shared_mem = 1; + if (omp_target_is_present (b, 0, d) != shared_mem) + abort (); + #pragma omp target enter data map (to: a) + if (omp_target_is_present (b, 0, d) == 0) + abort (); + #pragma omp target enter data map (alloc: b[:0]) + if (omp_target_is_present (b, 0, d) == 0) + abort (); + #pragma omp target exit data map (release: b[:0]) + if (omp_target_is_present (b, 0, d) == 0) + abort (); + #pragma omp target exit data map (release: b[:0]) + if (omp_target_is_present (b, 0, d) != shared_mem) + abort (); + #pragma omp target enter data map (to: a) + if (omp_target_is_present (b, 0, d) == 0) + abort (); + #pragma omp target enter data map (always, to: b[:0]) + if (omp_target_is_present (b, 0, d) == 0) + abort (); + #pragma omp target exit data map (delete: b[:0]) + if (omp_target_is_present (b, 0, d) != shared_mem) + abort (); + #pragma omp target exit data map (from: b[:0]) + return 0; +}