From patchwork Fri Feb 21 20:32:14 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 323015 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 3CFAC2C0333 for ; Sat, 22 Feb 2014 07:32:57 +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=rMsl3N6oUKyqQ20I HSy6k+DOujkDIVB4ACgzDeruqUbnJgm0XgCmS0FVa7suoxW6Oug26i912oEcAOB3 pxJYj1KuFUoPZici7M3Qpp+plNUwPvr7R4PorG1I9a8D9z/cdkzFmzhs/yWvk4iC dT+uplwvT7rjC0HEibNgkcMSzVA= 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=4sJJ+/31kVfMDs/tNHf0ra Q3bNg=; b=Sw6wtuEvaGg7hHKsZ6O9NmEhXK2SJ5XzyaNpydCJ2y0JxwvmsWE6UQ 74pYbZA4cZobiux9wy9cZfNTYYfNhY2gb3FNfoGIEIhyaOcIXC9fo2CN9Ne6kx7Q 8UtfsR5ET8r63OkmIa9YEuZP+r5t0OTR92lbthuNkT7pgQlgk9S2I= Received: (qmail 19452 invoked by alias); 21 Feb 2014 20:32:38 -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 19384 invoked by uid 89); 21 Feb 2014 20:32:38 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.4 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 21 Feb 2014 20:32:33 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1WGwlq-0000oE-4K from Thomas_Schwinge@mentor.com ; Fri, 21 Feb 2014 12:32:30 -0800 Received: from SVR-ORW-FEM-05.mgc.mentorg.com ([147.34.97.43]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Fri, 21 Feb 2014 12:32:30 -0800 Received: from build5-lucid-cs (147.34.91.1) by svr-orw-fem-05.mgc.mentorg.com (147.34.97.43) with Microsoft SMTP Server id 14.2.247.3; Fri, 21 Feb 2014 12:31:42 -0800 Received: by build5-lucid-cs (Postfix, from userid 49978) id 426F83219B8; Fri, 21 Feb 2014 12:32:28 -0800 (PST) From: Thomas Schwinge To: CC: Subject: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA. Date: Fri, 21 Feb 2014 21:32:14 +0100 Message-ID: <1393014736-19719-2-git-send-email-thomas@codesourcery.com> In-Reply-To: <1393014736-19719-1-git-send-email-thomas@codesourcery.com> References: <877g8os0vx.fsf@kepler.schwinge.homeip.net> <1393014736-19719-1-git-send-email-thomas@codesourcery.com> MIME-Version: 1.0 From: tschwinge gcc/ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA. (is_gimple_omp_oacc_specifically): Handle it. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise. * omp-low.c (scan_sharing_clauses, scan_omp_target) (expand_omp_target, lower_omp_target, lower_omp_1): Likewise. * gimple.def (GIMPLE_OMP_TARGET): Update comment. * gimple.c (gimple_build_omp_target): Likewise. (gimple_copy): Catch unimplemented case. * tree-inline.c (remap_gimple_stmt): Likewise. * tree-nested.c (convert_nonlocal_reference_stmt) (convert_local_reference_stmt, convert_gimple_call): Likewise. * oacc-builtins.def (BUILT_IN_GOACC_DATA_START) (BUILT_IN_GOACC_DATA_END): New builtins. libgomp/ * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New functions. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208016 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog.gomp | 15 ++++++ gcc/gimple-pretty-print.c | 3 ++ gcc/gimple.c | 4 +- gcc/gimple.def | 1 + gcc/gimple.h | 9 ++++ gcc/gimplify.c | 33 +++++++++--- gcc/oacc-builtins.def | 6 ++- gcc/omp-low.c | 132 ++++++++++++++++++++++++++++++++++++---------- gcc/tree-inline.c | 1 + gcc/tree-nested.c | 3 ++ libgomp/ChangeLog.gomp | 7 +++ libgomp/libgomp.map | 2 + libgomp/libgomp_g.h | 3 ++ libgomp/oacc-parallel.c | 34 +++++++++++- 14 files changed, 213 insertions(+), 40 deletions(-) diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index bd46f2e..824ec94 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,5 +1,20 @@ 2014-02-21 Thomas Schwinge + * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA. + (is_gimple_omp_oacc_specifically): Handle it. + * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. + * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise. + * omp-low.c (scan_sharing_clauses, scan_omp_target) + (expand_omp_target, lower_omp_target, lower_omp_1): Likewise. + * gimple.def (GIMPLE_OMP_TARGET): Update comment. + * gimple.c (gimple_build_omp_target): Likewise. + (gimple_copy): Catch unimplemented case. + * tree-inline.c (remap_gimple_stmt): Likewise. + * tree-nested.c (convert_nonlocal_reference_stmt) + (convert_local_reference_stmt, convert_gimple_call): Likewise. + * oacc-builtins.def (BUILT_IN_GOACC_DATA_START) + (BUILT_IN_GOACC_DATA_END): New builtins. + * omp-low.c (scan_sharing_clauses): Catch unexpected occurrences of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP. diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c index 91a3eb2..ad9369c 100644 --- gcc/gimple-pretty-print.c +++ gcc/gimple-pretty-print.c @@ -1289,6 +1289,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags) case GF_OMP_TARGET_KIND_UPDATE: kind = " update"; break; + case GF_OMP_TARGET_KIND_OACC_DATA: + kind = " oacc_data"; + break; default: gcc_unreachable (); } diff --git gcc/gimple.c gcc/gimple.c index 2a967aa..30561b1 100644 --- gcc/gimple.c +++ gcc/gimple.c @@ -1051,7 +1051,8 @@ gimple_build_omp_single (gimple_seq body, tree clauses) /* Build a GIMPLE_OMP_TARGET statement. BODY is the sequence of statements that will be executed. - CLAUSES are any of the OMP target construct's clauses. */ + KIND is the kind of target region. + CLAUSES are any of the construct's clauses. */ gimple gimple_build_omp_target (gimple_seq body, int kind, tree clauses) @@ -1747,6 +1748,7 @@ gimple_copy (gimple stmt) case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: copy_omp_body: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); new_seq = gimple_seq_copy (gimple_omp_body (stmt)); gimple_omp_set_body (copy, new_seq); break; diff --git gcc/gimple.def gcc/gimple.def index 2b78c06..ce800bd 100644 --- gcc/gimple.def +++ gcc/gimple.def @@ -360,6 +360,7 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE) DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT) /* GIMPLE_OMP_TARGET represents + #pragma acc data #pragma omp target {,data,update} BODY is the sequence of statements inside the target construct (NULL for target update). diff --git gcc/gimple.h gcc/gimple.h index 0d250ef..b4ee9fa 100644 --- gcc/gimple.h +++ gcc/gimple.h @@ -102,6 +102,7 @@ enum gf_mask { GF_OMP_TARGET_KIND_REGION = 0 << 0, GF_OMP_TARGET_KIND_DATA = 1 << 0, GF_OMP_TARGET_KIND_UPDATE = 2 << 0, + GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0, /* True on an GIMPLE_OMP_RETURN statement if the return does not require a thread synchronization via some sort of barrier. The exact barrier @@ -5684,6 +5685,14 @@ is_gimple_omp_oacc_specifically (const_gimple stmt) { case GIMPLE_OACC_PARALLEL: return true; + case GIMPLE_OMP_TARGET: + switch (gimple_omp_target_kind (stmt)) + { + case GF_OMP_TARGET_KIND_OACC_DATA: + return true; + default: + return false; + } default: return false; } diff --git gcc/gimplify.c gcc/gimplify.c index 9aa9301c..fd4305c 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7023,9 +7023,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p) return GS_ALL_DONE; } -/* Gimplify the gross structure of other OpenMP constructs. - In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA - and OMP_TEAMS. */ +/* Gimplify the gross structure of several OpenACC or OpenMP constructs. */ static void gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) @@ -7033,12 +7031,17 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) tree expr = *expr_p; gimple stmt; gimple_seq body = NULL; - enum omp_region_type ort = ORT_WORKSHARE; + enum omp_region_type ort; switch (TREE_CODE (expr)) { + case OACC_DATA: + ort = (enum omp_region_type) (ORT_TARGET + | ORT_TARGET_MAP_FORCE); + break; case OMP_SECTIONS: case OMP_SINGLE: + ort = ORT_WORKSHARE; break; case OMP_TARGET: ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD); @@ -7063,9 +7066,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) pop_gimplify_context (NULL); if (!(ort & ORT_TARGET_OFFLOAD)) { - gimple_seq cleanup = NULL; - tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA); + enum built_in_function end_ix; + switch (TREE_CODE (expr)) + { + case OACC_DATA: + end_ix = BUILT_IN_GOACC_DATA_END; + break; + case OMP_TARGET_DATA: + end_ix = BUILT_IN_GOMP_TARGET_END_DATA; + break; + default: + gcc_unreachable (); + } + tree fn = builtin_decl_explicit (end_ix); g = gimple_build_call (fn, 0); + gimple_seq cleanup = NULL; gimple_seq_add_stmt (&cleanup, g); g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY); body = NULL; @@ -7078,6 +7093,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) switch (TREE_CODE (expr)) { + case OACC_DATA: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA, + OACC_DATA_CLAUSES (expr)); + break; case OMP_SECTIONS: stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); break; @@ -8047,7 +8066,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, break; case OACC_KERNELS: - case OACC_DATA: case OACC_HOST_DATA: case OACC_DECLARE: case OACC_UPDATE: @@ -8076,6 +8094,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, ret = gimplify_omp_for (expr_p, pre_p); break; + case OACC_DATA: case OMP_SECTIONS: case OMP_SINGLE: case OMP_TARGET: diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def index a75e42d..eaf3228 100644 --- gcc/oacc-builtins.def +++ gcc/oacc-builtins.def @@ -1,7 +1,7 @@ /* This file contains the definitions and documentation for the OpenACC builtins used in the GNU compiler. - Copyright (C) 2013 Free Software Foundation, Inc. + Copyright (C) 2013-2014 Free Software Foundation, Inc. Contributed by Thomas Schwinge . @@ -29,3 +29,7 @@ along with GCC; see the file COPYING3. If not see DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel", BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start", + BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST) +DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end", + BT_FN_VOID, ATTR_NOTHROW_LIST) diff --git gcc/omp-low.c gcc/omp-low.c index bca4599..6dec687 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) { tree c, decl; bool scan_array_reductions = false; + bool offloaded; + switch (gimple_code (ctx->stmt)) + { + case GIMPLE_OACC_PARALLEL: + offloaded = true; + break; + case GIMPLE_OMP_TARGET: + switch (gimple_omp_target_kind (ctx->stmt)) + { + case GF_OMP_TARGET_KIND_REGION: + offloaded = true; + break; + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_UPDATE: + case GF_OMP_TARGET_KIND_OACC_DATA: + offloaded = false; + break; + default: + gcc_unreachable (); + } + break; + default: + offloaded = false; + } for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { @@ -1669,11 +1693,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER) { /* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in - #pragma omp target data, there is nothing to map for + target regions that are not offloaded; there is nothing to map for those. */ - if (!gimple_code_is_oacc (ctx->stmt) - && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA - && !POINTER_TYPE_P (TREE_TYPE (decl))) + if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl))) break; } if (DECL_P (decl)) @@ -1698,9 +1720,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) install_var_field (decl, true, 7, ctx); else install_var_field (decl, true, 3, ctx); - if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL - || (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_REGION)) + if (offloaded) install_var_local (decl, ctx); } } @@ -1824,8 +1844,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET || (gimple_omp_target_kind (ctx->stmt) != GF_OMP_TARGET_KIND_UPDATE)); - if (!gimple_code_is_oacc (ctx->stmt) - && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA) + if (!offloaded) break; decl = OMP_CLAUSE_DECL (c); if (DECL_P (decl) @@ -2340,7 +2359,7 @@ scan_omp_single (gimple stmt, omp_context *outer_ctx) layout_type (ctx->record_type); } -/* Scan an OpenMP target{, data, update} directive. */ +/* Scan a GIMPLE_OMP_TARGET. */ static void scan_omp_target (gimple stmt, omp_context *outer_ctx) @@ -2349,6 +2368,12 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx) tree name; int kind = gimple_omp_target_kind (stmt); + if (kind == GF_OMP_TARGET_KIND_OACC_DATA) + { + gcc_assert (taskreg_nesting_level == 0); + gcc_assert (target_nesting_level == 0); + } + ctx = new_omp_context (stmt, outer_ctx); ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0); ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED; @@ -8218,7 +8243,7 @@ expand_omp_atomic (struct omp_region *region) } -/* Expand the OpenMP target{, data, update} directive starting at REGION. */ +/* Expand the GIMPLE_OMP_TARGET starting at REGION. */ static void expand_omp_target (struct omp_region *region) @@ -8401,12 +8426,23 @@ expand_omp_target (struct omp_region *region) clauses = gimple_omp_target_clauses (entry_stmt); - if (kind == GF_OMP_TARGET_KIND_REGION) - start_ix = BUILT_IN_GOMP_TARGET; - else if (kind == GF_OMP_TARGET_KIND_DATA) - start_ix = BUILT_IN_GOMP_TARGET_DATA; - else - start_ix = BUILT_IN_GOMP_TARGET_UPDATE; + switch (kind) + { + case GF_OMP_TARGET_KIND_REGION: + start_ix = BUILT_IN_GOMP_TARGET; + break; + case GF_OMP_TARGET_KIND_DATA: + start_ix = BUILT_IN_GOMP_TARGET_DATA; + break; + case GF_OMP_TARGET_KIND_UPDATE: + start_ix = BUILT_IN_GOMP_TARGET_UPDATE; + break; + case GF_OMP_TARGET_KIND_OACC_DATA: + start_ix = BUILT_IN_GOACC_DATA_START; + break; + default: + gcc_unreachable (); + } /* By default, the value of DEVICE is -1 (let runtime library choose) and there is no conditional. */ @@ -8414,10 +8450,12 @@ expand_omp_target (struct omp_region *region) device = build_int_cst (integer_type_node, -1); c = find_omp_clause (clauses, OMP_CLAUSE_IF); + gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA); if (c) cond = OMP_CLAUSE_IF_EXPR (c); c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE); + gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA); if (c) { device = OMP_CLAUSE_DEVICE_ID (c); @@ -8433,6 +8471,7 @@ expand_omp_target (struct omp_region *region) (cond ? device : -2). */ if (cond) { + gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA); cond = gimple_boolify (cond); basic_block cond_bb, then_bb, else_bb; @@ -8523,7 +8562,9 @@ expand_omp_target (struct omp_region *region) gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET); gsi_remove (&gsi, true); } - if (kind == GF_OMP_TARGET_KIND_DATA && region->exit) + if ((kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) + && region->exit) { gsi = gsi_last_bb (region->exit); g = gsi_stmt (gsi); @@ -10277,7 +10318,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) } } -/* Lower the OpenMP target directive in the current statement +/* Lower the GIMPLE_OMP_TARGET in the current statement in GSI_P. CTX holds context information for the directive. */ static void @@ -10298,7 +10339,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt)); tgt_body = gimple_bind_body (tgt_bind); } - else if (kind == GF_OMP_TARGET_KIND_DATA) + else if (kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) tgt_body = gimple_omp_body (stmt); child_fn = ctx->cb.dst_fn; @@ -10322,6 +10364,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_MAP_TOFROM: case OMP_CLAUSE_MAP_POINTER: break; + case OMP_CLAUSE_MAP_FORCE_ALLOC: + case OMP_CLAUSE_MAP_FORCE_TO: + case OMP_CLAUSE_MAP_FORCE_FROM: + case OMP_CLAUSE_MAP_FORCE_TOFROM: + case OMP_CLAUSE_MAP_FORCE_PRESENT: + case OMP_CLAUSE_MAP_FORCE_DEALLOC: + case OMP_CLAUSE_MAP_FORCE_DEVICEPTR: + gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA); + break; default: gcc_unreachable (); } @@ -10330,6 +10381,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) + gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA); var = OMP_CLAUSE_DECL (c); if (!DECL_P (var)) { @@ -10373,7 +10426,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) lower_omp (&tgt_body, ctx); target_nesting_level--; } - else if (kind == GF_OMP_TARGET_KIND_DATA) + else if (kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) lower_omp (&tgt_body, ctx); if (kind == GF_OMP_TARGET_KIND_REGION) @@ -10400,9 +10454,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1; TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1; TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1; + tree tkind_type; + int talign_shift; + switch (kind) + { + case GF_OMP_TARGET_KIND_REGION: + case GF_OMP_TARGET_KIND_DATA: + case GF_OMP_TARGET_KIND_UPDATE: + tkind_type = unsigned_char_type_node; + talign_shift = 3; + break; + case GF_OMP_TARGET_KIND_OACC_DATA: + tkind_type = short_unsigned_type_node; + talign_shift = 8; + break; + default: + gcc_unreachable (); + } TREE_VEC_ELT (t, 2) - = create_tmp_var (build_array_type_nelts (unsigned_char_type_node, - map_cnt), + = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt), ".omp_data_kinds"); DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1; TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1; @@ -10515,7 +10585,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (TREE_CODE (s) != INTEGER_CST) TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0; - unsigned char tkind = 0; + unsigned HOST_WIDE_INT tkind; switch (OMP_CLAUSE_CODE (c)) { case OMP_CLAUSE_MAP: @@ -10530,14 +10600,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) default: gcc_unreachable (); } - unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); + gcc_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift)); + unsigned HOST_WIDE_INT talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar)); if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign) talign = DECL_ALIGN_UNIT (ovar); talign = ceil_log2 (talign); - tkind |= talign << 3; + tkind |= talign << talign_shift; + gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type))); CONSTRUCTOR_APPEND_ELT (vkind, purpose, - build_int_cst (unsigned_char_type_node, - tkind)); + build_int_cstu (tkind_type, tkind)); if (nc && nc != c) c = nc; } @@ -10589,7 +10660,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_seq (&new_body, tgt_body); new_body = maybe_catch_exception (new_body); } - else if (kind == GF_OMP_TARGET_KIND_DATA) + else if (kind == GF_OMP_TARGET_KIND_DATA + || kind == GF_OMP_TARGET_KIND_OACC_DATA) new_body = tgt_body; if (kind != GF_OMP_TARGET_KIND_UPDATE) { @@ -10810,6 +10882,8 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GIMPLE_OMP_TARGET: ctx = maybe_lookup_ctx (stmt); gcc_assert (ctx); + if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA) + gcc_assert (!ctx->cancellable); lower_omp_target (gsi_p, ctx); break; case GIMPLE_OMP_TEAMS: diff --git gcc/tree-inline.c gcc/tree-inline.c index 99903333..61c1cc8 100644 --- gcc/tree-inline.c +++ gcc/tree-inline.c @@ -1397,6 +1397,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id) break; case GIMPLE_OMP_TARGET: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); s1 = remap_gimple_seq (gimple_omp_body (stmt), id); copy = gimple_build_omp_target (s1, gimple_omp_target_kind (stmt), diff --git gcc/tree-nested.c gcc/tree-nested.c index 8933d02..afa7abb 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1307,6 +1307,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); save_suppress = info->suppress_expansion; convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, @@ -1769,6 +1770,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); save_suppress = info->suppress_expansion; convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi); walk_body (convert_local_reference_stmt, convert_local_reference_op, @@ -2184,6 +2186,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p, case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_CRITICAL: + gcc_assert (!is_gimple_omp_oacc_specifically (stmt)); walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt)); break; diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 3dffde4..5c15656 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,10 @@ +2014-02-21 Thomas Schwinge + + * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start. + * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes. + * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New + functions. + 2014-02-20 Thomas Schwinge * target.c (gomp_load_plugin_for_device): Don't call dlcose if diff --git libgomp/libgomp.map libgomp/libgomp.map index 2b64d05..cb52e45 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -233,5 +233,7 @@ OACC_2.0 { GOACC_2.0 { global: + GOACC_data_end; + GOACC_data_start; GOACC_parallel; }; diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h index 7c24317..b9083a5 100644 --- libgomp/libgomp_g.h +++ libgomp/libgomp_g.h @@ -218,5 +218,8 @@ extern void GOMP_teams (unsigned int, unsigned int); extern void GOACC_parallel (int, void (*) (void *), const void *, size_t, void **, size_t *, unsigned short *); +extern void GOACC_data_start (int, const void *, + size_t, void **, size_t *, unsigned short *); +extern void GOACC_data_end (void); #endif /* LIBGOMP_G_H */ diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c index bf7b74c..3ac7e39 100644 --- libgomp/oacc-parallel.c +++ libgomp/oacc-parallel.c @@ -1,4 +1,4 @@ -/* Copyright (C) 2013 Free Software Foundation, Inc. +/* Copyright (C) 2013-2014 Free Software Foundation, Inc. Contributed by Thomas Schwinge . @@ -23,7 +23,7 @@ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see . */ -/* This file handles the OpenACC parallel construct. */ +/* This file handles the OpenACC data and parallel constructs. */ #include "libgomp.h" #include "libgomp_g.h" @@ -51,3 +51,33 @@ GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target, } GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds_); } + + +void +GOACC_data_start (int device, const void *openmp_target, size_t mapnum, + void **hostaddrs, size_t *sizes, unsigned short *kinds) +{ + unsigned char kinds_[mapnum]; + size_t i; + + /* TODO. Eventually, we'll be interpreting all mapping kinds according to + the OpenACC semantics; for now we're re-using what is implemented for + OpenMP. */ + for (i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i]; + unsigned char align = kinds[i] >> 8; + if (kind > 4) + gomp_fatal ("memory mapping kind %x for %zd is not yet supported", + kind, i); + + kinds_[i] = kind | align << 3; + } + GOMP_target_data (device, openmp_target, mapnum, hostaddrs, sizes, kinds_); +} + +void +GOACC_data_end (void) +{ + GOMP_target_end_data (); +}