From patchwork Wed Oct 29 18:47:57 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: David Malcolm X-Patchwork-Id: 404764 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 C4CEC14003E for ; Thu, 30 Oct 2014 05:53:11 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:cc:subject:date:message-id:in-reply-to:references; q=dns; s= default; b=GzLlUcQEIk1bbw0ketTy5HnmdzBB/NY++5500JO9NSX618tQlj7Sy SF8By+CVF6C0ARV0zc3NRQAhg55dKYER3KxKk/TFTa3AAgCXKgC7RKA+Nw4WomSG p4ojP5bIiYeQ+BAFSMZ9p/0qqjuqcuNSSGKZ7E66EVo/X1ErX85piU= 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:from :to:cc:subject:date:message-id:in-reply-to:references; s= default; bh=mKgBphFtRYFXQoaTLvo/L7YW4oo=; b=URD3bUrlGv1dNqpw+3yi 1YnAXMLTFu2Ve6DYF0ine9md4CpyrdsVoxbu0ugOXI3MJO485+XnT8f4R4wWgjeE cJ5WKO/vKK7FB89bk2wMRn493HzXZWeMT+rGx+qpMXXro1MkEZeTPeybvFbZNnT9 QNhXCdcHZPbi9P7ceCfaf0w= Received: (qmail 28516 invoked by alias); 29 Oct 2014 18:52:24 -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 28429 invoked by uid 89); 29 Oct 2014 18:52:23 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.3 required=5.0 tests=AWL, BAYES_00, RP_MATCHES_RCVD, SPF_HELO_PASS, SPF_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 Oct 2014 18:52:20 +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 (8.14.4/8.14.4) with ESMTP id s9TIqImG002969 (version=TLSv1/SSLv3 cipher=DHE-RSA-AES256-GCM-SHA384 bits=256 verify=FAIL) for ; Wed, 29 Oct 2014 14:52:18 -0400 Received: from surprise.redhat.com (vpn-235-224.phx2.redhat.com [10.3.235.224]) by int-mx13.intmail.prod.int.phx2.redhat.com (8.14.4/8.14.4) with ESMTP id s9TIqDvt028359; Wed, 29 Oct 2014 14:52:18 -0400 From: David Malcolm To: gcc-patches@gcc.gnu.org Cc: David Malcolm Subject: [gimple-classes, committed 07/10] Make remaining gimple_omp_target_ accessors typesafe Date: Wed, 29 Oct 2014 14:47:57 -0400 Message-Id: <1414608480-32461-8-git-send-email-dmalcolm@redhat.com> In-Reply-To: <1414608480-32461-1-git-send-email-dmalcolm@redhat.com> References: <1414608480-32461-1-git-send-email-dmalcolm@redhat.com> X-IsSubscribed: yes gcc/ChangeLog.gimple-classes: * gimple.h (gimple_omp_target_clauses): Strengthen param from const_gimple to const gomp_target *. (gimple_omp_target_kind): Likewise. (gimple_omp_target_clauses_ptr): Strengthen param from gimple to gomp_target *. * gimple-walk.c (walk_gimple_op): Add checked cast. * omp-low.c (scan_sharing_clauses): Add checked casts. (create_omp_child_function): Likewise. (check_omp_nesting_restrictions): Likewise. (expand_omp_target): Likewise. (build_omp_regions_1): Likewise. (make_gimple_omp_edges): Likewise. * tree-inline.c (remap_gimple_stmt): Within case GIMPLE_OMP_TARGET, introduce local "omp_target_stmt" via a checked cast and use it in place of "stmt" for typesafety. * tree-nested.c (convert_nonlocal_reference_stmt): Likewise. (convert_local_reference_stmt): Likewise. (convert_tramp_reference_stmt): Add checked cast. (convert_gimple_call): Within case GIMPLE_OMP_TARGET, introduce local "omp_target_stmt" via a checked cast and use it in place of "stmt" for typesafety. --- gcc/ChangeLog.gimple-classes | 25 +++++ gcc/gimple-walk.c | 4 +- gcc/gimple.h | 16 ++-- gcc/omp-low.c | 28 +++--- gcc/tree-inline.c | 12 ++- gcc/tree-nested.c | 224 ++++++++++++++++++++++++------------------- 6 files changed, 184 insertions(+), 125 deletions(-) diff --git a/gcc/ChangeLog.gimple-classes b/gcc/ChangeLog.gimple-classes index f265a25..cc8b97b 100644 --- a/gcc/ChangeLog.gimple-classes +++ b/gcc/ChangeLog.gimple-classes @@ -1,5 +1,30 @@ 2014-10-29 David Malcolm + * gimple.h (gimple_omp_target_clauses): Strengthen param from + const_gimple to const gomp_target *. + (gimple_omp_target_kind): Likewise. + (gimple_omp_target_clauses_ptr): Strengthen param from gimple to + gomp_target *. + + * gimple-walk.c (walk_gimple_op): Add checked cast. + * omp-low.c (scan_sharing_clauses): Add checked casts. + (create_omp_child_function): Likewise. + (check_omp_nesting_restrictions): Likewise. + (expand_omp_target): Likewise. + (build_omp_regions_1): Likewise. + (make_gimple_omp_edges): Likewise. + * tree-inline.c (remap_gimple_stmt): Within case + GIMPLE_OMP_TARGET, introduce local "omp_target_stmt" via a checked + cast and use it in place of "stmt" for typesafety. + * tree-nested.c (convert_nonlocal_reference_stmt): Likewise. + (convert_local_reference_stmt): Likewise. + (convert_tramp_reference_stmt): Add checked cast. + (convert_gimple_call): Within case GIMPLE_OMP_TARGET, introduce + local "omp_target_stmt" via a checked cast and use it in place of + "stmt" for typesafety. + +2014-10-29 David Malcolm + * gimple.h (gimple_omp_single_clauses): Strengthen param from const_gimple to const gomp_single *. (gimple_omp_single_clauses_ptr): Strengthen param from gimple to diff --git a/gcc/gimple-walk.c b/gcc/gimple-walk.c index 5882fe6..1ba48c1 100644 --- a/gcc/gimple-walk.c +++ b/gcc/gimple-walk.c @@ -429,8 +429,8 @@ walk_gimple_op (gimple stmt, walk_tree_fn callback_op, break; case GIMPLE_OMP_TARGET: - ret = walk_tree (gimple_omp_target_clauses_ptr (stmt), callback_op, wi, - pset); + ret = walk_tree (gimple_omp_target_clauses_ptr (as_a (stmt)), + callback_op, wi, pset); if (ret) return ret; break; diff --git a/gcc/gimple.h b/gcc/gimple.h index 26a9c71..3220f9a 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -4956,24 +4956,21 @@ gimple_omp_single_set_clauses (gomp_single *omp_single_stmt, tree clauses) } -/* Return the clauses associated with OMP_TARGET GS. */ +/* Return the clauses associated with OMP_TARGET OMP_TARGET_STMT. */ static inline tree -gimple_omp_target_clauses (const_gimple gs) +gimple_omp_target_clauses (const gomp_target *omp_target_stmt) { - const gomp_target *omp_target_stmt = - as_a (gs); return omp_target_stmt->clauses; } -/* Return a pointer to the clauses associated with OMP_TARGET GS. */ +/* Return a pointer to the clauses associated with OMP_TARGET + OMP_TARGET_STMT. */ static inline tree * -gimple_omp_target_clauses_ptr (gimple gs) +gimple_omp_target_clauses_ptr (gomp_target *omp_target_stmt) { - gomp_target *omp_target_stmt = - as_a (gs); return &omp_target_stmt->clauses; } @@ -4991,9 +4988,8 @@ gimple_omp_target_set_clauses (gomp_target *omp_target_stmt, /* Return the kind of OMP target statemement. */ static inline int -gimple_omp_target_kind (const_gimple g) +gimple_omp_target_kind (const gomp_target *g) { - GIMPLE_CHECK (g, GIMPLE_OMP_TARGET); return (gimple_omp_subcode (g) & GF_OMP_TARGET_KIND_MASK); } diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 07eed83..21eee0f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1647,7 +1647,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in #pragma omp target data, there is nothing to map for those. */ - if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA + if (gimple_omp_target_kind (as_a (ctx->stmt)) + == GF_OMP_TARGET_KIND_DATA && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } @@ -1673,7 +1674,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_field (decl, true, 7, ctx); else install_var_field (decl, true, 3, ctx); - if (gimple_omp_target_kind (ctx->stmt) + if (gimple_omp_target_kind (as_a (ctx->stmt)) == GF_OMP_TARGET_KIND_REGION) install_var_local (decl, ctx); } @@ -1774,7 +1775,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) break; case OMP_CLAUSE_MAP: - if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) + if (gimple_omp_target_kind (as_a (ctx->stmt)) + == GF_OMP_TARGET_KIND_DATA) break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) @@ -1943,7 +1945,7 @@ create_omp_child_function (omp_context *ctx, bool task_copy) omp_context *octx; for (octx = ctx; octx; octx = octx->outer) if (gimple_code (octx->stmt) == GIMPLE_OMP_TARGET - && gimple_omp_target_kind (octx->stmt) + && gimple_omp_target_kind (as_a (octx->stmt)) == GF_OMP_TARGET_KIND_REGION) { target_p = true; @@ -2665,7 +2667,8 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) case GIMPLE_OMP_TEAMS: if (ctx == NULL || gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET - || gimple_omp_target_kind (ctx->stmt) != GF_OMP_TARGET_KIND_REGION) + || (gimple_omp_target_kind (as_a (ctx->stmt)) + != GF_OMP_TARGET_KIND_REGION)) { error_at (gimple_location (stmt), "teams construct not closely nested inside of target " @@ -2676,10 +2679,11 @@ check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) case GIMPLE_OMP_TARGET: for (; ctx != NULL; ctx = ctx->outer) if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET - && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_REGION) + && (gimple_omp_target_kind (as_a (ctx->stmt)) + == GF_OMP_TARGET_KIND_REGION)) { const char *name; - switch (gimple_omp_target_kind (stmt)) + switch (gimple_omp_target_kind (as_a (stmt))) { case GF_OMP_TARGET_KIND_REGION: name = "target"; break; case GF_OMP_TARGET_KIND_DATA: name = "target data"; break; @@ -8406,8 +8410,8 @@ expand_omp_target (struct omp_region *region) gsi = gsi_last_bb (entry_bb); stmt = gsi_stmt (gsi); gcc_assert (stmt && gimple_code (stmt) == GIMPLE_OMP_TARGET - && gimple_omp_target_kind (stmt) - == GF_OMP_TARGET_KIND_REGION); + && (gimple_omp_target_kind (as_a (stmt)) + == GF_OMP_TARGET_KIND_REGION)); gsi_remove (&gsi, true); e = split_block (entry_bb, stmt); entry_bb = e->dest; @@ -8749,7 +8753,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, ; } else if (code == GIMPLE_OMP_TARGET - && gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_UPDATE) + && (gimple_omp_target_kind (as_a (stmt)) + == GF_OMP_TARGET_KIND_UPDATE)) new_omp_region (bb, code, parent); else { @@ -11068,7 +11073,8 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, case GIMPLE_OMP_TARGET: cur_region = new_omp_region (bb, code, cur_region); fallthru = true; - if (gimple_omp_target_kind (last) == GF_OMP_TARGET_KIND_UPDATE) + if (gimple_omp_target_kind (as_a (last)) + == GF_OMP_TARGET_KIND_UPDATE) cur_region = cur_region->outer; break; diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c index cc5c3bb..eceed15 100644 --- a/gcc/tree-inline.c +++ b/gcc/tree-inline.c @@ -1473,10 +1473,14 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id) break; case GIMPLE_OMP_TARGET: - s1 = remap_gimple_seq (gimple_omp_body (stmt), id); - copy = gimple_build_omp_target - (s1, gimple_omp_target_kind (stmt), - gimple_omp_target_clauses (stmt)); + { + gomp_target *omp_target_stmt = as_a (stmt); + s1 = remap_gimple_seq (gimple_omp_body (omp_target_stmt), id); + copy = gimple_build_omp_target ( + s1, + gimple_omp_target_kind (omp_target_stmt), + gimple_omp_target_clauses (omp_target_stmt)); + } break; case GIMPLE_OMP_TEAMS: diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c index 2ee8cfd..57acc50 100644 --- a/gcc/tree-nested.c +++ b/gcc/tree-nested.c @@ -1393,43 +1393,51 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION) - { - save_suppress = info->suppress_expansion; - convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), - wi); - info->suppress_expansion = save_suppress; - walk_body (convert_nonlocal_reference_stmt, - convert_nonlocal_reference_op, info, - gimple_omp_body_ptr (stmt)); - break; - } - save_suppress = info->suppress_expansion; - if (convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), - wi)) - { - tree c, decl; - decl = get_chain_decl (info); - c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP); - OMP_CLAUSE_DECL (c) = decl; - OMP_CLAUSE_MAP_KIND (c) = OMP_CLAUSE_MAP_TO; - OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); - OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt); - gimple_omp_target_set_clauses (as_a (stmt), c); - } + { + gomp_target *omp_target_stmt = as_a (stmt); + if (gimple_omp_target_kind (omp_target_stmt) + != GF_OMP_TARGET_KIND_REGION) + { + save_suppress = info->suppress_expansion; + convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr ( + omp_target_stmt), + wi); + info->suppress_expansion = save_suppress; + walk_body (convert_nonlocal_reference_stmt, + convert_nonlocal_reference_op, info, + gimple_omp_body_ptr (omp_target_stmt)); + break; + } + save_suppress = info->suppress_expansion; + if (convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr ( + omp_target_stmt), + wi)) + { + tree c, decl; + decl = get_chain_decl (info); + c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (c) = decl; + OMP_CLAUSE_MAP_KIND (c) = OMP_CLAUSE_MAP_TO; + OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); + OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (omp_target_stmt); + gimple_omp_target_set_clauses (omp_target_stmt, c); + } - save_local_var_chain = info->new_local_var_chain; - info->new_local_var_chain = NULL; + save_local_var_chain = info->new_local_var_chain; + info->new_local_var_chain = NULL; - walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, - info, gimple_omp_body_ptr (stmt)); + walk_body (convert_nonlocal_reference_stmt, + convert_nonlocal_reference_op, + info, gimple_omp_body_ptr (omp_target_stmt)); - if (info->new_local_var_chain) - declare_vars (info->new_local_var_chain, - gimple_seq_first_stmt (gimple_omp_body (stmt)), - false); - info->new_local_var_chain = save_local_var_chain; - info->suppress_expansion = save_suppress; + if (info->new_local_var_chain) + declare_vars (info->new_local_var_chain, + gimple_seq_first_stmt (gimple_omp_body ( + omp_target_stmt)), + false); + info->new_local_var_chain = save_local_var_chain; + info->suppress_expansion = save_suppress; + } break; case GIMPLE_OMP_TEAMS: @@ -1982,39 +1990,50 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION) - { - save_suppress = info->suppress_expansion; - convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); - info->suppress_expansion = save_suppress; - walk_body (convert_local_reference_stmt, convert_local_reference_op, - info, gimple_omp_body_ptr (stmt)); - break; - } - save_suppress = info->suppress_expansion; - if (convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi)) - { - tree c; - (void) get_frame_type (info); - c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP); - OMP_CLAUSE_DECL (c) = info->frame_decl; - OMP_CLAUSE_MAP_KIND (c) = OMP_CLAUSE_MAP_TOFROM; - OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (info->frame_decl); - OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt); - gimple_omp_target_set_clauses (as_a (stmt), c); + { + gomp_target *omp_target_stmt = as_a (stmt); + if (gimple_omp_target_kind (omp_target_stmt) + != GF_OMP_TARGET_KIND_REGION) + { + save_suppress = info->suppress_expansion; + convert_local_omp_clauses (gimple_omp_target_clauses_ptr ( + omp_target_stmt), + wi); + info->suppress_expansion = save_suppress; + walk_body (convert_local_reference_stmt, convert_local_reference_op, + info, gimple_omp_body_ptr (omp_target_stmt)); + break; } + save_suppress = info->suppress_expansion; + if (convert_local_omp_clauses (gimple_omp_target_clauses_ptr ( + omp_target_stmt), + wi)) + { + tree c; + (void) get_frame_type (info); + c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (c) = info->frame_decl; + OMP_CLAUSE_MAP_KIND (c) = OMP_CLAUSE_MAP_TOFROM; + OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (info->frame_decl); + OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (omp_target_stmt); + gimple_omp_target_set_clauses (omp_target_stmt, c); + } - save_local_var_chain = info->new_local_var_chain; - info->new_local_var_chain = NULL; + save_local_var_chain = info->new_local_var_chain; + info->new_local_var_chain = NULL; - walk_body (convert_local_reference_stmt, convert_local_reference_op, info, - gimple_omp_body_ptr (stmt)); + walk_body (convert_local_reference_stmt, convert_local_reference_op, + info, + gimple_omp_body_ptr (omp_target_stmt)); - if (info->new_local_var_chain) - declare_vars (info->new_local_var_chain, - gimple_seq_first_stmt (gimple_omp_body (stmt)), false); - info->new_local_var_chain = save_local_var_chain; - info->suppress_expansion = save_suppress; + if (info->new_local_var_chain) + declare_vars (info->new_local_var_chain, + gimple_seq_first_stmt (gimple_omp_body ( + omp_target_stmt)), + false); + info->new_local_var_chain = save_local_var_chain; + info->suppress_expansion = save_suppress; + } break; case GIMPLE_OMP_TEAMS: @@ -2320,7 +2339,8 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, } case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION) + if (gimple_omp_target_kind (as_a (stmt)) + != GF_OMP_TARGET_KIND_REGION) { *handled_ops_p = false; return NULL_TREE; @@ -2419,40 +2439,48 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (stmt) != GF_OMP_TARGET_KIND_REGION) - { - walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt)); - break; - } - save_static_chain_added = info->static_chain_added; - info->static_chain_added = 0; - walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt)); - for (i = 0; i < 2; i++) - { - tree c, decl; - if ((info->static_chain_added & (1 << i)) == 0) - continue; - decl = i ? get_chain_decl (info) : info->frame_decl; - /* Don't add CHAIN.* or FRAME.* twice. */ - for (c = gimple_omp_target_clauses (stmt); - c; - c = OMP_CLAUSE_CHAIN (c)) - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && OMP_CLAUSE_DECL (c) == decl) - break; - if (c == NULL) - { - c = build_omp_clause (gimple_location (stmt), OMP_CLAUSE_MAP); - OMP_CLAUSE_DECL (c) = decl; - OMP_CLAUSE_MAP_KIND (c) - = i ? OMP_CLAUSE_MAP_TO : OMP_CLAUSE_MAP_TOFROM; - OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); - OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (stmt); - gimple_omp_target_set_clauses (as_a (stmt), - c); - } - } - info->static_chain_added |= save_static_chain_added; + { + gomp_target *omp_target_stmt = as_a (stmt); + if (gimple_omp_target_kind (omp_target_stmt) + != GF_OMP_TARGET_KIND_REGION) + { + walk_body (convert_gimple_call, NULL, info, + gimple_omp_body_ptr (omp_target_stmt)); + break; + } + save_static_chain_added = info->static_chain_added; + info->static_chain_added = 0; + walk_body (convert_gimple_call, NULL, info, + gimple_omp_body_ptr (omp_target_stmt)); + for (i = 0; i < 2; i++) + { + tree c, decl; + if ((info->static_chain_added & (1 << i)) == 0) + continue; + decl = i ? get_chain_decl (info) : info->frame_decl; + /* Don't add CHAIN.* or FRAME.* twice. */ + for (c = gimple_omp_target_clauses (omp_target_stmt); + c; + c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_DECL (c) == decl) + break; + if (c == NULL) + { + c = build_omp_clause (gimple_location (omp_target_stmt), + OMP_CLAUSE_MAP); + OMP_CLAUSE_DECL (c) = decl; + OMP_CLAUSE_MAP_KIND (c) + = i ? OMP_CLAUSE_MAP_TO : OMP_CLAUSE_MAP_TOFROM; + OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl); + OMP_CLAUSE_CHAIN (c) = + gimple_omp_target_clauses (omp_target_stmt); + gimple_omp_target_set_clauses (omp_target_stmt, + c); + } + } + info->static_chain_added |= save_static_chain_added; + } break; case GIMPLE_OMP_FOR: