From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 122169 invoked by alias); 5 Nov 2015 21:57:50 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 122159 invoked by uid 89); 5 Nov 2015 21:57:50 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.5 required=5.0 tests=AWL,BAYES_50,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; Thu, 05 Nov 2015 21:57:38 +0000 Received: from relay2.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 932BCAABB; Thu, 5 Nov 2015 21:57:15 +0000 (UTC) Date: Thu, 05 Nov 2015 21:57:00 -0000 From: Martin Jambor To: GCC Patches Cc: Jakub Jelinek Subject: [hsa 4/12] OpenMP lowering/expansion changes (gridification) Message-ID: <20151105215733.GG9264@virgil.suse.cz> Mail-Followup-To: GCC Patches , Jakub Jelinek References: <20151105215108.GC9264@virgil.suse.cz> MIME-Version: 1.0 Content-Type: text/plain; charset=utf-8 Content-Disposition: inline In-Reply-To: <20151105215108.GC9264@virgil.suse.cz> User-Agent: Mutt/1.5.24 (2015-08-30) X-IsSubscribed: yes X-SW-Source: 2015-11/txt/msg00523.txt.bz2 Hi, the patch in this email contains the changes to make our OpenMP lowering and expansion machinery produce GPU kernels for a certain limited class of loops. The plan is to make that class quite a big bigger, but only the following is ready for submission now. Basically, whenever the compiler configured for HSAIL generation encounters the following pattern: #pragma omp target #pragma omp teams thread_limit(workgroup_size) // thread_limit is optional #pragma omp distribute parallel for firstprivate(n) private(i) other_sharing_clauses() for (i = 0; i < n; i++) some_loop_body it creates a copy of the entire target body and expands it slightly differently for concurrent execution on a GPU. Note that both teams and distribute constructs are mandatory. Moreover, currently the distribute has to be in a combined statement with the inner for construct. And there are quite a few other restrictions which I hope to alleviate over the next year, most notably implement reductions. A few days ago I hoped to finish writing support for collapse(2) and collapse(3) clauses in time for stage1 but now I am a bit sceptical. The first phase of the "gridification" process is run before omp "scanning" phase. We look for the pattern above, and if we encounter one, we copy its entire body into a new gimple statement GIMPLE_OMP_GPUKERNEL. Within it, we mark the teams, distribute and parallel constructs with a new flag "kernel_phony." This flag will then make OMP lowering phase process their sharing clauses like usual, but the statements representing the constructs will be removed at lowering (and thus will never be expanded). The resulting wasteful repackaging of data is nicely cleaned by our optimizers even at -O1. At expansion time, we identify gomp_target statements with a kernel and expand the kernel into a special function, with the loop represented by the GPU grid and not control flow. Afterwards, the normal body of the target is expanded as usual. Finally, we need to take the grid dimensions stored within new fields of the target statement by the first phase, store in a structure and pass them to libgomp in a new parameter of GOMP_target_41. Originally, when I started with the above pattern matching, I did not allow any other gimple statements in between the respective omp constructs. That however proved to be too restrictive for two reasons. First, statements in pre-bodies of both distribute and for loops needed to be accounted for when calculating the kernel grid size (which is done before the target statement itself) and second, Fortran parameter dereferences happily result in interleaving statements when there were none in the user source code. Therefore, I allow register-type stores to local non-addressable variables in pre-bodies and also in between the OMP constructs. All of them are copied in front of the target statement and either used for grid size calculation or removed as useless by later optimizations. For convenience of anybody reviewing the code, I'm attaching a very simple testcase with selection of dumps that illustrate the whole process. While we have also been experimenting quite a bit with dynamic parallelism, we have only been able to achieve any good performance via this process of gridification. The user can be notified whether a particular target construct was gridified or not via our process of dumping notes, which however only appear in the detailed dump. I am seriously considering emitting some kind of warning, when HSA-enabled compiler is about to produce a non-gridified target code. I hope that eventually I managed to write the gridification in a way that interferes very little with the rest of the OMP pipeline and yet only re-implement the bare necessary minimum of functionality that is already there. I'll be grateful for any feedback regarding the approach. Thanks, Martin 2015-11-05 Martin Jambor * builtin-types.def (BT_FN_VOID_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. * fortran/types.def (BT_FN_VOID_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. * gimple-low.c (lower_stmt): Handle GIMPLE_OMP_GPUKERNEL. * gimple-pretty-print.c (dump_gimple_omp_for): Likewise. (dump_gimple_omp_block): Handle GF_OMP_FOR_KIND_KERNEL_BODY (pp_gimple_stmt_1): Handle GIMPLE_OMP_GPUKERNEL. * gimple-walk.c (walk_gimple_stmt): Likewise. * gimple.c (gimple_build_omp_gpukernel): New function. (gimple_omp_target_init_dimensions): Likewise. (gimple_copy): Handle GIMPLE_OMP_GPUKERNEL. * gimple.def (GIMPLE_OMP_TEAMS): Moved into its own layout. (GIMPLE_OMP_GPUKERNEL): New. * gimple.h (gf_mask): New element GF_OMP_FOR_KIND_KERNEL_BODY. (gomp_for): New field kernel_phony. (gimple_omp_target_grid_dim): New type. (gimple_statement_omp_parallel_layout): New fields dimensions, kernel_dim, kernel_phony. (gomp_teams): New field kernel_phony. (gimple_build_omp_gpukernel): Declare. (gimple_omp_target_init_dimensions): Likewise. (gimple_has_substatements): Handle GIMPLE_OMP_GPUKERNEL. (gimple_omp_for_kernel_phony): New function. (gimple_omp_for_set_kernel_phony): Likewise. (gimple_omp_parallel_kernel_phony): Likewise. (gimple_omp_parallel_set_kernel_phony): Likewise. (gimple_omp_target_dimensions): Likewise. (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. (gimple_omp_teams_kernel_phony): Likewise. (gimple_omp_teams_set_kernel_phony): Likewise. (CASE_GIMPLE_OMP): Handle GIMPLE_OMP_GPUKERNEL. * gsstruct.def (GSS_OMP_TEAMS_LAYOUT): New. * omp-builtins.def (BUILT_IN_GOMP_OFFLOAD_REGISTER): Likewise. (BUILT_IN_GOMP_OFFLOAD_UNREGISTER): Likewise. (BUILT_IN_GOMP_TARGET): Changed type. * omp-low.c: Include symbol-summary.h and hsa.h. (adjust_for_condition): New function. (get_omp_for_step_from_incr): Likewise. (extract_omp_for_data): Moved parts to adjust_for_condition and get_omp_for_step_from_incr. (build_outer_var_ref): Handle GIMPLE_OMP_GPUKERNEL. (fixup_child_record_type): Bail out if receiver_decl is NULL. (scan_omp_parallel): Do not create child functions for phony constructs. (scan_omp_target): Scan target dimensions. (check_omp_nesting_restrictions): Handle GIMPLE_OMP_GPUKERNEL. (scan_omp_1_stmt): Likewise. (region_needs_kernel_p): New function. (expand_parallel_call): Register apprpriate parallel child functions as HSA kernels. (kernel_dim_array_type, kernel_lattrs_dimnum_decl): New variables. (kernel_lattrs_grid_decl, kernel_lattrs_group_decl): Likewise. (kernel_launch_attributes_type): Likewise. (create_kernel_launch_attr_types): New function. (insert_store_range_dim): Likewise. (get_kernel_launch_attributes): Likewise. (expand_omp_target): Fill in kernel dimensions, if any. (expand_omp_for_kernel): New function. (arg_decl_map): New type. (remap_kernel_arg_accesses): New function. (expand_omp): New forward declaration. (expand_target_kernel_body): New function. (expand_omp): Call it. (lower_omp_for): Do not emit phony constructs. (lower_omp_for): Likewise. (lower_omp_taskreg): Do not emit phony constructs but create for them a temporary variable receiver_decl. (lower_omp_taskreg): Do not emit phony constructs. (lower_omp_teams): Likewise. (lower_omp_gpukernel): New function. (lower_omp_1): Call it. (reg_assignment_to_local_var_p): New function. (seq_only_contains_local_assignments): Likewise. (find_single_omp_among_assignments_1): Likewise. (find_single_omp_among_assignments): Likewise. (find_ungridifiable_statement): Likewise. (target_follows_gridifiable_pattern): Likewise. (remap_prebody_decls): Likewise. (copy_leading_local_assignments): Likewise. (process_kernel_body_copy): Likewise. (attempt_target_gridification): Likewise. (create_target_gpukernel_stmt): Likewise. (create_target_gpukernels): Likewise. (execute_lower_omp): Call create_target_gpukernels. (make_gimple_omp_edges): Handle GIMPLE_OMP_GPUKERNEL. diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def index b561436..e2fa418 100644 --- a/gcc/builtin-types.def +++ b/gcc/builtin-types.def @@ -450,6 +450,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONG_ULONG_ULONGPTR, BT_BOOL, BT_ULONG, BT_ULONG, BT_PTR_ULONG) DEF_FUNCTION_TYPE_3 (BT_FN_BOOL_ULONGLONG_ULONGLONG_ULONGLONGPTR, BT_BOOL, BT_ULONGLONG, BT_ULONGLONG, BT_PTR_ULONGLONG) +DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_PTR, BT_VOID, BT_PTR, BT_INT, BT_PTR) DEF_FUNCTION_TYPE_4 (BT_FN_SIZE_CONST_PTR_SIZE_SIZE_FILEPTR, BT_SIZE, BT_CONST_PTR, BT_SIZE, BT_SIZE, BT_FILEPTR) @@ -547,13 +548,13 @@ 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) +DEF_FUNCTION_TYPE_9 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR, + BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, + BT_PTR, BT_PTR, BT_UINT, BT_PTR, 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, BT_PTR_FN_VOID_PTR_PTR, BT_LONG, BT_LONG, diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def index ca75654..a9cfc84 100644 --- a/gcc/fortran/types.def +++ b/gcc/fortran/types.def @@ -145,6 +145,7 @@ DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I2_INT, BT_VOID, BT_VOLATILE_PTR, BT_I2, BT DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I4_INT, BT_VOID, BT_VOLATILE_PTR, BT_I4, BT_INT) DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I8_INT, BT_VOID, BT_VOLATILE_PTR, BT_I8, BT_INT) DEF_FUNCTION_TYPE_3 (BT_FN_VOID_VPTR_I16_INT, BT_VOID, BT_VOLATILE_PTR, BT_I16, BT_INT) +DEF_FUNCTION_TYPE_3 (BT_FN_VOID_PTR_INT_PTR, BT_VOID, BT_PTR, BT_INT, BT_PTR) DEF_FUNCTION_TYPE_4 (BT_FN_VOID_OMPFN_PTR_UINT_UINT, BT_VOID, BT_PTR_FN_VOID_PTR, BT_PTR, BT_UINT, BT_UINT) @@ -215,9 +216,9 @@ DEF_FUNCTION_TYPE_7 (BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_UINT_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, +DEF_FUNCTION_TYPE_9 (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR, BT_VOID, BT_INT, BT_PTR_FN_VOID_PTR, BT_SIZE, BT_PTR, - BT_PTR, BT_PTR, BT_UINT, BT_PTR) + BT_PTR, BT_PTR, BT_UINT, BT_PTR, 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, diff --git a/gcc/gimple-low.c b/gcc/gimple-low.c index 4994918..d2a6a80 100644 --- a/gcc/gimple-low.c +++ b/gcc/gimple-low.c @@ -358,6 +358,7 @@ lower_stmt (gimple_stmt_iterator *gsi, struct lower_data *data) case GIMPLE_OMP_TASK: case GIMPLE_OMP_TARGET: case GIMPLE_OMP_TEAMS: + case GIMPLE_OMP_GPUKERNEL: data->cannot_fallthru = false; lower_omp_directive (gsi, data); data->cannot_fallthru = false; diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index 7b50cdf..83498bc 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1187,6 +1187,9 @@ dump_gimple_omp_for (pretty_printer *buffer, gomp_for *gs, int spc, int flags) case GF_OMP_FOR_KIND_CILKSIMD: pp_string (buffer, "#pragma simd"); break; + case GF_OMP_FOR_KIND_KERNEL_BODY: + pp_string (buffer, "#pragma omp for kernel"); + break; default: gcc_unreachable (); } @@ -1488,6 +1491,9 @@ dump_gimple_omp_block (pretty_printer *buffer, gimple *gs, int spc, int flags) case GIMPLE_OMP_SECTION: pp_string (buffer, "#pragma omp section"); break; + case GIMPLE_OMP_GPUKERNEL: + pp_string (buffer, "#pragma omp gpukernel"); + break; default: gcc_unreachable (); } @@ -2270,6 +2276,7 @@ pp_gimple_stmt_1 (pretty_printer *buffer, gimple *gs, int spc, int flags) case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_SECTION: + case GIMPLE_OMP_GPUKERNEL: dump_gimple_omp_block (buffer, gs, spc, flags); break; diff --git a/gcc/gimple-walk.c b/gcc/gimple-walk.c index 850cf57..695592d 100644 --- a/gcc/gimple-walk.c +++ b/gcc/gimple-walk.c @@ -644,6 +644,7 @@ walk_gimple_stmt (gimple_stmt_iterator *gsi, walk_stmt_fn callback_stmt, case GIMPLE_OMP_SINGLE: case GIMPLE_OMP_TARGET: case GIMPLE_OMP_TEAMS: + case GIMPLE_OMP_GPUKERNEL: ret = walk_gimple_seq_mod (gimple_omp_body_ptr (stmt), callback_stmt, callback_op, wi); if (ret) diff --git a/gcc/gimple.c b/gcc/gimple.c index 4ce38da..9eba126 100644 --- a/gcc/gimple.c +++ b/gcc/gimple.c @@ -953,6 +953,19 @@ gimple_build_omp_master (gimple_seq body) return p; } +/* Build a GIMPLE_OMP_GPUKERNEL statement. + + BODY is the sequence of statements to be executed by the kernel. */ + +gimple * +gimple_build_omp_gpukernel (gimple_seq body) +{ + gimple *p = gimple_alloc (GIMPLE_OMP_GPUKERNEL, 0); + if (body) + gimple_omp_set_body (p, body); + + return p; +} /* Build a GIMPLE_OMP_TASKGROUP statement. @@ -1084,6 +1097,16 @@ 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. @@ -1804,6 +1827,7 @@ gimple_copy (gimple *stmt) case GIMPLE_OMP_SECTION: case GIMPLE_OMP_MASTER: case GIMPLE_OMP_TASKGROUP: + case GIMPLE_OMP_GPUKERNEL: copy_omp_body: new_seq = gimple_seq_copy (gimple_omp_body (stmt)); gimple_omp_set_body (copy, new_seq); diff --git a/gcc/gimple.def b/gcc/gimple.def index d3ca402..30f0111 100644 --- a/gcc/gimple.def +++ b/gcc/gimple.def @@ -369,13 +369,17 @@ DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL_LAYOUT) /* GIMPLE_OMP_TEAMS represents #pragma omp teams BODY is the sequence of statements inside the single section. CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ -DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_SINGLE_LAYOUT) +DEFGSCODE(GIMPLE_OMP_TEAMS, "gimple_omp_teams", GSS_OMP_TEAMS_LAYOUT) /* GIMPLE_OMP_ORDERED represents #pragma omp ordered. BODY is the sequence of statements to execute in the ordered section. CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */ DEFGSCODE(GIMPLE_OMP_ORDERED, "gimple_omp_ordered", GSS_OMP_SINGLE_LAYOUT) +/* GIMPLE_OMP_GPUKERNEL represents a parallel loop lowered for execution + on a GPU. It is an artificial statement created by omp lowering. */ +DEFGSCODE(GIMPLE_OMP_GPUKERNEL, "gimple_omp_gpukernel", GSS_OMP) + /* GIMPLE_PREDICT specifies a hint for branch prediction. PREDICT is one of the predictors from predict.def. diff --git a/gcc/gimple.h b/gcc/gimple.h index 781801b..a32d83c 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -153,6 +153,7 @@ enum gf_mask { GF_OMP_FOR_KIND_TASKLOOP = 2, GF_OMP_FOR_KIND_CILKFOR = 3, GF_OMP_FOR_KIND_OACC_LOOP = 4, + GF_OMP_FOR_KIND_KERNEL_BODY = 5, /* Flag for SIMD variants of OMP_FOR kinds. */ GF_OMP_FOR_SIMD = 1 << 3, GF_OMP_FOR_KIND_SIMD = GF_OMP_FOR_SIMD | 0, @@ -621,8 +622,24 @@ struct GTY((tag("GSS_OMP_FOR"))) /* [ WORD 11 ] Pre-body evaluated before the loop body begins. */ gimple_seq pre_body; + + /* [ WORD 12 ] + 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. */ + bool kernel_phony; }; +/* Descriptor of one dimension of a kernel grid. */ + +struct GTY(()) gimple_omp_target_grid_dim +{ + /* Size of the whole grid in the respective dimension. */ + tree grid_size; + + /* Size of the workgroup in the respective dimension. */ + tree workgroup_size; +}; /* GIMPLE_OMP_PARALLEL, GIMPLE_OMP_TARGET, GIMPLE_OMP_TASK */ @@ -642,6 +659,26 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) /* [ WORD 10 ] 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. */ + bool kernel_phony; }; /* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */ @@ -724,14 +761,14 @@ struct GTY((tag("GSS_OMP_CONTINUE"))) tree control_use; }; -/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_TEAMS, GIMPLE_OMP_ORDERED */ +/* GIMPLE_OMP_SINGLE, GIMPLE_OMP_ORDERED */ struct GTY((tag("GSS_OMP_SINGLE_LAYOUT"))) gimple_statement_omp_single_layout : public gimple_statement_omp { /* [ WORD 1-7 ] : base class */ - /* [ WORD 7 ] */ + /* [ WORD 8 ] */ tree clauses; }; @@ -742,11 +779,18 @@ struct GTY((tag("GSS_OMP_SINGLE_LAYOUT"))) stmt->code == GIMPLE_OMP_SINGLE. */ }; -struct GTY((tag("GSS_OMP_SINGLE_LAYOUT"))) +/* GIMPLE_OMP_TEAMS */ + +struct GTY((tag("GSS_OMP_TEAMS_LAYOUT"))) gomp_teams : public gimple_statement_omp_single_layout { - /* No extra fields; adds invariant: - stmt->code == GIMPLE_OMP_TEAMS. */ + /* [ WORD 1-8 ] : base class */ + + /* [ WORD 9 ] + 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. */ + bool kernel_phony; }; struct GTY((tag("GSS_OMP_SINGLE_LAYOUT"))) @@ -1450,6 +1494,7 @@ gomp_task *gimple_build_omp_task (gimple_seq, tree, tree, tree, tree, tree, tree); gimple *gimple_build_omp_section (gimple_seq); gimple *gimple_build_omp_master (gimple_seq); +gimple *gimple_build_omp_gpukernel (gimple_seq); gimple *gimple_build_omp_taskgroup (gimple_seq); gomp_continue *gimple_build_omp_continue (tree, tree); gomp_ordered *gimple_build_omp_ordered (gimple_seq, tree); @@ -1458,6 +1503,7 @@ 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); @@ -1708,6 +1754,7 @@ gimple_has_substatements (gimple *g) case GIMPLE_OMP_CRITICAL: case GIMPLE_WITH_CLEANUP_EXPR: case GIMPLE_TRANSACTION: + case GIMPLE_OMP_GPUKERNEL: return true; default: @@ -5077,6 +5124,21 @@ gimple_omp_for_set_pre_body (gimple *gs, gimple_seq pre_body) omp_for_stmt->pre_body = pre_body; } +/* Return the kernel_phony of OMP_FOR statement. */ + +static inline bool +gimple_omp_for_kernel_phony (const gomp_for *omp_for) +{ + return omp_for->kernel_phony; +} + +/* Set kernel_phony flag of OMP_FOR to VALUE. */ + +static inline void +gimple_omp_for_set_kernel_phony (gomp_for *omp_for, bool value) +{ + omp_for->kernel_phony = value; +} /* Return the clauses associated with OMP_PARALLEL GS. */ @@ -5163,6 +5225,22 @@ gimple_omp_parallel_set_data_arg (gomp_parallel *omp_parallel_stmt, omp_parallel_stmt->data_arg = data_arg; } +/* Return the kernel_phony flag of OMP_PARALLEL_STMT. */ + +static inline bool +gimple_omp_parallel_kernel_phony (const gomp_parallel *omp_parallel_stmt) +{ + return omp_parallel_stmt->kernel_phony; +} + +/* Set kernel_phony flag of OMP_PARALLEL_STMT to VALUE. */ + +static inline void +gimple_omp_parallel_set_kernel_phony (gomp_parallel *omp_parallel_stmt, + bool value) +{ + omp_parallel_stmt->kernel_phony = value; +} /* Return the clauses associated with OMP_TASK GS. */ @@ -5607,6 +5685,72 @@ 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. */ @@ -5636,6 +5780,21 @@ gimple_omp_teams_set_clauses (gomp_teams *omp_teams_stmt, tree clauses) omp_teams_stmt->clauses = clauses; } +/* Return the kernel_phony flag of an OMP_TEAMS_STMT. */ + +static inline bool +gimple_omp_teams_kernel_phony (const gomp_teams *omp_teams_stmt) +{ + return omp_teams_stmt->kernel_phony; +} + +/* Set kernel_phony flag of an OMP_TEAMS_STMT to VALUE. */ + +static inline void +gimple_omp_teams_set_kernel_phony (gomp_teams *omp_teams_stmt, bool value) +{ + omp_teams_stmt->kernel_phony = value; +} /* Return the clauses associated with OMP_SECTIONS GS. */ @@ -5965,7 +6124,8 @@ gimple_return_set_retbnd (gimple *gs, tree retval) case GIMPLE_OMP_RETURN: \ case GIMPLE_OMP_ATOMIC_LOAD: \ case GIMPLE_OMP_ATOMIC_STORE: \ - case GIMPLE_OMP_CONTINUE + case GIMPLE_OMP_CONTINUE: \ + case GIMPLE_OMP_GPUKERNEL static inline bool is_gimple_omp (const gimple *stmt) diff --git a/gcc/gsstruct.def b/gcc/gsstruct.def index d84e098..9d6b0ef 100644 --- a/gcc/gsstruct.def +++ b/gcc/gsstruct.def @@ -47,6 +47,7 @@ DEFGSSTRUCT(GSS_OMP_PARALLEL_LAYOUT, gimple_statement_omp_parallel_layout, false DEFGSSTRUCT(GSS_OMP_TASK, gomp_task, false) DEFGSSTRUCT(GSS_OMP_SECTIONS, gomp_sections, false) DEFGSSTRUCT(GSS_OMP_SINGLE_LAYOUT, gimple_statement_omp_single_layout, false) +DEFGSSTRUCT(GSS_OMP_TEAMS_LAYOUT, gomp_teams, false) DEFGSSTRUCT(GSS_OMP_CONTINUE, gomp_continue, false) DEFGSSTRUCT(GSS_OMP_ATOMIC_LOAD, gomp_atomic_load, false) DEFGSSTRUCT(GSS_OMP_ATOMIC_STORE_LAYOUT, gomp_atomic_store, false) diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index ea9cf0d..59c677b 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -302,8 +302,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_SINGLE_COPY_START, "GOMP_single_copy_start", BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST) 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_OFFLOAD_REGISTER, "GOMP_offload_register", + BT_FN_VOID_PTR_INT_PTR, ATTR_NOTHROW_LIST) +DEF_GOMP_BUILTIN (BUILT_IN_GOMP_OFFLOAD_UNREGISTER, "GOMP_offload_unregister", + BT_FN_VOID_PTR_INT_PTR, ATTR_NOTHROW_LIST) DEF_GOMP_BUILTIN (BUILT_IN_GOMP_TARGET, "GOMP_target_41", - BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR, + BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_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) diff --git a/gcc/omp-low.c b/gcc/omp-low.c index d0264e9..379535c 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -80,6 +80,8 @@ along with GCC; see the file COPYING3. If not see #include "lto-section-names.h" #include "gomp-constants.h" #include "gimple-pretty-print.h" +#include "symbol-summary.h" +#include "hsa.h" /* Lowering of OMP parallel and workshare constructs proceeds in two phases. The first phase scans the function looking for OMP statements @@ -510,6 +512,63 @@ is_combined_parallel (struct omp_region *region) return region->is_combined_parallel; } +/* Adjust *COND_CODE and *N2 so that the former is either LT_EXPR or + GT_EXPR. */ + +static void +adjust_for_condition (location_t loc, enum tree_code *cond_code, tree *n2) +{ + switch (*cond_code) + { + case LT_EXPR: + case GT_EXPR: + case NE_EXPR: + break; + case LE_EXPR: + if (POINTER_TYPE_P (TREE_TYPE (*n2))) + *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, 1); + else + *n2 = fold_build2_loc (loc, PLUS_EXPR, TREE_TYPE (*n2), *n2, + build_int_cst (TREE_TYPE (*n2), 1)); + *cond_code = LT_EXPR; + break; + case GE_EXPR: + if (POINTER_TYPE_P (TREE_TYPE (*n2))) + *n2 = fold_build_pointer_plus_hwi_loc (loc, *n2, -1); + else + *n2 = fold_build2_loc (loc, MINUS_EXPR, TREE_TYPE (*n2), *n2, + build_int_cst (TREE_TYPE (*n2), 1)); + *cond_code = GT_EXPR; + break; + default: + gcc_unreachable (); + } +} + +/* Return the looping step from INCR, extracted from the step of a gimple omp + for statement. */ + +static tree +get_omp_for_step_from_incr (location_t loc, tree incr) +{ + tree step; + switch (TREE_CODE (incr)) + { + case PLUS_EXPR: + step = TREE_OPERAND (incr, 1); + break; + case POINTER_PLUS_EXPR: + step = fold_convert (ssizetype, TREE_OPERAND (incr, 1)); + break; + case MINUS_EXPR: + step = TREE_OPERAND (incr, 1); + step = fold_build1_loc (loc, NEGATE_EXPR, TREE_TYPE (step), step); + break; + default: + gcc_unreachable (); + } + return step; +} /* Extract the header elements of parallel loop FOR_STMT and store them into *FD. */ @@ -634,58 +693,14 @@ extract_omp_for_data (gomp_for *for_stmt, struct omp_for_data *fd, loop->cond_code = gimple_omp_for_cond (for_stmt, i); loop->n2 = gimple_omp_for_final (for_stmt, i); - switch (loop->cond_code) - { - case LT_EXPR: - case GT_EXPR: - break; - case NE_EXPR: - gcc_assert (gimple_omp_for_kind (for_stmt) - == GF_OMP_FOR_KIND_CILKSIMD - || (gimple_omp_for_kind (for_stmt) - == GF_OMP_FOR_KIND_CILKFOR)); - break; - case LE_EXPR: - if (POINTER_TYPE_P (TREE_TYPE (loop->n2))) - loop->n2 = fold_build_pointer_plus_hwi_loc (loc, loop->n2, 1); - else - loop->n2 = fold_build2_loc (loc, - PLUS_EXPR, TREE_TYPE (loop->n2), loop->n2, - build_int_cst (TREE_TYPE (loop->n2), 1)); - loop->cond_code = LT_EXPR; - break; - case GE_EXPR: - if (POINTER_TYPE_P (TREE_TYPE (loop->n2))) - loop->n2 = fold_build_pointer_plus_hwi_loc (loc, loop->n2, -1); - else - loop->n2 = fold_build2_loc (loc, - MINUS_EXPR, TREE_TYPE (loop->n2), loop->n2, - build_int_cst (TREE_TYPE (loop->n2), 1)); - loop->cond_code = GT_EXPR; - break; - default: - gcc_unreachable (); - } + gcc_assert (loop->cond_code != NE_EXPR + || gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_CILKSIMD + || gimple_omp_for_kind (for_stmt) == GF_OMP_FOR_KIND_CILKFOR); + adjust_for_condition (loc, &loop->cond_code, &loop->n2); t = gimple_omp_for_incr (for_stmt, i); gcc_assert (TREE_OPERAND (t, 0) == var); - switch (TREE_CODE (t)) - { - case PLUS_EXPR: - loop->step = TREE_OPERAND (t, 1); - break; - case POINTER_PLUS_EXPR: - loop->step = fold_convert (ssizetype, TREE_OPERAND (t, 1)); - break; - case MINUS_EXPR: - loop->step = TREE_OPERAND (t, 1); - loop->step = fold_build1_loc (loc, - NEGATE_EXPR, TREE_TYPE (loop->step), - loop->step); - break; - default: - gcc_unreachable (); - } + loop->step = get_omp_for_step_from_incr (loc, t); if (simd || (fd->sched_kind == OMP_CLAUSE_SCHEDULE_STATIC @@ -1389,7 +1404,16 @@ build_outer_var_ref (tree var, omp_context *ctx, bool lastprivate = false) } } else if (ctx->outer) - x = lookup_decl (var, ctx->outer); + { + omp_context *outer = ctx->outer; + if (gimple_code (outer->stmt) == GIMPLE_OMP_GPUKERNEL) + { + outer = outer->outer; + gcc_assert (outer + && gimple_code (outer->stmt) != GIMPLE_OMP_GPUKERNEL); + } + x = lookup_decl (var, outer); + } else if (is_reference (var)) /* This can happen with orphaned constructs. If var is reference, it is possible it is shared and as such valid. */ @@ -1837,6 +1861,8 @@ fixup_child_record_type (omp_context *ctx) { tree f, type = ctx->record_type; + if (!ctx->receiver_decl) + return; /* ??? It isn't sufficient to just call remap_type here, because variably_modified_type_p doesn't work the way we expect for record types. Testing each field for whether it needs remapping @@ -2730,8 +2756,11 @@ scan_omp_parallel (gimple_stmt_iterator *gsi, omp_context *outer_ctx) DECL_NAMELESS (name) = 1; TYPE_NAME (ctx->record_type) = name; TYPE_ARTIFICIAL (ctx->record_type) = 1; - create_omp_child_function (ctx, false); - gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn); + if (!gimple_omp_parallel_kernel_phony (stmt)) + { + create_omp_child_function (ctx, false); + gimple_omp_parallel_set_child_fn (stmt, ctx->cb.dst_fn); + } scan_sharing_clauses (gimple_omp_parallel_clauses (stmt), ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); @@ -3156,6 +3185,13 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx) DECL_NAMELESS (name) = 1; 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) { if (is_gimple_omp_oacc (stmt)) @@ -3205,6 +3241,11 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { tree c; + if (ctx && gimple_code (ctx->stmt) == GIMPLE_OMP_GPUKERNEL) + /* GPUKERNEL is an artificial construct, nesting rules will be checked in + the original copy of its contents. */ + return true; + /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP builtin) inside an OpenACC CTX. */ if (!(is_gimple_omp (stmt) @@ -3831,6 +3872,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_ORDERED: case GIMPLE_OMP_CRITICAL: + case GIMPLE_OMP_GPUKERNEL: ctx = new_omp_context (stmt, ctx); scan_omp (gimple_omp_body_ptr (stmt), ctx); break; @@ -6082,6 +6124,35 @@ gimple_build_cond_empty (tree cond) return gimple_build_cond (pred_code, lhs, rhs, NULL_TREE, NULL_TREE); } +/* Return true if a parallel REGION is within a declare target function or + within a target region and is not a part of a gridified kernel. */ + +static bool +region_needs_kernel_p (struct omp_region *region) +{ + bool indirect = false; + for (region = region->outer; region; region = region->outer) + { + if (region->type == GIMPLE_OMP_PARALLEL) + indirect = true; + else if (region->type == GIMPLE_OMP_TARGET) + { + gomp_target *tgt_stmt; + tgt_stmt = as_a (last_stmt (region->entry)); + if (gimple_omp_target_dimensions (tgt_stmt)) + return indirect; + else + return true; + } + } + + if (lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (current_function_decl))) + return true; + + return false; +} + static void expand_omp_build_assign (gimple_stmt_iterator *, tree, tree, bool = false); @@ -6236,7 +6307,8 @@ expand_parallel_call (struct omp_region *region, basic_block bb, t1 = null_pointer_node; else t1 = build_fold_addr_expr (t); - t2 = build_fold_addr_expr (gimple_omp_parallel_child_fn (entry_stmt)); + tree child_fndecl = gimple_omp_parallel_child_fn (entry_stmt); + t2 = build_fold_addr_expr (child_fndecl); vec_alloc (args, 4 + vec_safe_length (ws_args)); args->quick_push (t2); @@ -6251,6 +6323,13 @@ expand_parallel_call (struct omp_region *region, basic_block bb, force_gimple_operand_gsi (&gsi, t, true, NULL_TREE, false, GSI_CONTINUE_LINKING); + + if (hsa_gen_requested_p () + && region_needs_kernel_p (region)) + { + cgraph_node *child_cnode = cgraph_node::get (child_fndecl); + hsa_register_kernel (child_cnode); + } } /* Insert a function call whose name is FUNC_NAME with the information from @@ -12092,6 +12171,98 @@ get_oacc_fn_attrib (tree fn) return lookup_attribute (OACC_FN_ATTRIB, DECL_ATTRIBUTES (fn)); } +/* Types used to pass grid and wortkgroup sizes to kernel invocation. */ + +static GTY(()) tree kernel_dim_array_type; +static GTY(()) tree kernel_lattrs_dimnum_decl; +static GTY(()) tree kernel_lattrs_grid_decl; +static GTY(()) tree kernel_lattrs_group_decl; +static GTY(()) tree kernel_launch_attributes_type; + +/* Create types used to pass kernel launch attributes to target. */ + +static void +create_kernel_launch_attr_types (void) +{ + if (kernel_launch_attributes_type) + return; + + tree dim_arr_index_type; + dim_arr_index_type = build_index_type (build_int_cst (integer_type_node, 2)); + kernel_dim_array_type = build_array_type (uint32_type_node, + dim_arr_index_type); + + kernel_launch_attributes_type = make_node (RECORD_TYPE); + kernel_lattrs_dimnum_decl = build_decl (BUILTINS_LOCATION, FIELD_DECL, + get_identifier ("ndim"), + uint32_type_node); + DECL_CHAIN (kernel_lattrs_dimnum_decl) = NULL_TREE; + + kernel_lattrs_grid_decl = build_decl (BUILTINS_LOCATION, FIELD_DECL, + get_identifier ("grid_size"), + kernel_dim_array_type); + DECL_CHAIN (kernel_lattrs_grid_decl) = kernel_lattrs_dimnum_decl; + kernel_lattrs_group_decl = build_decl (BUILTINS_LOCATION, FIELD_DECL, + get_identifier ("group_size"), + kernel_dim_array_type); + DECL_CHAIN (kernel_lattrs_group_decl) = kernel_lattrs_grid_decl; + finish_builtin_struct (kernel_launch_attributes_type, + "__gomp_kernel_launch_attributes", + kernel_lattrs_group_decl, NULL_TREE); +} + +/* Insert before the current statement in GSI a store of VALUE to INDEX of + array (of type kernel_dim_array_type) FLD_DECL of RANGE_VAR. VALUE must be + of type uint32_type_node. */ + +static void +insert_store_range_dim (gimple_stmt_iterator *gsi, tree range_var, + tree fld_decl, int index, tree value) +{ + tree ref = build4 (ARRAY_REF, uint32_type_node, + build3 (COMPONENT_REF, kernel_dim_array_type, + range_var, fld_decl, NULL_TREE), + build_int_cst (integer_type_node, index), + NULL_TREE, NULL_TREE); + gsi_insert_before (gsi, gimple_build_assign (ref, value), GSI_SAME_STMT); +} + +/* Return a tree representation of a pointer to a structure with grid and + work-group size information. Statements filling that information will be + inserted before GSI, TGT_STMT is the target statement which has the + necessary information in it. */ + +static tree +get_kernel_launch_attributes (gimple_stmt_iterator *gsi, gomp_target *tgt_stmt) +{ + create_kernel_launch_attr_types (); + tree u32_one = build_one_cst (uint32_type_node); + tree lattrs = create_tmp_var (kernel_launch_attributes_type, + "__kernel_launch_attrs"); + 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); + 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); +} + /* Expand the GIMPLE_OMP_TARGET starting at REGION. */ static void @@ -12485,6 +12657,10 @@ expand_omp_target (struct omp_region *region) else depend = build_int_cst (ptr_type_node, 0); args.quick_push (depend); + if (gimple_omp_target_dimensions (entry_stmt)) + args.quick_push (get_kernel_launch_attributes (&gsi, entry_stmt)); + else + args.quick_push (build_zero_cst (ptr_type_node)); break; case BUILT_IN_GOACC_PARALLEL: { @@ -12588,6 +12764,255 @@ expand_omp_target (struct omp_region *region) } } +/* Expand KFOR loop as a GPGPU kernel, i.e. as a body only with iteration + variable derived from the thread number. */ + +static void +expand_omp_for_kernel (struct omp_region *kfor) +{ + tree t, threadid; + tree type, itype; + gimple_stmt_iterator gsi; + tree n1, step; + struct omp_for_data fd; + + gomp_for *for_stmt = as_a (last_stmt (kfor->entry)); + gcc_checking_assert (gimple_omp_for_kind (for_stmt) + == GF_OMP_FOR_KIND_KERNEL_BODY); + basic_block body_bb = FALLTHRU_EDGE (kfor->entry)->dest; + + gcc_assert (gimple_omp_for_collapse (for_stmt) == 1); + gcc_assert (kfor->cont); + extract_omp_for_data (for_stmt, &fd, NULL); + + itype = type = TREE_TYPE (fd.loop.v); + if (POINTER_TYPE_P (type)) + itype = signed_type_for (type); + + gsi = gsi_start_bb (body_bb); + + n1 = fd.loop.n1; + step = fd.loop.step; + n1 = force_gimple_operand_gsi (&gsi, fold_convert (type, n1), + true, NULL_TREE, true, GSI_SAME_STMT); + step = force_gimple_operand_gsi (&gsi, fold_convert (itype, step), + true, NULL_TREE, true, GSI_SAME_STMT); + threadid = build_call_expr (builtin_decl_explicit + (BUILT_IN_OMP_GET_THREAD_NUM), 0); + threadid = fold_convert (itype, threadid); + threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE, + true, GSI_CONTINUE_LINKING); + + tree startvar = fd.loop.v; + t = fold_build2 (MULT_EXPR, itype, threadid, step); + if (POINTER_TYPE_P (type)) + t = fold_build_pointer_plus (n1, t); + else + t = fold_build2 (PLUS_EXPR, type, t, n1); + t = fold_convert (type, t); + t = force_gimple_operand_gsi (&gsi, t, + DECL_P (startvar) + && TREE_ADDRESSABLE (startvar), + NULL_TREE, true, GSI_CONTINUE_LINKING); + gassign *assign_stmt = gimple_build_assign (startvar, t); + gsi_insert_after (&gsi, assign_stmt, GSI_CONTINUE_LINKING); + + /* Remove the omp for statement */ + gsi = gsi_last_bb (kfor->entry); + gsi_remove (&gsi, true); + + /* Remove the GIMPLE_OMP_CONTINUE statement. */ + gsi = gsi_last_bb (kfor->cont); + gcc_assert (!gsi_end_p (gsi) + && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_CONTINUE); + gsi_remove (&gsi, true); + + /* Replace the GIMPLE_OMP_RETURN with a real return. */ + gsi = gsi_last_bb (kfor->exit); + gcc_assert (!gsi_end_p (gsi) + && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN); + gsi_remove (&gsi, true); + + /* Fixup the much simpler CFG. */ + remove_edge (find_edge (kfor->cont, body_bb)); + + if (kfor->cont != body_bb) + set_immediate_dominator (CDI_DOMINATORS, kfor->cont, body_bb); + set_immediate_dominator (CDI_DOMINATORS, kfor->exit, kfor->cont); +} + +/* Structure passed to remap_kernel_arg_accesses so that it can remap + argument_decls. */ + +struct arg_decl_map +{ + tree old_arg; + tree new_arg; +}; + +/* Invoked through walk_gimple_op, will remap all PARM_DECLs to the ones + pertaining to kernel function. */ + +static tree +remap_kernel_arg_accesses (tree *tp, int *walk_subtrees, void *data) +{ + struct walk_stmt_info *wi = (struct walk_stmt_info *) data; + struct arg_decl_map *adm = (struct arg_decl_map *) wi->info; + tree t = *tp; + + if (t == adm->old_arg) + *tp = adm->new_arg; + *walk_subtrees = !TYPE_P (t) && !DECL_P (t); + return NULL_TREE; +} + +static void expand_omp (struct omp_region *region); + +/* If TARGET region contains a kernel body for loop, remove its region from the + TARGET and expand it in GPGPU kernel fashion. */ + +static void +expand_target_kernel_body (struct omp_region *target) +{ + if (!hsa_gen_requested_p ()) + return; + + gomp_target *tgt_stmt = as_a (last_stmt (target->entry)); + struct omp_region **pp; + + for (pp = &target->inner; *pp; pp = &(*pp)->next) + if ((*pp)->type == GIMPLE_OMP_GPUKERNEL) + break; + + struct omp_region *gpukernel = *pp; + + tree orig_child_fndecl = gimple_omp_target_child_fn (tgt_stmt); + if (!gpukernel) + { + /* HSA cannot handle OACC stuff. */ + 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)); + cgraph_node *n = cgraph_node::get (orig_child_fndecl); + + hsa_register_kernel (n); + return; + } + + gcc_assert (gimple_omp_target_dimensions (tgt_stmt)); + tree inside_block = gimple_block (first_stmt (single_succ (gpukernel->entry))); + *pp = gpukernel->next; + for (pp = &gpukernel->inner; *pp; pp = &(*pp)->next) + if ((*pp)->type == GIMPLE_OMP_FOR) + break; + + struct omp_region *kfor = *pp; + gcc_assert (kfor); + gcc_assert (gimple_omp_for_kind (last_stmt ((kfor)->entry)) + == GF_OMP_FOR_KIND_KERNEL_BODY); + *pp = kfor->next; + if (kfor->inner) + expand_omp (kfor->inner); + if (gpukernel->inner) + expand_omp (gpukernel->inner); + + tree kern_fndecl = copy_node (orig_child_fndecl); + DECL_NAME (kern_fndecl) = clone_function_name (kern_fndecl, "kernel"); + SET_DECL_ASSEMBLER_NAME (kern_fndecl, DECL_NAME (kern_fndecl)); + tree tgtblock = gimple_block (tgt_stmt); + tree fniniblock = make_node (BLOCK); + BLOCK_ABSTRACT_ORIGIN (fniniblock) = tgtblock; + BLOCK_SOURCE_LOCATION (fniniblock) = BLOCK_SOURCE_LOCATION (tgtblock); + BLOCK_SOURCE_END_LOCATION (fniniblock) = BLOCK_SOURCE_END_LOCATION (tgtblock); + DECL_INITIAL (kern_fndecl) = fniniblock; + push_struct_function (kern_fndecl); + cfun->function_end_locus = gimple_location (tgt_stmt); + pop_cfun (); + + tree old_parm_decl = DECL_ARGUMENTS (kern_fndecl); + gcc_assert (!DECL_CHAIN (old_parm_decl)); + tree new_parm_decl = copy_node (DECL_ARGUMENTS (kern_fndecl)); + DECL_CONTEXT (new_parm_decl) = kern_fndecl; + DECL_ARGUMENTS (kern_fndecl) = new_parm_decl; + struct function *kern_cfun = DECL_STRUCT_FUNCTION (kern_fndecl); + kern_cfun->curr_properties = cfun->curr_properties; + + remove_edge (BRANCH_EDGE (kfor->entry)); + expand_omp_for_kernel (kfor); + + /* Remove the omp for statement */ + gimple_stmt_iterator gsi = gsi_last_bb (gpukernel->entry); + gsi_remove (&gsi, true); + /* Replace the GIMPLE_OMP_RETURN at the end of the kernel region with a real + return. */ + gsi = gsi_last_bb (gpukernel->exit); + gcc_assert (!gsi_end_p (gsi) + && gimple_code (gsi_stmt (gsi)) == GIMPLE_OMP_RETURN); + gimple *ret_stmt = gimple_build_return (NULL); + gsi_insert_after (&gsi, ret_stmt, GSI_SAME_STMT); + gsi_remove (&gsi, true); + + /* Statements in the first BB in the target construct have been produced by + target lowering and must be copied inside the GPUKERNEL, with the two + exceptions of the first OMP statement and the OMP_DATA assignment + statement. */ + gsi = gsi_start_bb (single_succ (gpukernel->entry)); + tree data_arg = gimple_omp_target_data_arg (tgt_stmt); + tree sender = data_arg ? TREE_VEC_ELT (data_arg, 0) : NULL; + for (gimple_stmt_iterator tsi = gsi_start_bb (single_succ (target->entry)); + !gsi_end_p (tsi); gsi_next (&tsi)) + { + gimple *stmt = gsi_stmt (tsi); + if (is_gimple_omp (stmt)) + break; + if (sender + && is_gimple_assign (stmt) + && TREE_CODE (gimple_assign_rhs1 (stmt)) == ADDR_EXPR + && TREE_OPERAND (gimple_assign_rhs1 (stmt), 0) == sender) + continue; + gimple *copy = gimple_copy (stmt); + gsi_insert_before (&gsi, copy, GSI_SAME_STMT); + gimple_set_block (copy, fniniblock); + } + + move_sese_region_to_fn (kern_cfun, single_succ (gpukernel->entry), + gpukernel->exit, inside_block); + + cgraph_node *kcn = cgraph_node::get_create (kern_fndecl); + kcn->mark_force_output (); + cgraph_node *orig_child = cgraph_node::get (orig_child_fndecl); + + hsa_register_kernel (kcn, orig_child); + + cgraph_node::add_new_function (kern_fndecl, true); + push_cfun (kern_cfun); + cgraph_edge::rebuild_edges (); + + /* Re-map any mention of the PARM_DECL of the original function to the + PARM_DECL of the new one. + + TODO: It would be great if lowering produced references into the GPU + kernel decl straight away and we did not have to do this. */ + struct arg_decl_map adm; + adm.old_arg = old_parm_decl; + adm.new_arg = new_parm_decl; + basic_block bb; + FOR_EACH_BB_FN (bb, kern_cfun) + { + for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi)) + { + gimple *stmt = gsi_stmt (gsi); + struct walk_stmt_info wi; + memset (&wi, 0, sizeof (wi)); + wi.info = &adm; + walk_gimple_op (stmt, remap_kernel_arg_accesses, &wi); + } + } + pop_cfun (); + + return; +} /* Expand the parallel region tree rooted at REGION. Expansion proceeds in depth-first order. Innermost regions are expanded @@ -12607,6 +13032,8 @@ expand_omp (struct omp_region *region) region. */ if (region->type == GIMPLE_OMP_PARALLEL) determine_parallel_type (region); + else if (region->type == GIMPLE_OMP_TARGET) + expand_target_kernel_body (region); if (region->type == GIMPLE_OMP_FOR && gimple_omp_for_combined_p (last_stmt (region->entry))) @@ -14402,11 +14829,13 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) ctx); } - gimple_seq_add_stmt (&body, stmt); + if (!gimple_omp_for_kernel_phony (stmt)) + gimple_seq_add_stmt (&body, stmt); gimple_seq_add_seq (&body, gimple_omp_body (stmt)); - gimple_seq_add_stmt (&body, gimple_build_omp_continue (fd.loop.v, - fd.loop.v)); + if (!gimple_omp_for_kernel_phony (stmt)) + gimple_seq_add_stmt (&body, gimple_build_omp_continue (fd.loop.v, + fd.loop.v)); /* After the loop, add exit clauses. */ lower_reduction_clauses (gimple_omp_for_clauses (stmt), &body, ctx); @@ -14418,9 +14847,12 @@ lower_omp_for (gimple_stmt_iterator *gsi_p, omp_context *ctx) body = maybe_catch_exception (body); - /* Region exit marker goes at the end of the loop body. */ - gimple_seq_add_stmt (&body, gimple_build_omp_return (fd.have_nowait)); - maybe_add_implicit_barrier_cancel (ctx, &body); + if (!gimple_omp_for_kernel_phony (stmt)) + { + /* Region exit marker goes at the end of the loop body. */ + gimple_seq_add_stmt (&body, gimple_build_omp_return (fd.have_nowait)); + maybe_add_implicit_barrier_cancel (ctx, &body); + } /* Add OpenACC joining and reduction markers just after the loop. */ if (oacc_tail) @@ -14863,6 +15294,14 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) par_olist = NULL; par_ilist = NULL; par_rlist = NULL; + bool phony_construct = is_a (stmt) + && gimple_omp_parallel_kernel_phony (as_a (stmt)); + if (phony_construct && ctx->record_type) + { + gcc_checking_assert (!ctx->receiver_decl); + ctx->receiver_decl = create_tmp_var + (build_reference_type (ctx->record_type), ".omp_rec"); + } lower_rec_input_clauses (clauses, &par_ilist, &par_olist, ctx, NULL); lower_omp (&par_body, ctx); if (gimple_code (stmt) == GIMPLE_OMP_PARALLEL) @@ -14921,13 +15360,19 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq_add_stmt (&new_body, gimple_build_omp_continue (integer_zero_node, integer_zero_node)); - gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); - gimple_omp_set_body (stmt, new_body); + if (!phony_construct) + { + gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false)); + gimple_omp_set_body (stmt, new_body); + } bind = gimple_build_bind (NULL, NULL, gimple_bind_block (par_bind)); gsi_replace (gsi_p, dep_bind ? dep_bind : bind, true); gimple_bind_add_seq (bind, ilist); - gimple_bind_add_stmt (bind, stmt); + if (!phony_construct) + gimple_bind_add_stmt (bind, stmt); + else + gimple_bind_add_seq (bind, new_body); gimple_bind_add_seq (bind, olist); pop_gimplify_context (NULL); @@ -16001,19 +16446,22 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx) &bind_body, &dlist, ctx, NULL); lower_omp (gimple_omp_body_ptr (teams_stmt), ctx); lower_reduction_clauses (gimple_omp_teams_clauses (teams_stmt), &olist, ctx); - gimple_seq_add_stmt (&bind_body, teams_stmt); - - location_t loc = gimple_location (teams_stmt); - tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS); - gimple *call = gimple_build_call (decl, 2, num_teams, thread_limit); - gimple_set_location (call, loc); - gimple_seq_add_stmt (&bind_body, call); + if (!gimple_omp_teams_kernel_phony (teams_stmt)) + { + gimple_seq_add_stmt (&bind_body, teams_stmt); + location_t loc = gimple_location (teams_stmt); + tree decl = builtin_decl_explicit (BUILT_IN_GOMP_TEAMS); + gimple *call = gimple_build_call (decl, 2, num_teams, thread_limit); + gimple_set_location (call, loc); + gimple_seq_add_stmt (&bind_body, call); + } gimple_seq_add_seq (&bind_body, gimple_omp_body (teams_stmt)); gimple_omp_set_body (teams_stmt, NULL); gimple_seq_add_seq (&bind_body, olist); gimple_seq_add_seq (&bind_body, dlist); - gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true)); + if (!gimple_omp_teams_kernel_phony (teams_stmt)) + gimple_seq_add_stmt (&bind_body, gimple_build_omp_return (true)); gimple_bind_set_body (bind, bind_body); pop_gimplify_context (bind); @@ -16024,6 +16472,17 @@ lower_omp_teams (gimple_stmt_iterator *gsi_p, omp_context *ctx) TREE_USED (block) = 1; } +/* Expand code within an artificial GPUKERNELS OMP construct. */ + +static void +lower_omp_gpukernel (gimple_stmt_iterator *gsi_p, omp_context *ctx) +{ + gimple *stmt = gsi_stmt (*gsi_p); + lower_omp (gimple_omp_body_ptr (stmt), ctx); + gimple_seq_add_stmt (gimple_omp_body_ptr (stmt), + gimple_build_omp_return (false)); +} + /* Callback for lower_omp_1. Return non-NULL if *tp needs to be regimplified. If DATA is non-NULL, lower_omp_1 is outside @@ -16235,6 +16694,11 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_assert (ctx); lower_omp_teams (gsi_p, ctx); break; + case GIMPLE_OMP_GPUKERNEL: + ctx = maybe_lookup_ctx (stmt); + gcc_assert (ctx); + lower_omp_gpukernel (gsi_p, ctx); + break; case GIMPLE_CALL: tree fndecl; call_stmt = as_a (stmt); @@ -16324,7 +16788,647 @@ lower_omp (gimple_seq *body, omp_context *ctx) fold_stmt (&gsi); input_location = saved_location; } - + +/* Returen true if STMT is an assignment of a register-type into a local + VAR_DECL. */ + +static bool +reg_assignment_to_local_var_p (gimple *stmt) +{ + gassign *assign = dyn_cast (stmt); + if (!assign) + return false; + tree lhs = gimple_assign_lhs (assign); + if (TREE_CODE (lhs) != VAR_DECL + || !is_gimple_reg_type (TREE_TYPE (lhs)) + || is_global_var (lhs)) + return false; + return true; +} + +/* Return true if all statements in SEQ are assignments to local register-type + variables. */ + +static bool +seq_only_contains_local_assignments (gimple_seq seq) +{ + if (!seq) + return true; + + gimple_stmt_iterator gsi; + for (gsi = gsi_start (seq); !gsi_end_p (gsi); gsi_next (&gsi)) + if (!reg_assignment_to_local_var_p (gsi_stmt (gsi))) + return false; + return true; +} + + +/* Scan statements in SEQ and call itself recursively on any bind. If during + whole search only assignments to register-type local variables and one + single OMP statement is encountered, return true, otherwise return false. + 8RET is where we store any OMP statement encountered. TARGET_LOC and NAME + are used for dumping a note about a failure. */ + +static bool +find_single_omp_among_assignments_1 (gimple_seq seq, location_t target_loc, + const char *name, gimple **ret) +{ + gimple_stmt_iterator gsi; + for (gsi = gsi_start (seq); !gsi_end_p (gsi); gsi_next (&gsi)) + { + gimple *stmt = gsi_stmt (gsi); + + if (reg_assignment_to_local_var_p (stmt)) + continue; + if (gbind *bind = dyn_cast (stmt)) + { + if (!find_single_omp_among_assignments_1 (gimple_bind_body (bind), + target_loc, name, ret)) + return false; + } + else if (is_gimple_omp (stmt)) + { + if (*ret) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, target_loc, + "Will not turn target construct into a simple " + "GPGPU kernel because %s construct contains " + "multiple OpenMP constructs\n", name); + return false; + } + *ret = stmt; + } + else + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, target_loc, + "Will not turn target construct into a simple " + "GPGPU kernel because %s construct contains " + "a complex statement\n", name); + return false; + } + } + return true; +} + +/* Scan statements in SEQ and make sure that it and any binds in it contain + only assignments to local register-type variables and one OMP construct. If + so, return that construct, otherwise return NULL. If dumping is enabled and + function fails, use TARGET_LOC and NAME to dump a note with the reason for + failure. */ + +static gimple * +find_single_omp_among_assignments (gimple_seq seq, location_t target_loc, + const char *name) +{ + if (!seq) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, target_loc, + "Will not turn target construct into a simple " + "GPGPU kernel because %s construct has empty " + "body\n", + name); + return NULL; + } + + gimple *ret = NULL; + if (find_single_omp_among_assignments_1 (seq, target_loc, name, &ret)) + { + if (!ret && dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, target_loc, + "Will not turn target construct into a simple " + "GPGPU kernel because %s construct does not contain" + "any other OpenMP construct\n", name); + return ret; + } + else + return NULL; +} + +/* Walker function looking for statements there is no point gridifying (and for + noreturn function calls which we cannot do). Return non-NULL if such a + function is found. */ + +static tree +find_ungridifiable_statement (gimple_stmt_iterator *gsi, bool *handled_ops_p, + struct walk_stmt_info *) +{ + *handled_ops_p = false; + gimple *stmt = gsi_stmt (*gsi); + switch (gimple_code (stmt)) + { + case GIMPLE_CALL: + if (gimple_call_noreturn_p (as_a (stmt))) + { + *handled_ops_p = true; + return error_mark_node; + } + break; + + /* We may reduce the following list if we find a way to implement the + clauses, but now there is no point trying further. */ + case GIMPLE_OMP_CRITICAL: + case GIMPLE_OMP_TASKGROUP: + case GIMPLE_OMP_TASK: + case GIMPLE_OMP_SECTION: + case GIMPLE_OMP_SECTIONS: + case GIMPLE_OMP_SECTIONS_SWITCH: + case GIMPLE_OMP_TARGET: + case GIMPLE_OMP_ORDERED: + *handled_ops_p = true; + return error_mark_node; + + default: + break; + } + return NULL; +} + + +/* If TARGET follows a pattern that can be turned into a gridified GPGPU + kernel, return true, otherwise return false. In the case of success, also + fill in GROUP_SIZE_P with the requested group size or NULL if there is + none. */ + +static bool +target_follows_gridifiable_pattern (gomp_target *target, tree *group_size_p) +{ + if (gimple_omp_target_kind (target) != GF_OMP_TARGET_KIND_REGION) + return false; + + location_t tloc = gimple_location (target); + gimple *stmt = find_single_omp_among_assignments (gimple_omp_body (target), + tloc, "target"); + if (!stmt) + return false; + gomp_teams *teams = dyn_cast (stmt); + tree group_size = NULL; + if (!teams) + { + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a simple " + "GPGPU kernel because it does not have a sole teams " + "construct in it.\n"); + return false; + } + + tree clauses = gimple_omp_teams_clauses (teams); + while (clauses) + { + switch (OMP_CLAUSE_CODE (clauses)) + { + case OMP_CLAUSE_NUM_TEAMS: + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a " + "gridified GPGPU kernel because we cannot " + "handle num_teams clause of teams " + "construct\n "); + return false; + + case OMP_CLAUSE_REDUCTION: + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a " + "gridified GPGPU kernel because a reduction " + "clause is present\n "); + return false; + + case OMP_CLAUSE_THREAD_LIMIT: + group_size = OMP_CLAUSE_OPERAND (clauses, 0); + break; + + default: + break; + } + clauses = OMP_CLAUSE_CHAIN (clauses); + } + + stmt = find_single_omp_among_assignments (gimple_omp_body (teams), tloc, + "teams"); + if (!stmt) + return false; + gomp_for *dist = dyn_cast (stmt); + if (!dist) + { + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a simple " + "GPGPU kernel because the teams construct does not have " + "a sole distribute construct in it.\n"); + return false; + } + + gcc_assert (gimple_omp_for_kind (dist) == GF_OMP_FOR_KIND_DISTRIBUTE); + if (!gimple_omp_for_combined_p (dist)) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified GPGPU " + "kernel because we cannot handle a standalone " + "distribute construct\n "); + return false; + } + if (dist->collapse > 1) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified GPGPU " + "kernel because the distribute construct contains " + "collapse clause\n"); + return false; + } + struct omp_for_data fd; + extract_omp_for_data (dist, &fd, NULL); + if (fd.chunk_size) + { + if (group_size && !operand_equal_p (group_size, fd.chunk_size, 0)) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a " + "gridified GPGPU kernel because the teams " + "thread limit is different from distribute " + "schedule chunk\n"); + return false; + } + group_size = fd.chunk_size; + } + stmt = find_single_omp_among_assignments (gimple_omp_body (dist), tloc, + "distribute"); + gomp_parallel *par; + if (!stmt || !(par = dyn_cast (stmt))) + return false; + + clauses = gimple_omp_parallel_clauses (par); + while (clauses) + { + switch (OMP_CLAUSE_CODE (clauses)) + { + case OMP_CLAUSE_NUM_THREADS: + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified" + "GPGPU kernel because there is a num_threads " + "clause of the parallel construct\n"); + return false; + case OMP_CLAUSE_REDUCTION: + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a " + "gridified GPGPU kernel because a reduction " + "clause is present\n "); + return false; + default: + break; + } + clauses = OMP_CLAUSE_CHAIN (clauses); + } + + stmt = find_single_omp_among_assignments (gimple_omp_body (par), tloc, + "parallel"); + gomp_for *gfor; + if (!stmt || !(gfor = dyn_cast (stmt))) + return false; + + if (gimple_omp_for_kind (gfor) != GF_OMP_FOR_KIND_FOR) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified GPGPU " + "kernel because the inner loop is not a simple for " + "loop\n"); + return false; + } + if (gfor->collapse > 1) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified GPGPU " + "kernel because the inner loop contains collapse " + "clause\n"); + return false; + } + + if (!seq_only_contains_local_assignments (gimple_omp_for_pre_body (gfor))) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified GPGPU " + "kernel because the inner loop pre_body contains" + "a complex instruction\n"); + return false; + } + + clauses = gimple_omp_for_clauses (gfor); + while (clauses) + { + switch (OMP_CLAUSE_CODE (clauses)) + { + case OMP_CLAUSE_SCHEDULE: + if (OMP_CLAUSE_SCHEDULE_KIND (clauses) != OMP_CLAUSE_SCHEDULE_AUTO) + { + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a " + "gridified GPGPU kernel because the inner " + "loop has a non-automatic scheduling clause\n"); + return false; + } + break; + + case OMP_CLAUSE_REDUCTION: + if (dump_enabled_p ()) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a " + "gridified GPGPU kernel because a reduction " + "clause is present\n "); + return false; + + default: + break; + } + clauses = OMP_CLAUSE_CHAIN (clauses); + } + + struct walk_stmt_info wi; + memset (&wi, 0, sizeof (wi)); + if (gimple *bad = walk_gimple_seq (gimple_omp_body (gfor), + find_ungridifiable_statement, + NULL, &wi)) + { + if (dump_enabled_p ()) + { + if (is_gimple_call (bad)) + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified " + " GPGPU kernel because the inner loop contains " + "call to a noreturn function\n"); + else + dump_printf_loc (MSG_NOTE, tloc, + "Will not turn target construct into a gridified " + "GPGPU kernel because the inner loop contains " + "statement %s which cannot be transformed\n", + gimple_code_name[(int) gimple_code (bad)]); + } + return false; + } + + *group_size_p = group_size; + return true; +} + +/* Operand walker, used to remap pre-body declarations according to a hash map + provided in DATA. */ + +static tree +remap_prebody_decls (tree *tp, int *walk_subtrees, void *data) +{ + tree t = *tp; + + if (DECL_P (t) || TYPE_P (t)) + *walk_subtrees = 0; + else + *walk_subtrees = 1; + + if (TREE_CODE (t) == VAR_DECL) + { + struct walk_stmt_info *wi = (struct walk_stmt_info *) data; + hash_map *declmap = (hash_map *) wi->info; + tree *repl = declmap->get (t); + if (repl) + *tp = *repl; + } + return NULL_TREE; +} + +/* Copy leading register-type assignments to local variables in SRC to just + before DST, Creating temporaries, adjusting mapping of operands in WI and + remapping operands as necessary. Add any new temporaries to TGT_BIND. + Return the first statement that does not conform to + reg_assignment_to_local_var_p or NULL. */ + +static gimple * +copy_leading_local_assignments (gimple_seq src, gimple_stmt_iterator *dst, + gbind *tgt_bind, struct walk_stmt_info *wi) +{ + hash_map *declmap = (hash_map *) wi->info; + gimple_stmt_iterator gsi; + for (gsi = gsi_start (src); !gsi_end_p (gsi); gsi_next (&gsi)) + { + gimple *stmt = gsi_stmt (gsi); + if (gbind *bind = dyn_cast (stmt)) + { + gimple *r = copy_leading_local_assignments (gimple_bind_body (bind), + dst, tgt_bind, wi); + if (r) + return r; + else + continue; + } + if (!reg_assignment_to_local_var_p (stmt)) + return stmt; + tree lhs = gimple_assign_lhs (as_a (stmt)); + tree repl = copy_var_decl (lhs, create_tmp_var_name (NULL), + TREE_TYPE (lhs)); + DECL_CONTEXT (repl) = current_function_decl; + gimple_bind_append_vars (tgt_bind, repl); + + declmap->put (lhs, repl); + gassign *copy = as_a (gimple_copy (stmt)); + walk_gimple_op (copy, remap_prebody_decls, wi); + gsi_insert_before (dst, copy, GSI_SAME_STMT); + } + return NULL; +} + +/* Given freshly copied top level kernel SEQ, identify the individual OMP + components, mark them as part of kernel and return the inner loop, and copy + assignment leading to them just before DST, remapping them using WI and + adding new temporaries to TGT_BIND. */ + +static gomp_for * +process_kernel_body_copy (gimple_seq seq, gimple_stmt_iterator *dst, + gbind *tgt_bind, struct walk_stmt_info *wi) +{ + gimple *stmt = copy_leading_local_assignments (seq, dst, tgt_bind, wi); + gomp_teams *teams = dyn_cast (stmt); + gcc_assert (teams); + gimple_omp_teams_set_kernel_phony (teams, true); + stmt = copy_leading_local_assignments (gimple_omp_body (teams), dst, + tgt_bind, wi); + gcc_checking_assert (stmt); + gomp_for *dist = dyn_cast (stmt); + gcc_assert (dist); + gimple_seq prebody = gimple_omp_for_pre_body (dist); + if (prebody) + copy_leading_local_assignments (prebody, dst, tgt_bind, wi); + gimple_omp_for_set_kernel_phony (dist, true); + stmt = copy_leading_local_assignments (gimple_omp_body (dist), dst, + tgt_bind, wi); + gcc_checking_assert (stmt); + + gomp_parallel *parallel = as_a (stmt); + gimple_omp_parallel_set_kernel_phony (parallel, true); + stmt = copy_leading_local_assignments (gimple_omp_body (parallel), dst, + tgt_bind, wi); + gomp_for *inner_loop = as_a (stmt); + gimple_omp_for_set_kind (inner_loop, GF_OMP_FOR_KIND_KERNEL_BODY); + prebody = gimple_omp_for_pre_body (inner_loop); + if (prebody) + copy_leading_local_assignments (prebody, dst, tgt_bind, wi); + + return inner_loop; +} + +/* If TARGET points to a GOMP_TARGET which follows a gridifiable pattern, + create a GPU kernel for it. GSI must point to the same statement, TGT_BIND + is the bind into which temporaries inserted before TARGET should be + added. */ + +static void +attempt_target_gridification (gomp_target *target, gimple_stmt_iterator *gsi, + gbind *tgt_bind) +{ + tree group_size; + if (!target || !target_follows_gridifiable_pattern (target, &group_size)) + return; + + location_t loc = gimple_location (target); + if (dump_enabled_p ()) + dump_printf_loc (MSG_OPTIMIZED_LOCATIONS, loc, + "Target construct will be turned into a gridified GPGPU " + "kernel\n"); + + /* Copy target body to a GPUKERNEL construct: */ + gimple_seq kernel_seq = copy_gimple_seq_and_replace_locals + (gimple_omp_body (target)); + + hash_map *declmap = new hash_map; + struct walk_stmt_info wi; + memset (&wi, 0, sizeof (struct walk_stmt_info)); + wi.info = declmap; + + /* Copy assignments in between OMP statements before target, mark OMP + statements within copy appropriatly. */ + gomp_for *inner_loop = process_kernel_body_copy (kernel_seq, gsi, tgt_bind, + &wi); + + gbind *old_bind = as_a (gimple_seq_first (gimple_omp_body (target))); + gbind *new_bind = as_a (gimple_seq_first (kernel_seq)); + tree new_block = gimple_bind_block (new_bind); + tree enc_block = BLOCK_SUPERCONTEXT (gimple_bind_block (old_bind)); + BLOCK_CHAIN (new_block) = BLOCK_SUBBLOCKS (enc_block); + BLOCK_SUBBLOCKS (enc_block) = new_block; + BLOCK_SUPERCONTEXT (new_block) = enc_block; + gimple *gpukernel = gimple_build_omp_gpukernel (kernel_seq); + gimple_seq_add_stmt + (gimple_bind_body_ptr (as_a (gimple_omp_body (target))), + gpukernel); + + 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]; + walk_tree (&iter.initial, remap_prebody_decls, &wi, NULL); + walk_tree (&iter.final, remap_prebody_decls, &wi, NULL); + + tree itype, type = TREE_TYPE (iter.index); + if (POINTER_TYPE_P (type)) + itype = signed_type_for (type); + else + itype = type; + + enum tree_code cond_code = iter.cond; + tree n1 = iter.initial; + tree n2 = iter.final; + adjust_for_condition (loc, &cond_code, &n2); + tree step = get_omp_for_step_from_incr (loc, iter.incr); + n1 = force_gimple_operand_gsi (gsi, fold_convert (type, n1), true, + NULL_TREE, true, GSI_SAME_STMT); + n2 = force_gimple_operand_gsi (gsi, fold_convert (itype, n2), true, + NULL_TREE, + true, GSI_SAME_STMT); + tree t = build_int_cst (itype, (cond_code == LT_EXPR ? -1 : 1)); + t = fold_build2 (PLUS_EXPR, itype, step, t); + t = fold_build2 (PLUS_EXPR, itype, t, n2); + t = fold_build2 (MINUS_EXPR, itype, t, fold_convert (itype, n1)); + if (TYPE_UNSIGNED (itype) && cond_code == GT_EXPR) + t = fold_build2 (TRUNC_DIV_EXPR, itype, + fold_build1 (NEGATE_EXPR, itype, t), + fold_build1 (NEGATE_EXPR, itype, step)); + else + t = fold_build2 (TRUNC_DIV_EXPR, itype, t, step); + 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) + { + ws = fold_convert (uint32_type_node, group_size); + ws = force_gimple_operand_gsi (gsi, ws, true, NULL_TREE, true, + GSI_SAME_STMT); + } + else + ws = build_zero_cst (uint32_type_node); + gimple_omp_target_set_workgroup_size (target, i, ws); + } + + delete declmap; + return; +} + +/* Walker function doing all the work for create_target_kernels. */ + +static tree +create_target_gpukernel_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, + struct walk_stmt_info *incoming) +{ + *handled_ops_p = false; + + gimple *stmt = gsi_stmt (*gsi); + gomp_target *target = dyn_cast (stmt); + if (target) + { + gbind *tgt_bind = (gbind *) incoming->info; + gcc_checking_assert (tgt_bind); + attempt_target_gridification (target, gsi, tgt_bind); + return NULL_TREE; + } + gbind *bind = dyn_cast (stmt); + if (bind) + { + *handled_ops_p = true; + struct walk_stmt_info wi; + memset (&wi, 0, sizeof (wi)); + wi.info = bind; + walk_gimple_seq_mod (gimple_bind_body_ptr (bind), + create_target_gpukernel_stmt, NULL, &wi); + } + return NULL_TREE; +} + +/* Prepare all target constructs in BODY_P for GPU kernel generation, if they + follow a gridifiable pattern. All such targets will have their bodies + duplicated, with the new copy being put into a gpukernel. All + kernel-related construct within the gpukernel will be marked with phony + flags or kernel kinds. Moreover, some re-structuring is often needed, such + as copying pre-bodies before the target construct so that kernel grid sizes + can be computed. */ + +static void +create_target_gpukernels (gimple_seq *body_p) +{ + struct walk_stmt_info wi; + memset (&wi, 0, sizeof (wi)); + walk_gimple_seq_mod (body_p, create_target_gpukernel_stmt, NULL, &wi); +} + + /* Main entry point. */ static unsigned int @@ -16344,6 +17448,10 @@ execute_lower_omp (void) delete_omp_context); body = gimple_body (current_function_decl); + + if (hsa_gen_requested_p () && !flag_disable_hsa_gridification) + create_target_gpukernels (&body); + scan_omp (&body, NULL); gcc_assert (taskreg_nesting_level == 0); FOR_EACH_VEC_ELT (taskreg_contexts, i, ctx) @@ -16681,6 +17789,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region, case GIMPLE_OMP_TASKGROUP: case GIMPLE_OMP_CRITICAL: case GIMPLE_OMP_SECTION: + case GIMPLE_OMP_GPUKERNEL: cur_region = new_omp_region (bb, code, cur_region); fallthru = true; break;