b2b4005150
2016-01-19 Martin Jambor <mjambor@suse.cz> Martin Liska <mliska@suse.cz> Michael Matz <matz@suse.de> libgomp/ * plugin/Makefrag.am: Add HSA plugin requirements. * plugin/configfrag.ac (HSA_RUNTIME_INCLUDE): New variable. (HSA_RUNTIME_LIB): Likewise. (HSA_RUNTIME_CPPFLAGS): Likewise. (HSA_RUNTIME_INCLUDE): New substitution. (HSA_RUNTIME_LIB): Likewise. (HSA_RUNTIME_LDFLAGS): Likewise. (hsa-runtime): New configure option. (hsa-runtime-include): Likewise. (hsa-runtime-lib): Likewise. (PLUGIN_HSA): New substitution variable. Fill HSA_RUNTIME_INCLUDE and HSA_RUNTIME_LIB according to the new configure options. (PLUGIN_HSA_CPPFLAGS): Likewise. (PLUGIN_HSA_LDFLAGS): Likewise. (PLUGIN_HSA_LIBS): Likewise. Check that we have access to HSA run-time. * libgomp-plugin.h (offload_target_type): New element OFFLOAD_TARGET_TYPE_HSA. * libgomp.h (gomp_target_task): New fields firstprivate_copies and args. (bool gomp_create_target_task): Updated. (gomp_device_descr): Extra parameter of run_func and async_run_func, new field can_run_func. * libgomp_g.h (GOMP_target_ext): Update prototype. * oacc-host.c (host_run): Added a new parameter args. * target.c (calculate_firstprivate_requirements): New function. (copy_firstprivate_data): Likewise. (gomp_target_fallback_firstprivate): Use them. (gomp_target_unshare_firstprivate): New function. (gomp_get_target_fn_addr): Allow returning NULL for shared memory devices. (GOMP_target): Do host fallback for all shared memory devices. Do not pass any args to plugins. (GOMP_target_ext): Introduce device-specific argument parameter args. Allow host fallback if device shares memory. Do not remap data if device has shared memory. (gomp_target_task_fn): Likewise. Also treat shared memory devices like host fallback for mappings. (GOMP_target_data): Treat shared memory devices like host fallback. (GOMP_target_data_ext): Likewise. (GOMP_target_update): Likewise. (GOMP_target_update_ext): Likewise. Also pass NULL as args to gomp_create_target_task. (GOMP_target_enter_exit_data): Likewise. (omp_target_alloc): Treat shared memory devices like host fallback. (omp_target_free): Likewise. (omp_target_is_present): Likewise. (omp_target_memcpy): Likewise. (omp_target_memcpy_rect): Likewise. (omp_target_associate_ptr): Likewise. (gomp_load_plugin_for_device): Also load can_run. * task.c (GOMP_PLUGIN_target_task_completion): Free firstprivate_copies. (gomp_create_target_task): Accept new argument args and store it to ttask. * plugin/plugin-hsa.c: New file. gcc/ * Makefile.in (OBJS): Add new source files. (GTFILES): Add hsa.c. * common.opt (disable_hsa): New variable. (-Whsa): New warning. * config.in (ENABLE_HSA): New. * configure.ac: Treat hsa differently from other accelerators. (OFFLOAD_TARGETS): Define ENABLE_OFFLOADING according to $enable_offloading. (ENABLE_HSA): Define ENABLE_HSA according to $enable_hsa. * doc/install.texi (Configuration): Document --with-hsa-runtime, --with-hsa-runtime-include, --with-hsa-runtime-lib and --with-hsa-kmt-lib. * doc/invoke.texi (-Whsa): Document. (hsa-gen-debug-stores): Likewise. * lto-wrapper.c (compile_images_for_offload_targets): Do not attempt to invoke offload compiler for hsa acclerator. * opts.c (common_handle_option): Determine whether HSA offloading should be performed. * params.def (PARAM_HSA_GEN_DEBUG_STORES): New parameter. * builtin-types.def (BT_FN_VOID_UINT_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. * gimple-low.c (lower_stmt): Also handle GIMPLE_OMP_GRID_BODY. * gimple-pretty-print.c (dump_gimple_omp_for): Also handle GF_OMP_FOR_KIND_GRID_LOOP. (dump_gimple_omp_block): Also handle GIMPLE_OMP_GRID_BODY. (pp_gimple_stmt_1): Likewise. * gimple-walk.c (walk_gimple_stmt): Likewise. * gimple.c (gimple_build_omp_grid_body): New function. (gimple_copy): Also handle GIMPLE_OMP_GRID_BODY. * gimple.def (GIMPLE_OMP_GRID_BODY): New. * gimple.h (enum gf_mask): Added GF_OMP_PARALLEL_GRID_PHONY, GF_OMP_FOR_KIND_GRID_LOOP, GF_OMP_FOR_GRID_PHONY and GF_OMP_TEAMS_GRID_PHONY. (gimple_statement_omp_single_layout): Updated comments. (gimple_build_omp_grid_body): New function. (gimple_has_substatements): Also handle GIMPLE_OMP_GRID_BODY. (gimple_omp_for_grid_phony): New function. (gimple_omp_for_set_grid_phony): Likewise. (gimple_omp_parallel_grid_phony): Likewise. (gimple_omp_parallel_set_grid_phony): Likewise. (gimple_omp_teams_grid_phony): Likewise. (gimple_omp_teams_set_grid_phony): Likewise. (gimple_return_set_retbnd): Also handle GIMPLE_OMP_GRID_BODY. * omp-builtins.def (BUILT_IN_GOMP_OFFLOAD_REGISTER): New. (BUILT_IN_GOMP_OFFLOAD_UNREGISTER): Likewise. (BUILT_IN_GOMP_TARGET): Updated type. * omp-low.c: Include symbol-summary.h, hsa.h and params.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_GRID_BODY. (fixup_child_record_type): Bail out if receiver_decl is NULL. (scan_sharing_clauses): Handle OMP_CLAUSE__GRIDDIM_. (scan_omp_parallel): Do not create child functions for phony constructs. (check_omp_nesting_restrictions): Handle GIMPLE_OMP_GRID_BODY. (scan_omp_1_op): Checking assert we are not remapping to ERROR_MARK. Also also handle GIMPLE_OMP_GRID_BODY. (parallel_needs_hsa_kernel_p): New function. (expand_parallel_call): Register apprpriate parallel child functions as HSA kernels. (grid_launch_attributes_trees): New type. (grid_attr_trees): New variable. (grid_create_kernel_launch_attr_types): New function. (grid_insert_store_range_dim): Likewise. (grid_get_kernel_launch_attributes): Likewise. (get_target_argument_identifier_1): Likewise. (get_target_argument_identifier): Likewise. (get_target_argument_value): Likewise. (push_target_argument_according_to_value): Likewise. (get_target_arguments): Likewise. (expand_omp_target): Call get_target_arguments instead of looking up for teams and thread limit. (grid_expand_omp_for_loop): New function. (grid_arg_decl_map): New type. (grid_remap_kernel_arg_accesses): New function. (grid_expand_target_kernel_body): New function. (expand_omp): Call it. (lower_omp_for): Do not emit phony constructs. (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_grid_body): New function. (lower_omp_1): Call it. (grid_reg_assignment_to_local_var_p): New function. (grid_seq_only_contains_local_assignments): Likewise. (grid_find_single_omp_among_assignments_1): Likewise. (grid_find_single_omp_among_assignments): Likewise. (grid_find_ungridifiable_statement): Likewise. (grid_target_follows_gridifiable_pattern): Likewise. (grid_remap_prebody_decls): Likewise. (grid_copy_leading_local_assignments): Likewise. (grid_process_kernel_body_copy): Likewise. (grid_attempt_target_gridification): Likewise. (grid_gridify_all_targets_stmt): Likewise. (grid_gridify_all_targets): Likewise. (execute_lower_omp): Call grid_gridify_all_targets. (make_gimple_omp_edges): Handle GIMPLE_OMP_GRID_BODY. * tree-core.h (omp_clause_code): Added OMP_CLAUSE__GRIDDIM_. (tree_omp_clause): Added union field dimension. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE__GRIDDIM_. * tree.c (omp_clause_num_ops): Added number of arguments of OMP_CLAUSE__GRIDDIM_. (omp_clause_code_name): Added 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. * passes.def: Schedule pass_ipa_hsa and pass_gen_hsail. * tree-pass.h (make_pass_gen_hsail): Declare. (make_pass_ipa_hsa): Likewise. * ipa-hsa.c: New file. * lto-section-in.c (lto_section_name): Add hsa section name. * lto-streamer.h (lto_section_type): Add hsa section. * timevar.def (TV_IPA_HSA): New. * hsa-brig-format.h: New file. * hsa-brig.c: New file. * hsa-dump.c: Likewise. * hsa-gen.c: Likewise. * hsa.c: Likewise. * hsa.h: Likewise. * toplev.c (compile_file): Call hsa_output_brig. * hsa-regalloc.c: New file. gcc/fortran/ * types.def (BT_FN_VOID_UINT_PTR_INT_PTR): New. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_INT_INT): Removed. (BT_FN_VOID_INT_OMPFN_SIZE_PTR_PTR_PTR_UINT_PTR_PTR): New. gcc/lto/ * lto-partition.c: Include "hsa.h" (add_symbol_to_partition_1): Put hsa implementations into the same partition as host implementations. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_async_run): New unused parameter. (GOMP_OFFLOAD_run): Likewise. include/ * gomp-constants.h (GOMP_DEVICE_HSA): New macro. (GOMP_VERSION_HSA): Likewise. (GOMP_TARGET_ARG_DEVICE_MASK): Likewise. (GOMP_TARGET_ARG_DEVICE_ALL): Likewise. (GOMP_TARGET_ARG_SUBSEQUENT_PARAM): Likewise. (GOMP_TARGET_ARG_ID_MASK): Likewise. (GOMP_TARGET_ARG_NUM_TEAMS): Likewise. (GOMP_TARGET_ARG_THREAD_LIMIT): Likewise. (GOMP_TARGET_ARG_VALUE_SHIFT): Likewise. (GOMP_TARGET_ARG_HSA_KERNEL_ATTRIBUTES): Likewise. From-SVN: r232549
396 lines
16 KiB
Modula-2
396 lines
16 KiB
Modula-2
/* This file contains the definitions of the GIMPLE IR tuples used in GCC.
|
|
|
|
Copyright (C) 2007-2016 Free Software Foundation, Inc.
|
|
Contributed by Aldy Hernandez <aldyh@redhat.com>
|
|
|
|
This file is part of GCC.
|
|
|
|
GCC is free software; you can redistribute it and/or modify it under
|
|
the terms of the GNU General Public License as published by the Free
|
|
Software Foundation; either version 3, or (at your option) any later
|
|
version.
|
|
|
|
GCC is distributed in the hope that it will be useful, but WITHOUT ANY
|
|
WARRANTY; without even the implied warranty of MERCHANTABILITY or
|
|
FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
|
|
for more details.
|
|
|
|
You should have received a copy of the GNU General Public License
|
|
along with GCC; see the file COPYING3. If not see
|
|
<http://www.gnu.org/licenses/>. */
|
|
|
|
/* The format of this file is
|
|
DEFGSCODE(GIMPLE_symbol, printable name, GSS_symbol). */
|
|
|
|
|
|
/* Error marker. This is used in similar ways as ERROR_MARK in tree.def. */
|
|
DEFGSCODE(GIMPLE_ERROR_MARK, "gimple_error_mark", GSS_BASE)
|
|
|
|
/* IMPORTANT. Do not rearrange the codes between GIMPLE_COND and
|
|
GIMPLE_RETURN. The ordering is exposed by gimple_has_ops calls.
|
|
These are all the GIMPLE statements with register operands. */
|
|
|
|
/* GIMPLE_COND <COND_CODE, OP1, OP2, TRUE_LABEL, FALSE_LABEL>
|
|
represents the conditional jump:
|
|
|
|
if (OP1 COND_CODE OP2) goto TRUE_LABEL else goto FALSE_LABEL
|
|
|
|
COND_CODE is the tree code used as the comparison predicate. It
|
|
must be of class tcc_comparison.
|
|
|
|
OP1 and OP2 are the operands used in the comparison. They must be
|
|
accepted by is_gimple_operand.
|
|
|
|
TRUE_LABEL and FALSE_LABEL are the LABEL_DECL nodes used as the
|
|
jump target for the comparison. */
|
|
DEFGSCODE(GIMPLE_COND, "gimple_cond", GSS_WITH_OPS)
|
|
|
|
/* GIMPLE_DEBUG represents a debug statement. */
|
|
DEFGSCODE(GIMPLE_DEBUG, "gimple_debug", GSS_WITH_OPS)
|
|
|
|
/* GIMPLE_GOTO <TARGET> represents unconditional jumps.
|
|
TARGET is a LABEL_DECL or an expression node for computed GOTOs. */
|
|
DEFGSCODE(GIMPLE_GOTO, "gimple_goto", GSS_WITH_OPS)
|
|
|
|
/* GIMPLE_LABEL <LABEL> represents label statements. LABEL is a
|
|
LABEL_DECL representing a jump target. */
|
|
DEFGSCODE(GIMPLE_LABEL, "gimple_label", GSS_WITH_OPS)
|
|
|
|
/* GIMPLE_SWITCH <INDEX, DEFAULT_LAB, LAB1, ..., LABN> represents the
|
|
multiway branch:
|
|
|
|
switch (INDEX)
|
|
{
|
|
case LAB1: ...; break;
|
|
...
|
|
case LABN: ...; break;
|
|
default: ...
|
|
}
|
|
|
|
INDEX is the variable evaluated to decide which label to jump to.
|
|
|
|
DEFAULT_LAB, LAB1 ... LABN are the tree nodes representing case labels.
|
|
They must be CASE_LABEL_EXPR nodes. */
|
|
DEFGSCODE(GIMPLE_SWITCH, "gimple_switch", GSS_WITH_OPS)
|
|
|
|
/* IMPORTANT.
|
|
|
|
Do not rearrange the codes between GIMPLE_ASSIGN and GIMPLE_RETURN.
|
|
It's exposed by GIMPLE_RANGE_CHECK calls. These are all the GIMPLE
|
|
statements with memory and register operands. */
|
|
|
|
/* GIMPLE_ASSIGN <SUBCODE, LHS, RHS1[, RHS2]> represents the assignment
|
|
statement
|
|
|
|
LHS = RHS1 SUBCODE RHS2.
|
|
|
|
SUBCODE is the tree code for the expression computed by the RHS of the
|
|
assignment. It must be one of the tree codes accepted by
|
|
get_gimple_rhs_class. If LHS is not a gimple register according to
|
|
is_gimple_reg, SUBCODE must be of class GIMPLE_SINGLE_RHS.
|
|
|
|
LHS is the operand on the LHS of the assignment. It must be a tree node
|
|
accepted by is_gimple_lvalue.
|
|
|
|
RHS1 is the first operand on the RHS of the assignment. It must always be
|
|
present. It must be a tree node accepted by is_gimple_val.
|
|
|
|
RHS2 is the second operand on the RHS of the assignment. It must be a tree
|
|
node accepted by is_gimple_val. This argument exists only if SUBCODE is
|
|
of class GIMPLE_BINARY_RHS. */
|
|
DEFGSCODE(GIMPLE_ASSIGN, "gimple_assign", GSS_WITH_MEM_OPS)
|
|
|
|
/* GIMPLE_ASM <STRING, I1, ..., IN, O1, ... OM, C1, ..., CP>
|
|
represents inline assembly statements.
|
|
|
|
STRING is the string containing the assembly statements.
|
|
I1 ... IN are the N input operands.
|
|
O1 ... OM are the M output operands.
|
|
C1 ... CP are the P clobber operands.
|
|
L1 ... LQ are the Q label operands. */
|
|
DEFGSCODE(GIMPLE_ASM, "gimple_asm", GSS_ASM)
|
|
|
|
/* GIMPLE_CALL <FN, LHS, ARG1, ..., ARGN[, CHAIN]> represents function
|
|
calls.
|
|
|
|
FN is the callee. It must be accepted by is_gimple_call_addr.
|
|
|
|
LHS is the operand where the return value from FN is stored. It may
|
|
be NULL.
|
|
|
|
ARG1 ... ARGN are the arguments. They must all be accepted by
|
|
is_gimple_operand.
|
|
|
|
CHAIN is the optional static chain link for nested functions. */
|
|
DEFGSCODE(GIMPLE_CALL, "gimple_call", GSS_CALL)
|
|
|
|
/* GIMPLE_TRANSACTION <BODY, LABEL> represents __transaction_atomic and
|
|
__transaction_relaxed blocks.
|
|
BODY is the sequence of statements inside the transaction.
|
|
LABEL is a label for the statement immediately following the
|
|
transaction. This is before RETURN so that it has MEM_OPS,
|
|
so that it can clobber global memory. */
|
|
DEFGSCODE(GIMPLE_TRANSACTION, "gimple_transaction", GSS_TRANSACTION)
|
|
|
|
/* GIMPLE_RETURN <RETVAL> represents return statements.
|
|
|
|
RETVAL is the value to return or NULL. If a value is returned it
|
|
must be accepted by is_gimple_operand. */
|
|
DEFGSCODE(GIMPLE_RETURN, "gimple_return", GSS_WITH_MEM_OPS)
|
|
|
|
/* GIMPLE_BIND <VARS, BLOCK, BODY> represents a lexical scope.
|
|
VARS is the set of variables declared in that scope.
|
|
BLOCK is the symbol binding block used for debug information.
|
|
BODY is the sequence of statements in the scope. */
|
|
DEFGSCODE(GIMPLE_BIND, "gimple_bind", GSS_BIND)
|
|
|
|
/* GIMPLE_CATCH <TYPES, HANDLER> represents a typed exception handler.
|
|
TYPES is the type (or list of types) handled. HANDLER is the
|
|
sequence of statements that handle these types. */
|
|
DEFGSCODE(GIMPLE_CATCH, "gimple_catch", GSS_CATCH)
|
|
|
|
/* GIMPLE_EH_FILTER <TYPES, FAILURE> represents an exception
|
|
specification. TYPES is a list of allowed types and FAILURE is the
|
|
sequence of statements to execute on failure. */
|
|
DEFGSCODE(GIMPLE_EH_FILTER, "gimple_eh_filter", GSS_EH_FILTER)
|
|
|
|
/* GIMPLE_EH_MUST_NOT_THROW <DECL> represents an exception barrier.
|
|
DECL is a noreturn function decl taking no arguments that will
|
|
be invoked if an exception propagates to this point. */
|
|
DEFGSCODE(GIMPLE_EH_MUST_NOT_THROW, "gimple_eh_must_not_throw", GSS_EH_MNT)
|
|
|
|
/* GIMPLE_EH_ELSE <N_BODY, E_BODY> must be the sole contents of
|
|
a GIMPLE_TRY_FINALLY node. For all normal exits from the try block,
|
|
N_BODY is run; for all exception exits from the try block,
|
|
E_BODY is run. */
|
|
DEFGSCODE(GIMPLE_EH_ELSE, "gimple_eh_else", GSS_EH_ELSE)
|
|
|
|
/* GIMPLE_RESX resumes execution after an exception. */
|
|
DEFGSCODE(GIMPLE_RESX, "gimple_resx", GSS_EH_CTRL)
|
|
|
|
/* GIMPLE_EH_DISPATCH demultiplexes an exception edge based on
|
|
the FILTER argument. */
|
|
DEFGSCODE(GIMPLE_EH_DISPATCH, "gimple_eh_dispatch", GSS_EH_CTRL)
|
|
|
|
/* GIMPLE_PHI <RESULT, ARG1, ..., ARGN> represents the PHI node
|
|
|
|
RESULT = PHI <ARG1, ..., ARGN>
|
|
|
|
RESULT is the SSA name created by this PHI node.
|
|
|
|
ARG1 ... ARGN are the arguments to the PHI node. N must be
|
|
exactly the same as the number of incoming edges to the basic block
|
|
holding the PHI node. Every argument is either an SSA name or a
|
|
tree node of class tcc_constant. */
|
|
DEFGSCODE(GIMPLE_PHI, "gimple_phi", GSS_PHI)
|
|
|
|
/* GIMPLE_TRY <TRY_KIND, EVAL, CLEANUP>
|
|
represents a try/catch or a try/finally statement.
|
|
|
|
TRY_KIND is either GIMPLE_TRY_CATCH or GIMPLE_TRY_FINALLY.
|
|
|
|
EVAL is the sequence of statements to execute on entry to GIMPLE_TRY.
|
|
|
|
CLEANUP is the sequence of statements to execute according to
|
|
TRY_KIND. If TRY_KIND is GIMPLE_TRY_CATCH, CLEANUP is only exected
|
|
if an exception is thrown during execution of EVAL. If TRY_KIND is
|
|
GIMPLE_TRY_FINALLY, CLEANUP is always executed after executing EVAL
|
|
(regardless of whether EVAL finished normally, or jumped out or an
|
|
exception was thrown). */
|
|
DEFGSCODE(GIMPLE_TRY, "gimple_try", GSS_TRY)
|
|
|
|
/* GIMPLE_NOP represents the "do nothing" statement. */
|
|
DEFGSCODE(GIMPLE_NOP, "gimple_nop", GSS_BASE)
|
|
|
|
|
|
/* IMPORTANT.
|
|
|
|
Do not rearrange any of the GIMPLE_OMP_* codes. This ordering is
|
|
exposed by the range check in gimple_omp_subcode(). */
|
|
|
|
|
|
/* Tuples used for lowering of OMP_ATOMIC. Although the form of the OMP_ATOMIC
|
|
expression is very simple (just in form mem op= expr), various implicit
|
|
conversions may cause the expression to become more complex, so that it does
|
|
not fit the gimple grammar very well. To overcome this problem, OMP_ATOMIC
|
|
is rewritten as a sequence of two codes in gimplification:
|
|
|
|
GIMPLE_OMP_LOAD (tmp, mem)
|
|
val = some computations involving tmp;
|
|
GIMPLE_OMP_STORE (val). */
|
|
DEFGSCODE(GIMPLE_OMP_ATOMIC_LOAD, "gimple_omp_atomic_load",
|
|
GSS_OMP_ATOMIC_LOAD)
|
|
DEFGSCODE(GIMPLE_OMP_ATOMIC_STORE, "gimple_omp_atomic_store",
|
|
GSS_OMP_ATOMIC_STORE_LAYOUT)
|
|
|
|
/* GIMPLE_OMP_CONTINUE marks the location of the loop or sections
|
|
iteration in partially lowered OpenMP code. */
|
|
DEFGSCODE(GIMPLE_OMP_CONTINUE, "gimple_omp_continue", GSS_OMP_CONTINUE)
|
|
|
|
/* GIMPLE_OMP_CRITICAL <NAME, BODY> represents
|
|
|
|
#pragma omp critical [name]
|
|
|
|
NAME is the name given to the critical section.
|
|
BODY is the sequence of statements that are inside the critical section. */
|
|
DEFGSCODE(GIMPLE_OMP_CRITICAL, "gimple_omp_critical", GSS_OMP_CRITICAL)
|
|
|
|
/* GIMPLE_OMP_FOR <BODY, CLAUSES, INDEX, INITIAL, FINAL, COND, INCR, PRE_BODY>
|
|
represents
|
|
|
|
PRE_BODY
|
|
#pragma omp for [clause1 ... clauseN]
|
|
for (INDEX = INITIAL; INDEX COND FINAL; INDEX {+=,-=} INCR)
|
|
BODY
|
|
|
|
Likewise for:
|
|
#pragma acc loop [clause1 ... clauseN]
|
|
|
|
BODY is the loop body.
|
|
|
|
CLAUSES is the list of clauses.
|
|
|
|
INDEX must be an integer or pointer variable, which is implicitly thread
|
|
private. It must be accepted by is_gimple_operand.
|
|
|
|
INITIAL is the initial value given to INDEX. It must be
|
|
accepted by is_gimple_operand.
|
|
|
|
FINAL is the final value that INDEX should take. It must
|
|
be accepted by is_gimple_operand.
|
|
|
|
COND is the condition code for the controlling predicate. It must
|
|
be one of { <, >, <=, >= }
|
|
|
|
INCR is the loop index increment. It must be tree node of type
|
|
tcc_constant.
|
|
|
|
PRE_BODY is a landing pad filled by the gimplifier with things from
|
|
INIT, COND, and INCR that are technically part of the OMP_FOR
|
|
structured block, but are evaluated before the loop body begins.
|
|
|
|
INITIAL, FINAL and INCR are required to be loop invariant integer
|
|
expressions that are evaluated without any synchronization.
|
|
The evaluation order, frequency of evaluation and side-effects are
|
|
unspecified by the standards. */
|
|
DEFGSCODE(GIMPLE_OMP_FOR, "gimple_omp_for", GSS_OMP_FOR)
|
|
|
|
/* GIMPLE_OMP_MASTER <BODY> represents #pragma omp master.
|
|
BODY is the sequence of statements to execute in the master section. */
|
|
DEFGSCODE(GIMPLE_OMP_MASTER, "gimple_omp_master", GSS_OMP)
|
|
|
|
/* GIMPLE_OMP_TASKGROUP <BODY> represents #pragma omp taskgroup.
|
|
BODY is the sequence of statements to execute in the taskgroup section. */
|
|
DEFGSCODE(GIMPLE_OMP_TASKGROUP, "gimple_omp_taskgroup", GSS_OMP)
|
|
|
|
/* GIMPLE_OMP_PARALLEL <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents
|
|
|
|
#pragma omp parallel [CLAUSES]
|
|
BODY
|
|
|
|
BODY is a the sequence of statements to be executed by all threads.
|
|
|
|
CLAUSES is an OMP_CLAUSE chain with all the clauses.
|
|
|
|
CHILD_FN is set when outlining the body of the parallel region.
|
|
All the statements in BODY are moved into this newly created
|
|
function when converting OMP constructs into low-GIMPLE.
|
|
|
|
DATA_ARG is a local variable in the parent function containing data
|
|
to be shared with CHILD_FN. This is used to implement all the data
|
|
sharing clauses. */
|
|
DEFGSCODE(GIMPLE_OMP_PARALLEL, "gimple_omp_parallel", GSS_OMP_PARALLEL_LAYOUT)
|
|
|
|
/* GIMPLE_OMP_TASK <BODY, CLAUSES, CHILD_FN, DATA_ARG, COPY_FN,
|
|
ARG_SIZE, ARG_ALIGN> represents
|
|
|
|
#pragma omp task [CLAUSES]
|
|
BODY
|
|
|
|
BODY is a the sequence of statements to be executed by all threads.
|
|
|
|
CLAUSES is an OMP_CLAUSE chain with all the clauses.
|
|
|
|
CHILD_FN is set when outlining the body of the explicit task region.
|
|
All the statements in BODY are moved into this newly created
|
|
function when converting OMP constructs into low-GIMPLE.
|
|
|
|
DATA_ARG is a local variable in the parent function containing data
|
|
to be shared with CHILD_FN. This is used to implement all the data
|
|
sharing clauses.
|
|
|
|
COPY_FN is set when outlining the firstprivate var initialization.
|
|
All the needed statements are emitted into the newly created
|
|
function, or when only memcpy is needed, it is NULL.
|
|
|
|
ARG_SIZE and ARG_ALIGN are the size and alignment of the incoming
|
|
data area allocated by GOMP_task and passed to CHILD_FN. */
|
|
DEFGSCODE(GIMPLE_OMP_TASK, "gimple_omp_task", GSS_OMP_TASK)
|
|
|
|
/* OMP_RETURN marks the end of an OpenMP directive. */
|
|
DEFGSCODE(GIMPLE_OMP_RETURN, "gimple_omp_return", GSS_OMP_ATOMIC_STORE_LAYOUT)
|
|
|
|
/* OMP_SECTION <BODY> represents #pragma omp section.
|
|
BODY is the sequence of statements in the section body. */
|
|
DEFGSCODE(GIMPLE_OMP_SECTION, "gimple_omp_section", GSS_OMP)
|
|
|
|
/* OMP_SECTIONS <BODY, CLAUSES, CONTROL> represents #pragma omp sections.
|
|
|
|
BODY is the sequence of statements in the sections body.
|
|
CLAUSES is an OMP_CLAUSE chain holding the list of associated clauses.
|
|
CONTROL is a VAR_DECL used for deciding which of the sections
|
|
to execute. */
|
|
DEFGSCODE(GIMPLE_OMP_SECTIONS, "gimple_omp_sections", GSS_OMP_SECTIONS)
|
|
|
|
/* GIMPLE_OMP_SECTIONS_SWITCH is a marker placed immediately after
|
|
OMP_SECTIONS. It represents the GIMPLE_SWITCH used to decide which
|
|
branch is taken. */
|
|
DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE)
|
|
|
|
/* GIMPLE_OMP_SINGLE <BODY, CLAUSES> represents #pragma omp single
|
|
BODY is the sequence of statements inside the single section.
|
|
CLAUSES is an OMP_CLAUSE chain holding the associated clauses. */
|
|
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
|
|
|
|
/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
|
|
#pragma acc {kernels,parallel,data,enter data,exit data,update}
|
|
#pragma omp target {,data,update}
|
|
BODY is the sequence of statements inside the construct
|
|
(NULL for some variants).
|
|
CLAUSES is an OMP_CLAUSE chain holding the associated clauses.
|
|
CHILD_FN is set when outlining the body of the offloaded region.
|
|
All the statements in BODY are moved into this newly created
|
|
function when converting OMP constructs into low-GIMPLE.
|
|
DATA_ARG is a vec of 3 local variables in the parent function
|
|
containing data to be mapped to CHILD_FN. This is used to
|
|
implement the MAP clauses. */
|
|
DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL_LAYOUT)
|
|
|
|
/* GIMPLE_OMP_TEAMS <BODY, CLAUSES> 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)
|
|
|
|
/* GIMPLE_OMP_ORDERED <BODY, CLAUSES> 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_GRID_BODY <BODY> represents a parallel loop lowered for execution
|
|
on a GPU. It is an artificial statement created by omp lowering. */
|
|
DEFGSCODE(GIMPLE_OMP_GRID_BODY, "gimple_omp_gpukernel", GSS_OMP)
|
|
|
|
/* GIMPLE_PREDICT <PREDICT, OUTCOME> specifies a hint for branch prediction.
|
|
|
|
PREDICT is one of the predictors from predict.def.
|
|
|
|
OUTCOME is NOT_TAKEN or TAKEN. */
|
|
DEFGSCODE(GIMPLE_PREDICT, "gimple_predict", GSS_BASE)
|
|
|
|
/* This node represents a cleanup expression. It is ONLY USED INTERNALLY
|
|
by the gimplifier as a placeholder for cleanups, and its uses will be
|
|
cleaned up by the time gimplification is done.
|
|
|
|
This tuple should not exist outside of the gimplifier proper. */
|
|
DEFGSCODE(GIMPLE_WITH_CLEANUP_EXPR, "gimple_with_cleanup_expr", GSS_WCE)
|