From patchwork Tue Jan 14 15:10:05 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 310763 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 364A42C008F for ; Wed, 15 Jan 2014 02:11:45 +1100 (EST) 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 :mime-version:content-type; q=dns; s=default; b=IJ86txTP4WHeT5Ob 9nEfF09XygiyASrsX2MM1vG91v/u+EVYpaNCaHyMXG9ql13RDz+5Gg+xOuyav0lX xuhIfaavwHs4N55M1H3vS2h/16gHu1hZWAnGieJ79vdk2r6VY3dKK2jGdOjqtNGS irvZubxPkVP6KQObheMlyLNfnPM= 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 :mime-version:content-type; s=default; bh=lGAQlhJZ2jeND9WEPZzLTx 1vD/Y=; b=KfW4y7xyh+mkf4QGOnTPXl32sYWlHT3F0XOUGynWLbihxDBf/P1yhc ScN1EaNrMdkG4gGjQb28kzRZ7foS5jwzl9j4+uoJITg8I98tzNRPVBfOhsa53g4/ QwydlHx2O+ft9jBge4eAVjrxMH9icoxc4z8HtUiuOTN/fJaAiVA+A= Received: (qmail 12769 invoked by alias); 14 Jan 2014 15:10:41 -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 12656 invoked by uid 89); 14 Jan 2014 15:10:40 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.3 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-HELO: eggs.gnu.org Received: from eggs.gnu.org (HELO eggs.gnu.org) (208.118.235.92) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES256-SHA encrypted) ESMTPS; Tue, 14 Jan 2014 15:10:36 +0000 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1W35dJ-0006Qt-Ba for gcc-patches@gcc.gnu.org; Tue, 14 Jan 2014 10:10:34 -0500 Received: from relay1.mentorg.com ([192.94.38.131]:62394) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1W35dJ-0006N9-0A for gcc-patches@gcc.gnu.org; Tue, 14 Jan 2014 10:10:25 -0500 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1W35dC-0002Fv-FK from Thomas_Schwinge@mentor.com ; Tue, 14 Jan 2014 07:10:18 -0800 Received: from SVR-ORW-FEM-02.mgc.mentorg.com ([147.34.96.206]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Tue, 14 Jan 2014 07:10:18 -0800 Received: from build5-lucid-cs (147.34.91.1) by svr-orw-fem-02.mgc.mentorg.com (147.34.96.168) with Microsoft SMTP Server id 14.2.247.3; Tue, 14 Jan 2014 07:10:16 -0800 Received: by build5-lucid-cs (Postfix, from userid 49978) id BD9EA321ABB; Tue, 14 Jan 2014 07:10:16 -0800 (PST) From: To: , CC: Thomas Schwinge Subject: [gomp4 3/6] Initial support for OpenACC memory mapping semantics. Date: Tue, 14 Jan 2014 16:10:05 +0100 Message-ID: <1389712208-416-3-git-send-email-thomas@codesourcery.com> In-Reply-To: <1389712208-416-2-git-send-email-thomas@codesourcery.com> References: <87ppnuvbv6.fsf@schwinge.name> <1389712208-416-1-git-send-email-thomas@codesourcery.com> <1389712208-416-2-git-send-email-thomas@codesourcery.com> MIME-Version: 1.0 X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 From: Thomas Schwinge gcc/ * tree-core.h (omp_clause_map_kind): Add OMP_CLAUSE_MAP_FORCE, OMP_CLAUSE_MAP_FORCE_ALLOC, OMP_CLAUSE_MAP_FORCE_TO, OMP_CLAUSE_MAP_FORCE_FROM, OMP_CLAUSE_MAP_FORCE_TOFROM, OMP_CLAUSE_MAP_FORCE_PRESENT, OMP_CLAUSE_MAP_FORCE_DEALLOC, and OMP_CLAUSE_MAP_FORCE_DEVICEPTR. * tree-pretty-print.c (dump_omp_clause): Handle these. * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_FORCE. (omp_region_type): Add ORT_TARGET_MAP_FORCE. (omp_add_variable, omp_notice_threadprivate_variable) (omp_notice_variable, gimplify_scan_omp_clauses) (gimplify_adjust_omp_clauses_1): Extend accordingly. (gimplify_oacc_parallel): Add ORT_TARGET_MAP_FORCE to ORT_TARGET usage. * omp-low.c (install_var_field, scan_sharing_clauses) (lower_oacc_parallel, lower_omp_target): Extend accordingly. --- gcc/gimplify.c | 92 ++++++++++++++++++++++++++++++++++++++++++------- gcc/omp-low.c | 33 +++++++++++------- gcc/tree-core.h | 19 +++++++++- gcc/tree-pretty-print.c | 21 +++++++++++ 4 files changed, 140 insertions(+), 25 deletions(-) diff --git gcc/gimplify.c gcc/gimplify.c index 90507c2..633784f 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -69,7 +69,13 @@ enum gimplify_omp_var_data GOVD_PRIVATE_OUTER_REF = 1024, GOVD_LINEAR = 2048, GOVD_ALIGNED = 4096, + + /* Flags for GOVD_MAP. */ + /* Don't copy back. */ GOVD_MAP_TO_ONLY = 8192, + /* Force a specific behavior (or else, a run-time error). */ + GOVD_MAP_FORCE = 16384, + GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR | GOVD_LOCAL) @@ -86,7 +92,11 @@ enum omp_region_type ORT_UNTIED_TASK = 5, ORT_TEAMS = 8, ORT_TARGET_DATA = 16, - ORT_TARGET = 32 + ORT_TARGET = 32, + + /* Flags for ORT_TARGET. */ + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */ + ORT_TARGET_MAP_FORCE = 64 }; /* Gimplify hashtable helper. */ @@ -5430,9 +5440,20 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags) copy into or out of the context. */ if (!(flags & GOVD_LOCAL)) { - nflags = flags & GOVD_MAP - ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT - : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE; + if (flags & GOVD_MAP) + { + nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT; +#if 0 + /* Not sure if this is actually needed; haven't found a case + where this would change anything; TODO. */ + if (flags & GOVD_MAP_FORCE) + nflags |= OMP_CLAUSE_MAP_FORCE; +#endif + } + else if (flags & GOVD_PRIVATE) + nflags = GOVD_PRIVATE; + else + nflags = GOVD_FIRSTPRIVATE; nflags |= flags & GOVD_SEEN; t = DECL_VALUE_EXPR (decl); gcc_assert (TREE_CODE (t) == INDIRECT_REF); @@ -5501,6 +5522,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl, for (octx = ctx; octx; octx = octx->outer_context) if (octx->region_type & ORT_TARGET) { + gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE)); + n = splay_tree_lookup (octx->variables, (splay_tree_key)decl); if (n == NULL) { @@ -5562,19 +5585,45 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code) n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); if (ctx->region_type & ORT_TARGET) { + unsigned map_force; + if (ctx->region_type & ORT_TARGET_MAP_FORCE) + map_force = GOVD_MAP_FORCE; + else + map_force = 0; if (n == NULL) { if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl))) { error ("%qD referenced in target region does not have " "a mappable type", decl); - omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags); + omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT | flags); } else - omp_add_variable (ctx, decl, GOVD_MAP | flags); + omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags); } else - n->value |= flags; + { +#if 0 + /* The following fails for: + + int l = 10; + float c[l]; + #pragma acc parallel copy(c[2:4]) + { + #pragma acc parallel + { + int t = sizeof c; + } + } + + ..., which we currently don't have to care about (nesting + disabled), but eventually will have to; TODO. */ + if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT)) + gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force); +#endif + + n->value |= flags; + } ret = lang_hooks.decls.omp_disregard_value_expr (decl, true); goto do_outer; } @@ -5858,6 +5907,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, goto do_add; case OMP_CLAUSE_MAP: + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case OMP_CLAUSE_MAP_FORCE_PRESENT: + case OMP_CLAUSE_MAP_FORCE_DEALLOC: + case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: + input_location = OMP_CLAUSE_LOCATION (c); + /* TODO. */ + sorry ("data clause not yet implemented"); + remove = true; + break; + default: + break; + } if (OMP_CLAUSE_SIZE (c) && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) @@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1; else if (code == OMP_CLAUSE_MAP) { - OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY - ? OMP_CLAUSE_MAP_TO - : OMP_CLAUSE_MAP_TOFROM; + unsigned map_kind; + map_kind = (flags & GOVD_MAP_TO_ONLY + ? OMP_CLAUSE_MAP_TO + : OMP_CLAUSE_MAP_TOFROM); + if (flags & GOVD_MAP_FORCE) + map_kind |= OMP_CLAUSE_MAP_FORCE; + OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind; + if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST) { @@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p) tree expr = *expr_p; gimple g; gimple_seq body = NULL; + enum omp_region_type ort = + (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE); - gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, - ORT_TARGET); + gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort); push_gimplify_context (); diff --git gcc/omp-low.c gcc/omp-low.c index 899e970..8c7df1b 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1064,6 +1064,8 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) || !splay_tree_lookup (ctx->field_map, (splay_tree_key) var)); gcc_assert ((mask & 2) == 0 || !ctx->sfield_map || !splay_tree_lookup (ctx->sfield_map, (splay_tree_key) var)); + gcc_assert ((mask & 3) == 3 + || gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); type = TREE_TYPE (var); if (mask & 4) @@ -1611,6 +1613,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); case OMP_CLAUSE_MAP: if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); @@ -1630,11 +1633,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) { - gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL); /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in #pragma omp target data, there is nothing to map for those. */ - if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA + if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL + && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } @@ -8709,8 +8712,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_MAP: - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: var = OMP_CLAUSE_DECL (c); if (!DECL_P (var)) { @@ -8797,8 +8798,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_MAP: - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: nc = c; ovar = OMP_CLAUSE_DECL (c); if (!DECL_P (ovar)) @@ -8893,12 +8892,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_MAP: tkind = OMP_CLAUSE_MAP_KIND (c); break; - case OMP_CLAUSE_TO: - tkind = OMP_CLAUSE_MAP_TO; - break; - case OMP_CLAUSE_FROM: - tkind = OMP_CLAUSE_MAP_FROM; - break; default: gcc_unreachable (); } @@ -10179,6 +10172,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: break; case OMP_CLAUSE_MAP: +#ifdef ENABLE_CHECKING + /* First check what we're prepared to handle in the following. */ + switch (OMP_CLAUSE_MAP_KIND (c)) + { + case OMP_CLAUSE_MAP_ALLOC: + case OMP_CLAUSE_MAP_TO: + case OMP_CLAUSE_MAP_FROM: + case OMP_CLAUSE_MAP_TOFROM: + case OMP_CLAUSE_MAP_POINTER: + break; + default: + gcc_unreachable (); + } +#endif + /* FALLTHRU */ + case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: var = OMP_CLAUSE_DECL (c); diff --git gcc/tree-core.h gcc/tree-core.h index 3602b5f..0aedea3 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -1125,7 +1125,24 @@ enum omp_clause_map_kind /* The following kind is an internal only map kind, used for pointer based array sections. OMP_CLAUSE_SIZE for these is not the pointer size, which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */ - OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL + OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL, + /* The following are only valid for OpenACC. */ + /* Flag to force a specific behavior (or else, a run-time error). */ + OMP_CLAUSE_MAP_FORCE = 1 << 3, + /* Allocate. */ + OMP_CLAUSE_MAP_FORCE_ALLOC = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_ALLOC, + /* ..., and copy to device. */ + OMP_CLAUSE_MAP_FORCE_TO = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TO, + /* ..., and copy from device. */ + OMP_CLAUSE_MAP_FORCE_FROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_FROM, + /* ..., and copy to and from device. */ + OMP_CLAUSE_MAP_FORCE_TOFROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TOFROM, + /* Must already be present. */ + OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL, + /* Deallocate a mapping, without copying from device. */ + OMP_CLAUSE_MAP_FORCE_DEALLOC, + /* Is a device pointer. */ + OMP_CLAUSE_MAP_FORCE_DEVICEPTR }; enum omp_clause_proc_bind_kind diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index 320c35b..f75f181 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -506,6 +506,27 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags) case OMP_CLAUSE_MAP_TOFROM: pp_string (buffer, "tofrom"); break; + case OMP_CLAUSE_MAP_FORCE_ALLOC: + pp_string (buffer, "force_alloc"); + break; + case OMP_CLAUSE_MAP_FORCE_TO: + pp_string (buffer, "force_to"); + break; + case OMP_CLAUSE_MAP_FORCE_FROM: + pp_string (buffer, "force_from"); + break; + case OMP_CLAUSE_MAP_FORCE_TOFROM: + pp_string (buffer, "force_tofrom"); + break; + case OMP_CLAUSE_MAP_FORCE_PRESENT: + pp_string (buffer, "force_present"); + break; + case OMP_CLAUSE_MAP_FORCE_DEALLOC: + pp_string (buffer, "force_dealloc"); + break; + case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: + pp_string (buffer, "force_deviceptr"); + break; default: gcc_unreachable (); }