From patchwork Mon Nov 30 23:12:46 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 550551 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 72CEA1401B5 for ; Tue, 1 Dec 2015 10:13:01 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=DNtg0zd3; 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:subject:message-id:mime-version:content-type; q=dns; s= default; b=MkFUTzqViQ/r8O1Fkq9deZcWiPOlWyZB2kycD9KAQpzy56Oar3gmX ZRq9ieNg/Az6+nYkq1vgP2Z30RKtv+qJg2XCQbKtofeAjo9Y+GGdNPf9W3HeqL5k bhb9IrGP9FMIaVkRa0RlnHuAHzmoHH4e1fe7GqPTDCOOaBGl0Qi81E= 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:subject:message-id:mime-version:content-type; s= default; bh=Wh+91LvJA0Ql5Er6TLTxRNlEFh8=; b=DNtg0zd3VOuP+HVHIHeo z/3C93v1YMQ8iQwWk+ASKV1FhlNYZF0GwqL/JkyBpvIsApt2eFFjVLfhVakWhPkE CkpGl/4cPRRfv2vf6k0iPGScsSO3wl3pMPWgsOGUO3q32dylZ0WNdI716PlnzysY ST5a6y5YW5DOSZ4IXyPd+9E= Received: (qmail 11477 invoked by alias); 30 Nov 2015 23:12:53 -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 11459 invoked by uid 89); 30 Nov 2015 23:12:52 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.2 required=5.0 tests=AWL, BAYES_40, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.2 X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (CAMELLIA256-SHA encrypted) ESMTPS; Mon, 30 Nov 2015 23:12:50 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id E0939AAC4 for ; Mon, 30 Nov 2015 23:12:46 +0000 (UTC) Date: Tue, 1 Dec 2015 00:12:46 +0100 From: Martin Jambor To: GCC Patches Subject: [hsa] Describe grid with target clauses Message-ID: <20151130231246.GA19649@virgil.suse.cz> Mail-Followup-To: GCC Patches MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.5.24 (2015-08-30) X-IsSubscribed: yes Hi, Jakub requested that I remove the grid description from new fields of the classes representing gimple omp statement and put them into special artificial clauses instead. This patch implement that, with one target clause per dimension (so up to three clauses) and each one describing both the grid size and group size along that dimension (hence the new clause type has two parameters). Committed to the branch, I will be preparing a new diff against the trunk shortly. Thanks, Martin 2015-11-30 Martin Jambor * gimple.c (gimple_omp_target_init_dimensions): Removed. * gimple.h (gimple_statement_omp_parallel_layout): Removed fields dimensions and kernel_dim. (gimple_omp_target_dimensions): Removed. (gimple_omp_target_grid_size): Likewise. (gimple_omp_target_grid_size_ptr): Likewise. (gimple_omp_target_set_grid_size): Likewise. (gimple_omp_target_workgroup_size): Likewise. (gimple_omp_target_workgroup_size_ptr): Likewise. (gimple_omp_target_set_workgroup_size): Likewise. * omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE__GRIDDIM_. (scan_omp_target): Do not scan kernel_dim. (region_needs_kernel_p): Use clauses to recognize gridified kernels. (get_kernel_launch_attributes): Generate launch attributes from clauses. (get_target_arguments): Use clauses to recognize gridified kernels. (expand_target_kernel_body): Likewise. (attempt_target_gridification): Record grid description into clauses. * tree-core.h (omp_clause_code): New element OMP_CLAUSE__GRIDDIM_. (tree_omp_clause): New subcode dimension. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE__GRIDDIM_. * tree.c (omp_clause_num_ops): Add number of opernads of OMP_CLAUSE__GRIDDIM_. (omp_clause_code_name): Add name of OMP_CLAUSE__GRIDDIM_. (walk_tree_1): Handle OMP_CLAUSE__GRIDDIM_. * tree.h (OMP_CLAUSE_GRIDDIM_DIMENSION): New. (OMP_CLAUSE_SET_GRIDDIM_DIMENSION): Likewise. (OMP_CLAUSE_GRIDDIM_SIZE): Likewise. (OMP_CLAUSE_GRIDDIM_GROUP): Likewise. --- gcc/gimple.c | 11 ------- gcc/gimple.h | 82 ------------------------------------------------- gcc/omp-low.c | 72 ++++++++++++++++++++++++++----------------- gcc/tree-core.h | 9 +++++- gcc/tree-pretty-print.c | 12 ++++++++ gcc/tree.c | 5 ++- gcc/tree.h | 11 +++++++ 7 files changed, 79 insertions(+), 123 deletions(-) diff --git a/gcc/gimple.c b/gcc/gimple.c index d876e90..4658f29 100644 --- a/gcc/gimple.c +++ b/gcc/gimple.c @@ -1098,17 +1098,6 @@ gimple_build_omp_target (gimple_seq body, int kind, tree clauses) return p; } -/* Set dimensions of TARGET to NUM and allocate kernel_dim array of the - statement with the appropriate number of elements. */ - -void -gimple_omp_target_init_dimensions (gomp_target *target, size_t num) -{ - gcc_assert (num > 0); - target->dimensions = num; - target->kernel_dim = ggc_cleared_vec_alloc (num); -} - /* Build a GIMPLE_OMP_TEAMS statement. BODY is the sequence of statements that will be executed. diff --git a/gcc/gimple.h b/gcc/gimple.h index 14e6cf6..4c4c799 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -661,21 +661,7 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) Shared data argument. */ tree data_arg; - /* TODO: Revisit placement of the following two fields. On one hand, we - currently only use them on target construct. On the other, use on - parallel construct is also possible in the future. */ - /* [ WORD 11 ] */ - /* Number of elements in kernel_iter array. */ - size_t dimensions; - - /* [ WORD 12 ] */ - /* If target also contains a GPU kernel, it should be run with the - following grid sizes. */ - struct gimple_omp_target_grid_dim - * GTY((length ("%h.dimensions"))) kernel_dim; - - /* [ WORD 13 ] */ /* If set, this statement is part of a gridified kernel, its clauses need to be scanned and lowered but the statement should be discarded after lowering. */ @@ -1504,7 +1490,6 @@ gomp_sections *gimple_build_omp_sections (gimple_seq, tree); gimple *gimple_build_omp_sections_switch (void); gomp_single *gimple_build_omp_single (gimple_seq, tree); gomp_target *gimple_build_omp_target (gimple_seq, int, tree); -void gimple_omp_target_init_dimensions (gomp_target *, size_t); gomp_teams *gimple_build_omp_teams (gimple_seq, tree); gomp_atomic_load *gimple_build_omp_atomic_load (tree, tree); gomp_atomic_store *gimple_build_omp_atomic_store (tree); @@ -5683,73 +5668,6 @@ gimple_omp_target_set_data_arg (gomp_target *omp_target_stmt, omp_target_stmt->data_arg = data_arg; } -/* Return the number of dimensions of kernel grid. */ - -static inline size_t -gimple_omp_target_dimensions (gomp_target *omp_target_stmt) -{ - return omp_target_stmt->dimensions; -} - -/* Return the size of kernel grid of OMP_TARGET_STMT along dimension N. */ - -static inline tree -gimple_omp_target_grid_size (gomp_target *omp_target_stmt, unsigned n) -{ - gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n); - return omp_target_stmt->kernel_dim[n].grid_size; -} - -/* Return pointer to tree specifying the size of kernel grid of OMP_TARGET_STMT - along dimension N. */ - -static inline tree * -gimple_omp_target_grid_size_ptr (gomp_target *omp_target_stmt, unsigned n) -{ - gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n); - return &omp_target_stmt->kernel_dim[n].grid_size; -} - -/* Set the size of kernel grid of OMP_TARGET_STMT along dimension N to V */ - -static inline void -gimple_omp_target_set_grid_size (gomp_target *omp_target_stmt, unsigned n, - tree v) -{ - gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n); - omp_target_stmt->kernel_dim[n].grid_size = v; -} - -/* Return the size of kernel work group of OMP_TARGET_STMT along dimension N. */ - -static inline tree -gimple_omp_target_workgroup_size (gomp_target *omp_target_stmt, unsigned n) -{ - gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n); - return omp_target_stmt->kernel_dim[n].workgroup_size; -} - -/* Return pointer to tree specifying the size of kernel work group of - OMP_TARGET_STMT along dimension N. */ - -static inline tree * -gimple_omp_target_workgroup_size_ptr (gomp_target *omp_target_stmt, unsigned n) -{ - gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n); - return &omp_target_stmt->kernel_dim[n].workgroup_size; -} - -/* Set the size of kernel workgroup of OMP_TARGET_STMT along dimension N to - V */ - -static inline void -gimple_omp_target_set_workgroup_size (gomp_target *omp_target_stmt, unsigned n, - tree v) -{ - gcc_assert (gimple_omp_target_dimensions (omp_target_stmt) > n); - omp_target_stmt->kernel_dim[n].workgroup_size = v; -} - /* Return the clauses associated with OMP_TEAMS GS. */ static inline tree diff --git a/gcc/omp-low.c b/gcc/omp-low.c index f1d10a2..5933c60 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2140,6 +2140,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) } break; + case OMP_CLAUSE__GRIDDIM_: + if (ctx->outer) + { + scan_omp_op (&OMP_CLAUSE_GRIDDIM_SIZE (c), ctx->outer); + scan_omp_op (&OMP_CLAUSE_GRIDDIM_GROUP (c), ctx->outer); + } + break; + case OMP_CLAUSE_NOWAIT: case OMP_CLAUSE_ORDERED: case OMP_CLAUSE_COLLAPSE: @@ -2336,6 +2344,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE__GRIDDIM_: break; case OMP_CLAUSE_DEVICE_RESIDENT: @@ -3088,12 +3097,6 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; - for (size_t i = 0; i < gimple_omp_target_dimensions (stmt); i++) - { - scan_omp_op (gimple_omp_target_grid_size_ptr (stmt, i), ctx); - scan_omp_op (gimple_omp_target_workgroup_size_ptr (stmt, i), ctx); - } - if (offloaded) { create_omp_child_function (ctx, false); @@ -6310,7 +6313,9 @@ region_needs_kernel_p (struct omp_region *region) { gomp_target *tgt_stmt; tgt_stmt = as_a (last_stmt (region->entry)); - if (gimple_omp_target_dimensions (tgt_stmt)) + + if (find_omp_clause (gimple_omp_target_clauses (tgt_stmt), + OMP_CLAUSE__GRIDDIM_)) return indirect; else return true; @@ -12624,26 +12629,30 @@ get_kernel_launch_attributes (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt) tree u32_one = build_one_cst (uint32_type_node); tree lattrs = create_tmp_var (kernel_launch_attributes_type, "__kernel_launch_attrs"); + + unsigned max_dim = 0; + for (tree clause = gimple_omp_target_clauses (tgt_stmt); + clause; + clause = OMP_CLAUSE_CHAIN (clause)) + { + if (OMP_CLAUSE_CODE (clause) != OMP_CLAUSE__GRIDDIM_) + continue; + + unsigned dim = OMP_CLAUSE_GRIDDIM_DIMENSION (clause); + max_dim = MAX (dim, max_dim); + + insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, dim, + OMP_CLAUSE_GRIDDIM_SIZE (clause)); + insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, dim, + OMP_CLAUSE_GRIDDIM_GROUP (clause)); + } + tree dimref = build3 (COMPONENT_REF, uint32_type_node, lattrs, kernel_lattrs_dimnum_decl, NULL_TREE); /* At this moment we cannot gridify a loop with a collapse clause. */ /* TODO: Adjust when we support bigger collapse. */ - gcc_assert (gimple_omp_target_dimensions (tgt_stmt) == 1); + gcc_assert (max_dim == 0); gsi_insert_before (gsi, gimple_build_assign (dimref, u32_one), GSI_SAME_STMT); - - /* Calculation of grid size: */ - insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, 0, - gimple_omp_target_grid_size (tgt_stmt, 0)); - insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, 0, - gimple_omp_target_workgroup_size (tgt_stmt, 0)); - insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, 1, - u32_one); - insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, 2, - u32_one); - insert_store_range_dim (gsi, lattrs, kernel_lattrs_grid_decl, 2, - u32_one); - insert_store_range_dim (gsi, lattrs, kernel_lattrs_group_decl, 1, - u32_one); TREE_ADDRESSABLE (lattrs) = 1; return build_fold_addr_expr (lattrs); } @@ -12717,7 +12726,8 @@ get_target_arguments (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt) args.quick_push (t); /* Add HSA-specific grid sizes, if available. */ - if (gimple_omp_target_dimensions (tgt_stmt)) + if (find_omp_clause (gimple_omp_target_clauses (tgt_stmt), + OMP_CLAUSE__GRIDDIM_)) { t = get_target_argument_identifier (GOMP_DEVICE_HSA, true, GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES); @@ -13392,14 +13402,16 @@ expand_target_kernel_body (struct omp_region *target) if (gimple_omp_target_kind (tgt_stmt) != GF_OMP_TARGET_KIND_REGION) return; gcc_checking_assert (orig_child_fndecl); - gcc_assert (!gimple_omp_target_dimensions (tgt_stmt)); + gcc_assert (!find_omp_clause (gimple_omp_target_clauses (tgt_stmt), + OMP_CLAUSE__GRIDDIM_)); cgraph_node *n = cgraph_node::get (orig_child_fndecl); hsa_register_kernel (n); return; } - gcc_assert (gimple_omp_target_dimensions (tgt_stmt)); + gcc_assert (find_omp_clause (gimple_omp_target_clauses (tgt_stmt), + OMP_CLAUSE__GRIDDIM_)); tree inside_block = gimple_block (first_stmt (single_succ (gpukernel->entry))); *pp = gpukernel->next; for (pp = &gpukernel->inner; *pp; pp = &(*pp)->next) @@ -17470,7 +17482,6 @@ attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi, walk_tree (&group_size, remap_prebody_decls, &wi, NULL); size_t collapse = gimple_omp_for_collapse (inner_loop); - gimple_omp_target_init_dimensions (target, collapse); for (size_t i = 0; i < collapse; i++) { gimple_omp_for_iter iter = inner_loop->iter[i]; @@ -17506,7 +17517,6 @@ attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi, t = fold_convert (uint32_type_node, t); tree gs = force_gimple_operand_gsi (gsi, t, true, NULL_TREE, true, GSI_SAME_STMT); - gimple_omp_target_set_grid_size (target, i, gs); tree ws; if (i == 0 && group_size) { @@ -17516,7 +17526,13 @@ attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi, } else ws = build_zero_cst (uint32_type_node); - gimple_omp_target_set_workgroup_size (target, i, ws); + + tree c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE__GRIDDIM_); + OMP_CLAUSE_SET_GRIDDIM_DIMENSION (c, (unsigned int) i); + OMP_CLAUSE_GRIDDIM_SIZE (c) = gs; + OMP_CLAUSE_GRIDDIM_GROUP (c) = ws; + OMP_CLAUSE_CHAIN (c) = gimple_omp_target_clauses (target); + gimple_omp_target_set_clauses (target, c); } delete declmap; diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 9cc64d9..858f220 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -460,7 +460,11 @@ enum omp_clause_code { OMP_CLAUSE_VECTOR_LENGTH, /* OpenACC clause: tile ( size-expr-list ). */ - OMP_CLAUSE_TILE + OMP_CLAUSE_TILE, + + /* OpenMP internal-only clause to specify grid dimensions of a gridified + kernel. */ + OMP_CLAUSE__GRIDDIM_ }; #undef DEFTREESTRUCT @@ -1377,6 +1381,9 @@ struct GTY(()) tree_omp_clause { enum tree_code reduction_code; enum omp_clause_linear_kind linear_kind; enum tree_code if_modifier; + /* The dimension a OMP_CLAUSE__GRIDDIM_ clause of a gridified target + construct describes. */ + unsigned int dimension; } GTY ((skip)) subcode; /* The gimplification of OMP_CLAUSE_REDUCTION_{INIT,MERGE} for omp-low's diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index caec760..ad5cfdb 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -945,6 +945,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) pp_right_paren (pp); break; + case OMP_CLAUSE__GRIDDIM_: + pp_string (pp, "_griddim_("); + pp_unsigned_wide_integer (pp, OMP_CLAUSE_GRIDDIM_DIMENSION (clause)); + pp_colon (pp); + dump_generic_node (pp, OMP_CLAUSE_GRIDDIM_SIZE (clause), spc, flags, + false); + pp_comma (pp); + dump_generic_node (pp, OMP_CLAUSE_GRIDDIM_GROUP (clause), spc, flags, + false); + pp_right_paren (pp); + break; + default: /* Should never happen. */ dump_generic_node (pp, clause, spc, flags, false); diff --git a/gcc/tree.c b/gcc/tree.c index 2387deb..3a74982 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -329,6 +329,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_NUM_WORKERS */ 1, /* OMP_CLAUSE_VECTOR_LENGTH */ 1, /* OMP_CLAUSE_TILE */ + 2, /* OMP_CLAUSE__GRIDDIM_ */ }; const char * const omp_clause_code_name[] = @@ -400,7 +401,8 @@ const char * const omp_clause_code_name[] = "num_gangs", "num_workers", "vector_length", - "tile" + "tile", + "griddim" }; @@ -11603,6 +11605,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, switch (OMP_CLAUSE_CODE (*tp)) { case OMP_CLAUSE_GANG: + case OMP_CLAUSE__GRIDDIM_: WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 1)); /* FALLTHRU */ diff --git a/gcc/tree.h b/gcc/tree.h index 0c1602e..7b9bcb3 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1636,6 +1636,17 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_TILE_LIST(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_TILE), 0) +#define OMP_CLAUSE_GRIDDIM_DIMENSION(NODE) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_)\ + ->omp_clause.subcode.dimension) +#define OMP_CLAUSE_SET_GRIDDIM_DIMENSION(NODE, DIMENSION) \ + (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_)\ + ->omp_clause.subcode.dimension = (DIMENSION)) +#define OMP_CLAUSE_GRIDDIM_SIZE(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_), 0) +#define OMP_CLAUSE_GRIDDIM_GROUP(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__GRIDDIM_), 1) + /* SSA_NAME accessors. */ /* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE