This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[HSA] Introduce support for shared libraries and host fallback
- From: Martin LiÅka <mliska at suse dot cz>
- To: gcc-patches at gcc dot gnu dot org
- Date: Thu, 8 Oct 2015 18:47:26 +0200
- Subject: [HSA] Introduce support for shared libraries and host fallback
- Authentication-results: sourceware.org; auth=none
Hello.
Following patch set introduces HSA support for BRIG shared libraries. Apart from that,
it adds new warning option (-Whsa) that pops up when HSA code generation cannot expand
a function. Moreover, remaining of patches are follow-up of previous big changes.
Thanks,
Martin
>From 89d0a81c84cbbc18af05a6c144ec5f84fbd55a36 Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Fri, 2 Oct 2015 10:44:00 +0200
Subject: [PATCH 1/9] HSA: introduce warnings and libgomp host fallback
support.
gcc/ChangeLog:
2015-10-02 Martin Liska <mliska@suse.cz>
* common.opt: Add new Whsa option.
* hsa-gen.c (hsa_function_representation::hsa_function_representation):
Add new member seen_error.
(hsa_type_for_scalar_tree_type): Use HSA_SORRY_AT{V} instread of
sorry.
(hsa_type_for_tree_type): Likewise.
(hsa_op_immed::hsa_op_immed): Likewise.
(process_mem_base): Likewise.
(gen_hsa_addr): Likewise.
(gen_hsa_insns_for_load): Likewise.
(gen_hsa_insns_for_store): Likewise.
(gen_hsa_ctor_assignment): Likewise.
(gen_hsa_insns_for_single_assignment): Likewise.
(gen_hsa_cmp_insn_from_gimple): Likewise.
(gen_hsa_insns_for_operation_assignment): Likewise.
(verify_function_arguments): Likewise.
(gen_hsa_insns_for_direct_call): Likewise.
(gen_hsa_insns_for_return): Likewise.
(get_address_from_value): Likewise.
(gen_hsa_insns_for_call): Likewise.
(gen_hsa_insns_for_gimple_stmt): Likewise.
(gen_hsa_phi_from_gimple_phi): Likewise.
(gen_body_from_gimple): Use aforementioned hsa_cfun->seen_error.
(generate_hsa): Likewise.
* hsa.c (hsa_deinit_compilation_unit_data): Release
hsa_failed_functions.
(hsa_seen_error): New function.
* hsa.h (hsa_failed_functions): New variable.
libgomp/ChangeLog:
2015-10-02 Martin Liska <mliska@suse.cz>
* libgomp.h (struct gomp_device_descr): Add new function
pointer (can_run_func).
* plugin/plugin-hsa.c (init_enviroment_variables): Rename this
function from init_debug.
(hsa_warn): New function.
(struct kernel_info): Add new member variable
initialization_failed.
(struct agent_info): Add new member variable
prog_finalized_error.
(get_kernel_in_module): Do not call GOMP_PLUGIN_fatal.
(init_hsa_context): Use HSA_DEBUG macro.
(GOMP_OFFLOAD_init_device): Likewise.
(destroy_hsa_program): Likewise.
(GOMP_OFFLOAD_load_image): Produce warnings instread
of failures.
(GOMP_OFFLOAD_unload_image): Do not check if the agent
is initialized, the check is in the called function.
(GOMP_OFFLOAD_fini_device): Likewise.
(create_and_finalize_hsa_program): Likewise.
(release_kernel_dispatch): Use HSA_DEBUG macro.
(init_single_kernel): Produce warnings instread of failures.
(init_kernel): Use HSA_DEBUG macro.
(parse_launch_attributes): Likewise.
(GOMP_OFFLOAD_can_run): New function.
(GOMP_OFFLOAD_run): Remove part of initialization that
is moved to GOMP_OFFLOAD_can_run.
(GOMP_OFFLOAD_fini_device): Fix coding style.
* target.c (run_on_host): New function.
(GOMP_target): Use the function.
(gomp_load_plugin_for_device): Dynamically load the new hook.
---
gcc/common.opt | 4 +
gcc/hsa-gen.c | 238 ++++++++++++++++++++++++++++----------------
gcc/hsa.c | 14 ++-
gcc/hsa.h | 4 +
libgomp/libgomp.h | 1 +
libgomp/plugin/plugin-hsa.c | 230 ++++++++++++++++++++++++++++--------------
libgomp/target.c | 42 +++++---
7 files changed, 359 insertions(+), 174 deletions(-)
diff --git a/gcc/common.opt b/gcc/common.opt
index 7b0ec96..5ef6f46 100644
--- a/gcc/common.opt
+++ b/gcc/common.opt
@@ -581,6 +581,10 @@ Wfree-nonheap-object
Common Var(warn_free_nonheap_object) Init(1) Warning
Warn when attempting to free a non-heap object
+Whsa
+Common Var(warn_hsa) Init(1) Warning
+Warn when a function cannot be expanded to HSAIL
+
Winline
Common Var(warn_inline) Warning
Warn when an inlined function cannot be inlined
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 291d650..fb17a25 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -77,6 +77,32 @@ along with GCC; see the file COPYING3. If not see
#include "cfgloop.h"
#include "cfganal.h"
+/* Print a warning message and set that we have seen an error. */
+
+#define HSA_SORRY_MSG "could not emit HSAIL for the function"
+
+#define HSA_SORRY_ATV(location, message, ...) \
+ do \
+ { \
+ hsa_cfun->seen_error = true; \
+ if (warning_at (EXPR_LOCATION (hsa_cfun->decl), OPT_Whsa, \
+ HSA_SORRY_MSG)) \
+ inform (location, message, ##__VA_ARGS__); \
+ } \
+ while (false);
+
+/* Same as previous, but highlight a location. */
+
+#define HSA_SORRY_AT(location, message) \
+ do \
+ { \
+ hsa_cfun->seen_error = true; \
+ if (warning_at (EXPR_LOCATION (hsa_cfun->decl), OPT_Whsa, \
+ HSA_SORRY_MSG)) \
+ inform (location, message); \
+ } \
+ while (false);
+
/* Following structures are defined in the final version
of HSA specification. */
@@ -196,6 +222,7 @@ hsa_function_representation::hsa_function_representation ()
shadow_reg = NULL;
kernel_dispatch_count = 0;
maximum_omp_data_size = 0;
+ seen_error = false;
}
/* Destructor of class holding function/kernel-wide informaton and state. */
@@ -439,8 +466,9 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
if (!tree_fits_uhwi_p (TYPE_SIZE (base)))
{
- sorry ("Support for HSA does not implement huge or variable-sized type %T",
- type);
+ HSA_SORRY_ATV (EXPR_LOCATION (type),
+ "support for HSA does not implement huge or "
+ "variable-sized type %T", type);
return res;
}
@@ -468,7 +496,8 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
if (res == BRIG_TYPE_NONE)
{
- sorry ("Support for HSA does not implement type %T", type);
+ HSA_SORRY_ATV (EXPR_LOCATION (type),
+ "support for HSA does not implement type %T", type);
return res;
}
@@ -478,8 +507,9 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
if (bsize == tsize)
{
- sorry ("Support for HSA does not implement a vector type where "
- "a type and unit size are equal: %T", type);
+ HSA_SORRY_ATV (EXPR_LOCATION (type),
+ "support for HSA does not implement a vector type "
+ "where a type and unit size are equal: %T", type);
return res;
}
@@ -495,7 +525,8 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
res |= BRIG_TYPE_PACK_128;
break;
default:
- sorry ("Support for HSA does not implement type %T", type);
+ HSA_SORRY_ATV (EXPR_LOCATION (type),
+ "support for HSA does not implement type %T", type);
}
}
@@ -538,8 +569,8 @@ hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
gcc_checking_assert (TYPE_P (type));
if (!tree_fits_uhwi_p (TYPE_SIZE_UNIT (type)))
{
- sorry ("Support for HSA does not implement huge or variable-sized type %T",
- type);
+ HSA_SORRY_ATV (EXPR_LOCATION (type), "support for HSA does not "
+ "implement huge or variable-sized type %T", type);
return BRIG_TYPE_NONE;
}
@@ -566,8 +597,9 @@ hsa_type_for_tree_type (const_tree type, unsigned HOST_WIDE_INT *dim_p = NULL,
|| !tree_fits_shwi_p (TYPE_MIN_VALUE (domain))
|| !tree_fits_shwi_p (TYPE_MAX_VALUE (domain)))
{
- sorry ("Support for HSA does not implement array %T with unknown "
- "bounds", type);
+ HSA_SORRY_ATV (EXPR_LOCATION (type),
+ "support for HSA does not implement array %T with "
+ "unknown bounds", type);
return BRIG_TYPE_NONE;
}
HOST_WIDE_INT min = tree_to_shwi (TYPE_MIN_VALUE (domain));
@@ -775,7 +807,7 @@ hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
hsa_type_for_tree_type (TREE_TYPE (tree_val), NULL,
min32int))
{
- if (seen_error ())
+ if (hsa_seen_error ())
return;
gcc_checking_assert ((is_gimple_min_invariant (tree_val)
@@ -798,7 +830,8 @@ hsa_op_immed::hsa_op_immed (tree tree_val, bool min32int)
tree v = CONSTRUCTOR_ELT (tree_value, i)->value;
if (!CONSTANT_CLASS_P (v))
{
- sorry ("HSA ctor should have only constants");
+ HSA_SORRY_AT (EXPR_LOCATION (tree_val),
+ "HSA ctor should have only constants");
return;
}
}
@@ -1636,8 +1669,9 @@ process_mem_base (tree base, hsa_symbol **symbol, BrigType16_t *addrtype,
if (!DECL_P (decl) || TREE_CODE (decl) == FUNCTION_DECL)
{
- sorry ("Support for HSA does not implement a memory reference to "
- "a non-declaration type");
+ HSA_SORRY_AT (EXPR_LOCATION (base),
+ "support for HSA does not implement a memory reference "
+ "to a non-declaration type");
return;
}
@@ -1682,8 +1716,9 @@ gen_hsa_addr (tree ref, hsa_bb *hbb, vec <hsa_op_reg_p> *ssa_map,
&& ((tree_to_uhwi (TREE_OPERAND (ref, 1)) % BITS_PER_UNIT) != 0
|| (tree_to_uhwi (TREE_OPERAND (ref, 2)) % BITS_PER_UNIT) != 0))
{
- sorry ("Support for HSA does not implement bit field references "
- "such as %E", ref);
+ HSA_SORRY_ATV (EXPR_LOCATION (origref),
+ "support for HSA does not implement "
+ "bit field references such as %E", ref);
goto out;
}
@@ -1758,11 +1793,13 @@ gen_hsa_addr (tree ref, hsa_bb *hbb, vec <hsa_op_reg_p> *ssa_map,
offset += wi::to_offset (TMR_OFFSET (ref));
break;
case FUNCTION_DECL:
- sorry ("HSA does not support indirect calls");
+ HSA_SORRY_AT (EXPR_LOCATION (origref),
+ "support for HSA does not implement function pointers");
goto out;
case SSA_NAME:
default:
- sorry ("Support for HSA does not implement memory access to %E", origref);
+ HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does "
+ "not implement memory access to %E", origref);
goto out;
}
@@ -1797,8 +1834,8 @@ out:
bitsize = 0;
if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
- sorry ("Support for HSA does not implement unhandled bit field reference "
- "such as %E", ref);
+ HSA_SORRY_ATV (EXPR_LOCATION (origref), "support for HSA does not "
+ "implement unhandled bit field reference such as %E", ref);
if (output_bitsize != NULL && output_bitpos != NULL)
{
@@ -2014,8 +2051,9 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb,
{
if (dest->type != hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT))
{
- sorry ("Support for HSA does not implement conversion of %E to "
- "the requested non-pointer type.", rhs);
+ HSA_SORRY_ATV (EXPR_LOCATION (rhs),
+ "support for HSA does not implement conversion "
+ "of %E to the requested non-pointer type.", rhs);
return;
}
@@ -2100,8 +2138,9 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb,
/* Handle load of a bit field. */
if (bitsize > 64)
{
- sorry ("Support for HSA does not implement load from a bit field "
- "bigger than 64 bits");
+ HSA_SORRY_AT (EXPR_LOCATION (rhs),
+ "support for HSA does not implement load from a bit "
+ "field bigger than 64 bits");
return;
}
@@ -2119,8 +2158,9 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb,
}
}
else
- sorry ("Support for HSA does not implement loading of expression %E",
- rhs);
+ HSA_SORRY_ATV
+ (EXPR_LOCATION (rhs),
+ "support for HSA does not implement loading of expression %E", rhs);
}
/* Return number of bits necessary for representation of a bit field,
@@ -2158,8 +2198,9 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb,
/* Handle store to a bit field. */
if (bitsize > 64)
{
- sorry ("Support for HSA does not implement store to a bit field "
- "bigger than 64 bits");
+ HSA_SORRY_AT (EXPR_LOCATION (lhs),
+ "support for HSA does not implement store to a bit field "
+ "bigger than 64 bits");
return;
}
@@ -2366,7 +2407,8 @@ gen_hsa_ctor_assignment (hsa_op_address *addr_lhs, tree rhs, hsa_bb *hbb)
{
if (vec_safe_length (CONSTRUCTOR_ELTS (rhs)))
{
- sorry ("Support for HSA does not implement load from constructor");
+ HSA_SORRY_AT (EXPR_LOCATION (rhs),
+ "support for HSA does not implement load from constructor");
return;
}
@@ -2385,7 +2427,7 @@ gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb,
if (TREE_CODE (lhs) == SSA_NAME)
{
hsa_op_reg *dest = hsa_reg_for_gimple_ssa (lhs, ssa_map);
- if (seen_error ())
+ if (hsa_seen_error ())
return;
gen_hsa_insns_for_load (dest, rhs, TREE_TYPE (lhs), hbb, ssa_map);
@@ -2395,7 +2437,7 @@ gen_hsa_insns_for_single_assignment (tree lhs, tree rhs, hsa_bb *hbb,
{
/* Store to memory. */
hsa_op_base *src = hsa_reg_or_immed_for_gimple_op (rhs, hbb, ssa_map);
- if (seen_error ())
+ if (hsa_seen_error ())
return;
gen_hsa_insns_for_store (lhs, src, hbb, ssa_map);
@@ -2541,8 +2583,9 @@ gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
break;
default:
- sorry ("Support for HSA does not implement comparison tree code %s\n",
- get_tree_code_name (code));
+ HSA_SORRY_ATV (EXPR_LOCATION (lhs),
+ "support for HSA does not implement comparison tree "
+ "code %s\n", get_tree_code_name (code));
return;
}
@@ -2648,8 +2691,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb,
case CEIL_DIV_EXPR:
case FLOOR_DIV_EXPR:
case ROUND_DIV_EXPR:
- sorry ("Support for HSA does not implement CEIL_DIV_EXPR, FLOOR_DIV_EXPR "
- "or ROUND_DIV_EXPR");
+ HSA_SORRY_AT (gimple_location (assign),
+ "support for HSA does not implement CEIL_DIV_EXPR, "
+ "FLOOR_DIV_EXPR or ROUND_DIV_EXPR");
return;
case TRUNC_MOD_EXPR:
opcode = BRIG_OPCODE_REM;
@@ -2657,8 +2701,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb,
case CEIL_MOD_EXPR:
case FLOOR_MOD_EXPR:
case ROUND_MOD_EXPR:
- sorry ("Support for HSA does not implement CEIL_MOD_EXPR, FLOOR_MOD_EXPR "
- "or ROUND_MOD_EXPR");
+ HSA_SORRY_AT (gimple_location (assign),
+ "support for HSA does not implement CEIL_MOD_EXPR, "
+ "FLOOR_MOD_EXPR or ROUND_MOD_EXPR");
return;
case NEGATE_EXPR:
opcode = BRIG_OPCODE_NEG;
@@ -2850,8 +2895,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb,
}
default:
/* Implement others as we come across them. */
- sorry ("Support for HSA does not implement operation %s",
- get_tree_code_name (code));
+ HSA_SORRY_ATV (gimple_location (assign),
+ "support for HSA does not implement operation %s",
+ get_tree_code_name (code));
return;
}
@@ -3012,13 +3058,15 @@ verify_function_arguments (tree decl)
{
if (DECL_STATIC_CHAIN (decl))
{
- sorry ("HSA does not support nested functions: %D", decl);
+ HSA_SORRY_ATV (EXPR_LOCATION (decl),
+ "HSA does not support nested functions: %D", decl);
return;
}
else if (!TYPE_ARG_TYPES (TREE_TYPE (decl)))
{
- sorry ("HSA does not support functions with variadic arguments "
- "(or unknown return type): %D", decl);
+ HSA_SORRY_ATV (EXPR_LOCATION (decl),
+ "HSA does not support functions with variadic arguments "
+ "(or unknown return type): %D", decl);
return;
}
}
@@ -3034,7 +3082,7 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
{
tree decl = gimple_call_fndecl (stmt);
verify_function_arguments (decl);
- if (seen_error ())
+ if (hsa_seen_error ())
return;
hsa_insn_call *call_insn = new hsa_insn_call (decl);
@@ -3053,8 +3101,9 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
if (AGGREGATE_TYPE_P (TREE_TYPE (parm)))
{
- sorry ("Support for HSA does not implement an aggregate argument "
- "in a function call");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "support for HSA does not "
+ "implement an aggregate argument in a function call");
return;
}
@@ -3081,8 +3130,9 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
{
if (AGGREGATE_TYPE_P (result_type))
{
- sorry ("Support for HSA does not implement returning a value "
- "which is of an aggregate type: %T", result_type);
+ HSA_SORRY_ATV (gimple_location (stmt),
+ "support for HSA does not implement returning a value "
+ "which is of an aggregate type %T", result_type);
return;
}
@@ -3109,8 +3159,9 @@ gen_hsa_insns_for_direct_call (gimple *stmt, hsa_bb *hbb,
{
if (result)
{
- sorry ("Support for HSA does not implement an assignment of return "
- "value from a void function");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "support for HSA does not implement an assignment of "
+ "return value from a void function");
return;
}
@@ -3137,8 +3188,9 @@ gen_hsa_insns_for_return (greturn *stmt, hsa_bb *hbb,
{
if (AGGREGATE_TYPE_P (TREE_TYPE (retval)))
{
- sorry ("HSA does not support return statement with an aggregate "
- "value type");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "HSA does not support return "
+ "statement with an aggregate value type");
return;
}
@@ -3834,7 +3886,9 @@ get_address_from_value (tree val, hsa_bb *hbb, vec <hsa_op_reg_p> *ssa_map)
/* Otherwise fall-through */
default:
- sorry ("Support for HSA does not implement memory access to %E", val);
+ HSA_SORRY_ATV (EXPR_LOCATION (val),
+ "support for HSA does not implement memory access to %E",
+ val);
return new hsa_op_address (NULL, NULL, 0);
}
}
@@ -3968,15 +4022,17 @@ gen_hsa_insns_for_call (gimple *stmt, hsa_bb *hbb,
tree function_decl = gimple_call_fndecl (stmt);
if (function_decl == NULL_TREE)
{
- sorry ("HSA does not support indirect calls");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "support for HSA does not implement indirect calls");
return;
}
if (hsa_callable_function_p (function_decl))
gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
else if (!gen_hsa_insns_for_known_library_call (stmt, hbb, ssa_map))
- sorry ("HSA does support only call for functions with 'hsafunc' "
- "attribute");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "HSA does support only call of functions within omp "
+ "declare target or with 'hsafunc' attribute");
return;
}
@@ -4292,8 +4348,9 @@ specialop:
if (TREE_CODE (byte_size) != INTEGER_CST)
{
- sorry ("Support for HSA does not implement __builtin_memcpy with "
- "a non constant size");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "support for HSA does not implement __builtin_memcpy "
+ "with a non constant size");
return;
}
@@ -4302,9 +4359,11 @@ specialop:
/* TODO: fallback to call to memcpy library function. */
if (n > HSA_MEMORY_BUILTINS_LIMIT)
{
- sorry ("Support for HSA does implement __builtin_memcpy with a size"
- " bigger than %u bytes, %u bytes are requested",
- HSA_MEMORY_BUILTINS_LIMIT, n);
+ HSA_SORRY_ATV
+ (gimple_location (stmt),
+ "support for HSA does implement __builtin_memcpy with a size "
+ "bigger than %u bytes, %u bytes are requested",
+ HSA_MEMORY_BUILTINS_LIMIT, n);
return;
}
@@ -4329,8 +4388,10 @@ specialop:
if (TREE_CODE (c) != INTEGER_CST)
{
- sorry ("Support for HSA does not implement __builtin_memset with "
- "a non constant byte value that should be written");
+ HSA_SORRY_AT
+ (gimple_location (stmt),
+ "support for HSA does not implement __builtin_memset with a "
+ "non constant byte value that should be written");
return;
}
@@ -4338,8 +4399,9 @@ specialop:
if (TREE_CODE (byte_size) != INTEGER_CST)
{
- sorry ("Support for HSA does not implement __builtin_memset with "
- "a non constant size");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "support for HSA does not implement "
+ "__builtin_memset with a non constant size");
return;
}
@@ -4348,9 +4410,11 @@ specialop:
/* TODO: fallback to call to memset library function. */
if (n > HSA_MEMORY_BUILTINS_LIMIT)
{
- sorry ("Support for HSA does implement __builtin_memset with a size"
- " bigger than %u bytes, %u bytes are requested",
- HSA_MEMORY_BUILTINS_LIMIT, n);
+ HSA_SORRY_ATV
+ (gimple_location (stmt),
+ "support for HSA does implement __builtin_memset with a size "
+ "bigger than %u bytes, %u bytes are requested",
+ HSA_MEMORY_BUILTINS_LIMIT, n);
return;
}
@@ -4369,8 +4433,9 @@ specialop:
break;
}
default:
- sorry ("Support for HSA does not implement calls to builtin %D",
- gimple_call_fndecl (stmt));
+ HSA_SORRY_ATV (gimple_location (stmt),
+ "support for HSA does not implement calls to builtin %D",
+ gimple_call_fndecl (stmt));
return;
}
}
@@ -4413,8 +4478,9 @@ gen_hsa_insns_for_gimple_stmt (gimple *stmt, hsa_bb *hbb,
{
tree label = gimple_label_label (as_a <glabel *> (stmt));
if (FORCED_LABEL (label))
- sorry ("Support for HSA does not implement gimple label with address "
- "taken");
+ HSA_SORRY_AT (gimple_location (stmt),
+ "support for HSA does not implement gimple label with "
+ "address taken");
break;
}
@@ -4429,8 +4495,9 @@ gen_hsa_insns_for_gimple_stmt (gimple *stmt, hsa_bb *hbb,
break;
}
default:
- sorry ("Support for HSA does not implement gimple statement %s",
- gimple_code_name[(int) gimple_code (stmt)]);
+ HSA_SORRY_ATV (gimple_location (stmt),
+ "support for HSA does not implement gimple statement %s",
+ gimple_code_name[(int) gimple_code (stmt)]);
}
}
@@ -4490,8 +4557,9 @@ gen_hsa_phi_from_gimple_phi (gimple *phi_stmt, hsa_bb *hbb,
}
else
{
- sorry ("Support for HSA does not handle PHI nodes with constant "
- "address operands");
+ HSA_SORRY_AT (gimple_location (phi_stmt),
+ "support for HSA does not handle PHI nodes with "
+ "constant address operands");
return;
}
}
@@ -4607,7 +4675,9 @@ gen_body_from_gimple (vec <hsa_op_reg_p> *ssa_map)
to the same basic block. */
if (e->flags & EDGE_EH)
{
- sorry ("Support for HSA does not implement exception handling");
+ HSA_SORRY_AT
+ (UNKNOWN_LOCATION,
+ "support for HSA does not implement exception handling");
return;
}
}
@@ -4621,7 +4691,7 @@ gen_body_from_gimple (vec <hsa_op_reg_p> *ssa_map)
for (gsi = gsi_start_bb (bb); !gsi_end_p (gsi); gsi_next (&gsi))
{
gen_hsa_insns_for_gimple_stmt (gsi_stmt (gsi), hbb, ssa_map);
- if (seen_error ())
+ if (hsa_seen_error ())
return;
}
}
@@ -5061,16 +5131,16 @@ emit_hsa_module_variables (void)
static void
generate_hsa (bool kernel)
{
+ vec <hsa_op_reg_p> ssa_map = vNULL;
+
if (hsa_num_threads == NULL)
emit_hsa_module_variables ();
+ hsa_init_data_for_cfun ();
verify_function_arguments (cfun->decl);
- if (seen_error ())
- return;
-
- vec <hsa_op_reg_p> ssa_map = vNULL;
+ if (hsa_seen_error ())
+ goto fail;
- hsa_init_data_for_cfun ();
hsa_cfun->decl = cfun->decl;
hsa_cfun->kern_p = kernel;
@@ -5080,13 +5150,13 @@ generate_hsa (bool kernel)
hsa_sanitize_name (hsa_cfun->name);
gen_function_def_parameters (hsa_cfun, &ssa_map);
- if (seen_error ())
+ if (hsa_seen_error ())
goto fail;
init_omp_in_prologue ();
gen_body_from_gimple (&ssa_map);
- if (seen_error ())
+ if (hsa_seen_error ())
goto fail;
if (hsa_cfun->kern_p)
diff --git a/gcc/hsa.c b/gcc/hsa.c
index ce8ae45..c938248 100644
--- a/gcc/hsa.c
+++ b/gcc/hsa.c
@@ -138,10 +138,8 @@ hsa_init_compilation_unit_data (void)
void
hsa_deinit_compilation_unit_data (void)
{
- if (!compilation_unit_data_initialized)
- return;
-
- delete hsa_global_variable_symbols;
+ if (compilation_unit_data_initialized)
+ delete hsa_global_variable_symbols;
}
/* Return true if we are generating large HSA machine model. */
@@ -647,4 +645,12 @@ hsa_register_kernel (cgraph_node *gpu, cgraph_node *host)
hsa_summaries->link_functions (gpu, host, HSA_KERNEL);
}
+/* Return true if expansion of the current HSA function has already failed. */
+
+bool
+hsa_seen_error (void)
+{
+ return hsa_cfun->seen_error;
+}
+
#include "gt-hsa.h"
diff --git a/gcc/hsa.h b/gcc/hsa.h
index ce6414a..b400b39 100644
--- a/gcc/hsa.h
+++ b/gcc/hsa.h
@@ -973,6 +973,9 @@ public:
OMP data size is necessary memory that is used for copying before
a kernel dispatch. */
unsigned maximum_omp_data_size;
+
+ /* Return true if there's an HSA-specific warning already seen. */
+ bool seen_error;
};
enum hsa_function_kind
@@ -1064,6 +1067,7 @@ char *hsa_brig_function_name (const char *p);
const char *hsa_get_declaration_name (tree decl);
void hsa_register_kernel (cgraph_node *host);
void hsa_register_kernel (cgraph_node *gpu, cgraph_node *host);
+bool hsa_seen_error (void);
/* In hsa-gen.c. */
void hsa_build_append_simple_mov (hsa_op_reg *, hsa_op_base *, hsa_bb *);
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index bb0ffab..88c4143 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -758,6 +758,7 @@ struct gomp_device_descr
void *(*dev2host_func) (int, void *, const void *, size_t);
void *(*host2dev_func) (int, void *, const void *, size_t);
void (*run_func) (int, void *, void *, const void *);
+ bool (*can_run_func) (void *);
/* Splay tree containing information about mapped memory regions. */
struct splay_tree_s mem_map;
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index 76a3b45..60dac54 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -49,15 +49,60 @@ GOMP_OFFLOAD_version (void)
static bool debug;
-/* Initialize debug according to the environment. */
+/* Flag to decide if the runtime should suppress a possible fallback to host
+ execution. */
+
+static bool suppress_host_fallback;
+
+/* Initialize debug and supress_host_fallback according to the environment. */
static void
-init_debug (void)
+init_enviroment_variables (void)
{
if (getenv ("HSA_DEBUG"))
debug = true;
else
debug = false;
+
+ if (getenv ("HSA_SUPPRESS_HOST_FALLBACK"))
+ suppress_host_fallback = true;
+ else
+ suppress_host_fallback = false;
+}
+
+/* Print a debug message to stderr if DEBUG value is set to true. */
+
+#define HSA_DEBUG(...) \
+ do \
+ { \
+ if (debug) \
+ { \
+ fprintf (stderr, "HSA debug: "); \
+ fprintf (stderr, __VA_ARGS__); \
+ } \
+ } \
+ while (false);
+
+/* Print HSA warning STR with an HSA STATUS code. */
+
+static void
+hsa_warn (const char *str, hsa_status_t status)
+{
+ if (!debug)
+ return;
+
+ const char* hsa_error;
+ hsa_status_string (status, &hsa_error);
+
+ unsigned l = strlen (hsa_error);
+
+ char *err = GOMP_PLUGIN_malloc (sizeof (char) * l);
+ memcpy (err, hsa_error, l - 1);
+ err[l] = '\0';
+
+ fprintf (stderr, "HSA warning: %s (%s)\n", str, err);
+
+ free (err);
}
/* Report a fatal error STR together with the HSA error corresponding to STATUS
@@ -111,6 +156,8 @@ struct kernel_info
/* Flag indicating whether the kernel has been initialized and all fields
below it contain valid data. */
bool initialized;
+ /* Flag indicating that the kernel has a problem that blocks an execution. */
+ bool initialization_failed;
/* The object to be put into the dispatch queue. */
uint64_t object;
/* Required size of kernel arguments. */
@@ -175,6 +222,8 @@ struct agent_info
/* Flag whether the HSA program that consists of all the modules has been
finalized. */
bool prog_finalized;
+ /* Flag whether the program was finalized but with a failture. */
+ bool prog_finalized_error;
/* HSA executable - the finalized program that is used to locate kernels. */
hsa_executable_t executable;
};
@@ -204,7 +253,6 @@ get_kernel_in_module (struct module_info *module, const char *kernel_name)
if (strcmp (module->kernels[i].name, kernel_name) == 0)
return &module->kernels[i];
- GOMP_PLUGIN_fatal ("Could not find kernel dependency: %s\n", kernel_name);
return NULL;
}
@@ -271,17 +319,15 @@ init_hsa_context (void)
if (hsa_context.initialized)
return;
- init_debug ();
+ init_enviroment_variables ();
status = hsa_init ();
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Run-time could not be initialized", status);
- if (debug)
- fprintf (stderr, "HSA run-time initialized\n");
+ HSA_DEBUG ("HSA run-time initialized\n");
status = hsa_iterate_agents (count_gpu_agents, NULL);
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("HSA GPU devices could not be enumerated", status);
- if (debug)
- fprintf (stderr, "There are %i HSA GPU devices.\n", hsa_context.agent_count);
+ HSA_DEBUG ("There are %i HSA GPU devices.\n", hsa_context.agent_count);
hsa_context.agents
= GOMP_PLUGIN_malloc_cleared (hsa_context.agent_count
@@ -379,8 +425,7 @@ GOMP_OFFLOAD_init_device (int n)
if (agent->kernarg_region.handle == (uint64_t) -1)
GOMP_PLUGIN_fatal ("Could not find suitable memory region for kernel "
"arguments");
- if (debug)
- fprintf (stderr, "HSA agent initialized, queue has id %llu\n",
+ HSA_DEBUG ("HSA agent initialized, queue has id %llu\n",
(long long unsigned) agent->command_q->id);
agent->initialized = true;
}
@@ -431,10 +476,12 @@ remove_module_from_agent (struct agent_info *agent, struct module_info *module)
static void
destroy_hsa_program (struct agent_info *agent)
{
+ if (!agent->prog_finalized || agent->prog_finalized_error)
+ return;
+
hsa_status_t status;
- if (debug)
- fprintf (stderr, "Destroying the current HSA program.\n");
+ HSA_DEBUG ("Destroying the current HSA program.\n");
status = hsa_executable_destroy (agent->executable);
if (status != HSA_STATUS_SUCCESS)
@@ -473,8 +520,7 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version __attribute__ ((unused)),
if (agent->prog_finalized)
destroy_hsa_program (agent);
- if (debug)
- fprintf (stderr, "Encountered %d kernels in an image\n", kernel_count);
+ HSA_DEBUG ("Encountered %d kernels in an image\n", kernel_count);
pair = GOMP_PLUGIN_malloc (kernel_count * sizeof (struct addr_pair));
*target_table = pair;
module = (struct module_info *)
@@ -523,19 +569,15 @@ create_and_finalize_hsa_program (struct agent_info *agent)
if (pthread_mutex_lock (&agent->prog_mutex))
GOMP_PLUGIN_fatal ("Could not lock an HSA agent program mutex");
if (agent->prog_finalized)
- {
- if (pthread_mutex_unlock (&agent->prog_mutex))
- GOMP_PLUGIN_fatal ("Could not unlock an HSA agent program mutex");
- return;
- }
+ goto final;
status = hsa_ext_program_create (HSA_MACHINE_MODEL_LARGE, HSA_PROFILE_FULL,
HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
NULL, &prog_handle);
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not create an HSA program", status);
- if (debug)
- fprintf (stderr, "Created a finalizer program\n");
+
+ HSA_DEBUG ("Created a finalized program\n");
struct module_info *module = agent->first_module;
while (module)
@@ -544,8 +586,6 @@ create_and_finalize_hsa_program (struct agent_info *agent)
module->image_desc->brig_module);
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not add a module to the HSA program", status);
- if (debug)
- fprintf (stderr, "Added module %i to the HSA program\n", mi);
module = module->next;
mi++;
}
@@ -558,9 +598,12 @@ create_and_finalize_hsa_program (struct agent_info *agent)
HSA_CODE_OBJECT_TYPE_PROGRAM,
&code_object);
if (status != HSA_STATUS_SUCCESS)
- hsa_fatal ("Finalization of the HSA program failed", status);
- if (debug)
- fprintf (stderr, "Finalization done\n");
+ {
+ hsa_warn ("Finalization of the HSA program failed", status);
+ goto failure;
+ }
+
+ HSA_DEBUG ("Finalization done\n");
hsa_ext_program_destroy (prog_handle);
status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN,
@@ -576,9 +619,17 @@ create_and_finalize_hsa_program (struct agent_info *agent)
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not freeze the HSA executable", status);
- if (debug)
- fprintf (stderr, "Froze HSA executable with the finalized code object\n");
+ HSA_DEBUG ("Froze HSA executable with the finalized code object\n");
+
+ /* If all goes good, jump to final. */
+ goto final;
+
+failure:
+ agent->prog_finalized_error = true;
+
+final:
agent->prog_finalized = true;
+
if (pthread_mutex_unlock (&agent->prog_mutex))
GOMP_PLUGIN_fatal ("Could not unlock an HSA agent program mutex");
}
@@ -626,8 +677,7 @@ create_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size)
static void
release_kernel_dispatch (struct hsa_kernel_dispatch *shadow)
{
- if (debug)
- fprintf (stderr, "Released kernel dispatch: %p has value: %lu (%p)\n",
+ HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n",
shadow, shadow->debug, (void *)shadow->debug);
hsa_memory_free (shadow->kernarg_address);
@@ -657,9 +707,11 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
status = hsa_executable_get_symbol (agent->executable, NULL, kernel->name,
agent->id, 0, &kernel_symbol);
if (status != HSA_STATUS_SUCCESS)
- hsa_fatal ("Could not find symbol for kernel in the code object", status);
- if (debug)
- fprintf (stderr, "Located kernel %s\n", kernel->name);
+ {
+ hsa_warn ("Could not find symbol for kernel in the code object", status);
+ goto failure;
+ }
+ HSA_DEBUG ("Located kernel %s\n", kernel->name);
status = hsa_executable_symbol_get_info
(kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel->object);
if (status != HSA_STATUS_SUCCESS)
@@ -678,22 +730,18 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
(kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
&kernel->private_segment_size);
if (status != HSA_STATUS_SUCCESS)
- hsa_fatal ("Could not get info about kernel private segment size", status);
-
- if (debug)
- {
- fprintf (stderr, "Kernel structure for %s fully initialized with "
- "following segment sizes: \n",
- kernel->name);
- fprintf (stderr, " group_segment_size: %u\n",
- (unsigned) kernel->group_segment_size);
- fprintf (stderr, " private_segment_size: %u\n",
- (unsigned) kernel->private_segment_size);
- fprintf (stderr, " kernarg_segment_size: %u\n",
- (unsigned) kernel->kernarg_segment_size);
- fprintf (stderr, " omp_data_size: %u\n",
- kernel->omp_data_size);
- }
+ hsa_fatal ("Could not get info about kernel private segment size",
+ status);
+
+ HSA_DEBUG ("Kernel structure for %s fully initialized with "
+ "following segment sizes: \n", kernel->name);
+ HSA_DEBUG (" group_segment_size: %u\n",
+ (unsigned) kernel->group_segment_size);
+ HSA_DEBUG (" private_segment_size: %u\n",
+ (unsigned) kernel->private_segment_size);
+ HSA_DEBUG (" kernarg_segment_size: %u\n",
+ (unsigned) kernel->kernarg_segment_size);
+ HSA_DEBUG (" omp_data_size: %u\n", kernel->omp_data_size);
if (kernel->omp_data_size > *max_omp_data_size)
*max_omp_data_size = kernel->omp_data_size;
@@ -704,8 +752,23 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
{
struct kernel_info *dependency = get_kernel_in_module
(module, kernel->dependencies[i]);
+
+ if (dependency == NULL)
+ {
+ HSA_DEBUG ("Could not find a dependency for a kernel: %s, "
+ "dependency name: %s\n", kernel->name,
+ kernel->dependencies[i]);
+ goto failure;
+ }
+
+
init_single_kernel (dependency, max_omp_data_size);
}
+
+ return;
+
+failure:
+ kernel->initialization_failed = true;
}
/* Indent stream F by INDENT spaces. */
@@ -800,8 +863,8 @@ init_kernel (struct kernel_info *kernel)
dispatch operation. */
init_single_kernel (kernel, &kernel->max_omp_data_size);
- if (debug)
- fprintf (stderr, "\n");
+ if (!kernel->initialization_failed)
+ HSA_DEBUG ("\n");
kernel->initialized = true;
if (pthread_mutex_unlock (&kernel->init_mutex))
@@ -840,8 +903,7 @@ parse_launch_attributes (const void *input,
def->wdims[1] = 1;
def->wdims[2] = 1;
*result = def;
- if (debug)
- fprintf (stderr, "GOMP_OFFLOAD_run called with no launch attributes\n");
+ HSA_DEBUG ("GOMP_OFFLOAD_run called with no launch attributes\n");
return true;
}
@@ -854,13 +916,37 @@ parse_launch_attributes (const void *input,
if (kla->gdims[0] == 0)
return false;
- if (debug)
- fprintf (stderr, "GOMP_OFFLOAD_run called with grid size %u and group "
- "size %u\n", kla->gdims[0], kla->wdims[0]);
+ HSA_DEBUG ("GOMP_OFFLOAD_run called with grid size %u and group size %u\n",
+ kla->gdims[0], kla->wdims[0]);
return true;
}
+/* Return true if the HSA runtime can run function FN_PTR. */
+
+bool
+GOMP_OFFLOAD_can_run (void *fn_ptr)
+{
+ struct kernel_info *kernel = (struct kernel_info *) fn_ptr;
+ struct agent_info *agent = kernel->agent;
+ create_and_finalize_hsa_program (agent);
+
+ if (agent->prog_finalized_error)
+ goto failure;
+
+ init_kernel (kernel);
+ if (kernel->initialization_failed)
+ goto failure;
+
+ return true;
+
+failure:
+ if (suppress_host_fallback)
+ GOMP_PLUGIN_fatal ("HSA host fallback has been suppressed");
+ HSA_DEBUG ("HSA target cannot be launched, doing a host fallback\n");
+ return false;
+}
+
/* Part of the libgomp plugin interface. Run a kernel on a device N and pass
the it an array of pointers in VARS as a parameter. The kernel is
identified by FN_PTR which must point to a kernel_info structure. */
@@ -874,16 +960,18 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, const void* kern_launch)
const struct kernel_launch_attributes *kla;
if (!parse_launch_attributes (kern_launch, &def, &kla))
{
- if (debug)
- fprintf (stderr, "Will not run HSA kernel because the grid size is "
- "zero\n");
+ HSA_DEBUG ("Will not run HSA kernel because the grid size is zero\n");
return;
}
if (pthread_rwlock_rdlock (&agent->modules_rwlock))
GOMP_PLUGIN_fatal ("Unable to read-lock an HSA agent rwlock");
- create_and_finalize_hsa_program (agent);
- init_kernel (kernel) ;
+ if (!agent->initialized)
+ GOMP_PLUGIN_fatal ("Agent must be initialized");
+
+ if (!kernel->initialized)
+ GOMP_PLUGIN_fatal ("Called kernel must be initialized");
+
struct hsa_kernel_dispatch *shadow = create_kernel_dispatch_recursive
(kernel, kernel->max_omp_data_size);
@@ -894,8 +982,7 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, const void* kern_launch)
}
uint64_t index = hsa_queue_add_write_index_release (agent->command_q, 1);
- if (debug)
- fprintf (stderr, "Got AQL index %llu\n", (long long int) index);
+ HSA_DEBUG ("Got AQL index %llu\n", (long long int) index);
/* Wait until the queue is not full before writing the packet. */
while (index - hsa_queue_load_read_index_acquire(agent->command_q)
@@ -933,22 +1020,19 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, const void* kern_launch)
memcpy (shadow->kernarg_address + sizeof (vars), &shadow,
sizeof (struct hsa_kernel_runtime *));
- if (debug)
- fprintf (stderr, "Copying kernel runtime pointer to kernarg_address\n");
+ HSA_DEBUG ("Copying kernel runtime pointer to kernarg_address\n");
uint16_t header;
header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
- if (debug)
- fprintf (stderr, "Going to dispatch kernel %s\n", kernel->name);
+ HSA_DEBUG ("Going to dispatch kernel %s\n", kernel->name);
__atomic_store_n ((uint16_t*)(&packet->header), header, __ATOMIC_RELEASE);
hsa_signal_store_release (agent->command_q->doorbell_signal, index);
- if (debug)
- fprintf (stderr, "Kernel dispatched, waiting for completion\n");
+ HSA_DEBUG ("Kernel dispatched, waiting for completion\n");
hsa_signal_wait_acquire (s, HSA_SIGNAL_CONDITION_LT, 1,
UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
@@ -999,8 +1083,7 @@ GOMP_OFFLOAD_unload_image (int n, unsigned version __attribute__ ((unused)),
remove_module_from_agent (agent, module);
destroy_module (module);
free (module);
- if (agent->prog_finalized)
- destroy_hsa_program (agent);
+ destroy_hsa_program (agent);
if (pthread_rwlock_unlock (&agent->modules_rwlock))
GOMP_PLUGIN_fatal ("Unable to unlock an HSA agent rwlock");
}
@@ -1027,8 +1110,7 @@ GOMP_OFFLOAD_fini_device (int n)
free (module);
}
agent->first_module = NULL;
- if (agent->prog_finalized)
- destroy_hsa_program (agent);
+ destroy_hsa_program (agent);
hsa_status_t status = hsa_queue_destroy (agent->command_q);
if (status != HSA_STATUS_SUCCESS)
@@ -1037,7 +1119,7 @@ GOMP_OFFLOAD_fini_device (int n)
GOMP_PLUGIN_fatal ("Failed to destroy an HSA agent program mutex");
if (pthread_rwlock_destroy (&agent->modules_rwlock))
GOMP_PLUGIN_fatal ("Failed to destroy an HSA agent rwlock");
- agent->initialized = false;
+ agent->initialized = false;
}
/* Part of the libgomp plugin interface. Not implemented as it is not required
diff --git a/libgomp/target.c b/libgomp/target.c
index 7a60b66..a555c0f 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -945,6 +945,24 @@ gomp_fini_device (struct gomp_device_descr *devicep)
devicep->is_initialized = false;
}
+/* Execute a host function FN with HOSTADDRS. */
+
+static void
+run_on_host (void (*fn) (void *), void **hostaddrs)
+{
+ struct gomp_thread old_thr, *thr = gomp_thread ();
+ old_thr = *thr;
+ memset (thr, '\0', sizeof (*thr));
+ if (gomp_places_list)
+ {
+ thr->place = old_thr.place;
+ thr->ts.place_partition_len = gomp_places_list_len;
+ }
+ fn (hostaddrs);
+ gomp_free_thread (thr);
+ *thr = old_thr;
+}
+
/* Called when encountering a target directive. If DEVICE
is GOMP_DEVICE_ICV, it means use device-var ICV. If it is
GOMP_DEVICE_HOST_FALLBACK (or any value
@@ -966,17 +984,7 @@ GOMP_target (int device, void (*fn) (void *), const void *kernel_launch,
|| !(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400))
{
/* Host fallback. */
- struct gomp_thread old_thr, *thr = gomp_thread ();
- old_thr = *thr;
- memset (thr, '\0', sizeof (*thr));
- if (gomp_places_list)
- {
- thr->place = old_thr.place;
- thr->ts.place_partition_len = gomp_places_list_len;
- }
- fn (hostaddrs);
- gomp_free_thread (thr);
- *thr = old_thr;
+ run_on_host (fn, hostaddrs);
return;
}
@@ -1006,6 +1014,12 @@ GOMP_target (int device, void (*fn) (void *), const void *kernel_launch,
fn_addr = (void *) tgt_fn->tgt_offset;
}
+ if (devicep->can_run_func && !devicep->can_run_func (fn_addr))
+ {
+ run_on_host (fn, hostaddrs);
+ return;
+ }
+
struct target_mem_desc *tgt_vars;
if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
tgt_vars = NULL;
@@ -1020,6 +1034,7 @@ GOMP_target (int device, void (*fn) (void *), const void *kernel_launch,
thr->place = old_thr.place;
thr->ts.place_partition_len = gomp_places_list_len;
}
+
devicep->run_func (devicep->target_id, fn_addr,
tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs,
kernel_launch);
@@ -1160,7 +1175,10 @@ gomp_load_plugin_for_device (struct gomp_device_descr *device,
DLSYM (host2dev);
device->capabilities = device->get_caps_func ();
if (device->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400)
- DLSYM (run);
+ {
+ DLSYM (run);
+ DLSYM (can_run);
+ }
if (device->capabilities & GOMP_OFFLOAD_CAP_OPENACC_200)
{
if (!DLSYM_OPT (openacc.exec, openacc_parallel)
--
2.5.1
>From 35b9c822b935a9d10df2e1b93cc3179d9ca470f4 Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Mon, 5 Oct 2015 16:39:12 +0200
Subject: [PATCH 2/9] HSA: Add loading of default BRIG shared libraries
libgomp/ChangeLog:
2015-10-06 Martin Liska <mliska@suse.cz>
* plugin/plugin-hsa.c (struct brig_library_info): New structure.
(struct agent_info): Add BRIG shared libraries.
(add_shared_library): New function.
(release_agent_shared_libraries): Likewise.
(create_and_finalize_hsa_program): Add loading of shared
libraries.
(GOMP_OFFLOAD_fini_device): Release allocated memory.
gcc/ChangeLog:
2015-10-06 Martin Liska <mliska@suse.cz>
* hsa-gen.c (gen_hsa_insns_for_call): Expand unsupported
builtins to calls that are going to be resolved by BRIG shared
libraries.
---
gcc/hsa-gen.c | 35 ++++----------
libgomp/plugin/plugin-hsa.c | 109 +++++++++++++++++++++++++++++++++++++++++---
2 files changed, 112 insertions(+), 32 deletions(-)
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index fb17a25..3dd7613 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -4348,22 +4348,15 @@ specialop:
if (TREE_CODE (byte_size) != INTEGER_CST)
{
- HSA_SORRY_AT (gimple_location (stmt),
- "support for HSA does not implement __builtin_memcpy "
- "with a non constant size");
+ gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
return;
}
unsigned n = tree_to_uhwi (byte_size);
- /* TODO: fallback to call to memcpy library function. */
if (n > HSA_MEMORY_BUILTINS_LIMIT)
{
- HSA_SORRY_ATV
- (gimple_location (stmt),
- "support for HSA does implement __builtin_memcpy with a size "
- "bigger than %u bytes, %u bytes are requested",
- HSA_MEMORY_BUILTINS_LIMIT, n);
+ gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
return;
}
@@ -4388,10 +4381,7 @@ specialop:
if (TREE_CODE (c) != INTEGER_CST)
{
- HSA_SORRY_AT
- (gimple_location (stmt),
- "support for HSA does not implement __builtin_memset with a "
- "non constant byte value that should be written");
+ gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
return;
}
@@ -4399,22 +4389,15 @@ specialop:
if (TREE_CODE (byte_size) != INTEGER_CST)
{
- HSA_SORRY_AT (gimple_location (stmt),
- "support for HSA does not implement "
- "__builtin_memset with a non constant size");
+ gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
return;
}
unsigned n = tree_to_uhwi (byte_size);
- /* TODO: fallback to call to memset library function. */
if (n > HSA_MEMORY_BUILTINS_LIMIT)
{
- HSA_SORRY_ATV
- (gimple_location (stmt),
- "support for HSA does implement __builtin_memset with a size "
- "bigger than %u bytes, %u bytes are requested",
- HSA_MEMORY_BUILTINS_LIMIT, n);
+ gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
return;
}
@@ -4433,10 +4416,10 @@ specialop:
break;
}
default:
- HSA_SORRY_ATV (gimple_location (stmt),
- "support for HSA does not implement calls to builtin %D",
- gimple_call_fndecl (stmt));
- return;
+ {
+ gen_hsa_insns_for_direct_call (stmt, hbb, ssa_map);
+ return;
+ }
}
}
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index 60dac54..f1c0427 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -7,7 +7,7 @@
#include "hsa.h"
#include "hsa-traits.h"
#include "hsa_ext_finalize.h"
-
+#include "dlfcn.h"
/* Part of the libgomp plugin interface. Return the name of the accelerator,
which is "hsa". */
@@ -70,19 +70,28 @@ init_enviroment_variables (void)
suppress_host_fallback = false;
}
-/* Print a debug message to stderr if DEBUG value is set to true. */
+/* Print a logging message with PREFIX to stderr if HSA_DEBUG value
+ is set to true. */
-#define HSA_DEBUG(...) \
+#define HSA_LOG(prefix, ...) \
do \
{ \
if (debug) \
{ \
- fprintf (stderr, "HSA debug: "); \
+ fprintf (stderr, prefix); \
fprintf (stderr, __VA_ARGS__); \
} \
} \
while (false);
+/* Print a debugging message to stderr. */
+
+#define HSA_DEBUG(...) HSA_LOG ("HSA debug: ", __VA_ARGS__)
+
+/* Print a warning message to stderr. */
+
+#define HSA_WARNING(...) HSA_LOG ("HSA warning: ", __VA_ARGS__)
+
/* Print HSA warning STR with an HSA STATUS code. */
static void
@@ -190,6 +199,14 @@ struct module_info
struct kernel_info kernels[];
};
+/* Information about shared brig library. */
+
+struct brig_library_info
+{
+ char *file_name;
+ hsa_ext_module_t image;
+};
+
/* Description of an HSA GPU agent and the program associated with it. */
struct agent_info
@@ -226,6 +243,10 @@ struct agent_info
bool prog_finalized_error;
/* HSA executable - the finalized program that is used to locate kernels. */
hsa_executable_t executable;
+ /* List of BRIG libraries. */
+ struct brig_library_info **brig_libraries;
+ /* Number of loaded shared BRIG libraries. */
+ unsigned brig_libraries_count;
};
/* Information about the whole HSA environment and all of its agents. */
@@ -557,6 +578,50 @@ GOMP_OFFLOAD_load_image (int ord, unsigned version __attribute__ ((unused)),
return kernel_count;
}
+/* Add a shared BRIG library from a FILE_NAME to an AGENT. */
+
+static struct brig_library_info *
+add_shared_library (const char *file_name, struct agent_info *agent)
+{
+ struct brig_library_info *library = NULL;
+
+ void *f = dlopen (file_name, RTLD_NOW);
+ void *start = dlsym (f, "__brig_start");
+ void *end = dlsym (f, "__brig_end");
+
+ if (start == NULL || end == NULL)
+ return NULL;
+
+ unsigned size = end - start;
+ char *buf = (char *) malloc (size);
+ memcpy (buf, start, size);
+
+ library = GOMP_PLUGIN_malloc (sizeof (struct agent_info));
+ library->file_name = (char *) GOMP_PLUGIN_malloc
+ ((strlen (file_name) + 1) * sizeof (char));
+ strcpy (library->file_name, file_name);
+ library->image = (hsa_ext_module_t) buf;
+
+ return library;
+}
+
+/* Release memory used for BRIG shared libraries that correspond
+ to an AGENT. */
+
+static void
+release_agent_shared_libraries (struct agent_info *agent)
+{
+ for (unsigned i = 0; i < agent->brig_libraries_count; i++)
+ if (agent->brig_libraries[i])
+ {
+ free (agent->brig_libraries[i]->file_name);
+ free (agent->brig_libraries[i]->image);
+ free (agent->brig_libraries[i]);
+ }
+
+ free (agent->brig_libraries);
+}
+
/* Create and finalize the program consisting of all loaded modules. */
static void
@@ -582,13 +647,42 @@ create_and_finalize_hsa_program (struct agent_info *agent)
struct module_info *module = agent->first_module;
while (module)
{
- status = hsa_ext_program_add_module(prog_handle,
- module->image_desc->brig_module);
+ status = hsa_ext_program_add_module (prog_handle,
+ module->image_desc->brig_module);
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Could not add a module to the HSA program", status);
module = module->next;
mi++;
}
+
+ /* Load all shared libraries. */
+ const char *libraries[] = { "libhsamath.so", "libhsastd.so" };
+ const unsigned libraries_count = sizeof (libraries) / sizeof (const char *);
+
+ agent->brig_libraries_count = libraries_count;
+ agent->brig_libraries = GOMP_PLUGIN_malloc_cleared
+ (sizeof (struct brig_library_info) * libraries_count);
+
+ for (unsigned i = 0; i < libraries_count; i++)
+ {
+ struct brig_library_info *library = add_shared_library (libraries[i],
+ agent);
+ if (library == NULL)
+ {
+ HSA_WARNING ("Could not open a shared BRIG library: %s\n",
+ libraries[i]);
+ continue;
+ }
+
+ status = hsa_ext_program_add_module (prog_handle, library->image);
+ if (status != HSA_STATUS_SUCCESS)
+ hsa_warn ("Could not add a shared BRIG library the HSA program",
+ status);
+ else
+ HSA_DEBUG ("a shared BRIG library has been added to a program: %s\n",
+ libraries[i]);
+ }
+
hsa_ext_control_directives_t control_directives;
memset (&control_directives, 0, sizeof (control_directives));
hsa_code_object_t code_object;
@@ -625,6 +719,7 @@ create_and_finalize_hsa_program (struct agent_info *agent)
goto final;
failure:
+ release_agent_shared_libraries (agent);
agent->prog_finalized_error = true;
final:
@@ -1112,6 +1207,8 @@ GOMP_OFFLOAD_fini_device (int n)
agent->first_module = NULL;
destroy_hsa_program (agent);
+ release_agent_shared_libraries (agent);
+
hsa_status_t status = hsa_queue_destroy (agent->command_q);
if (status != HSA_STATUS_SUCCESS)
hsa_fatal ("Error destroying command queue", status);
--
2.5.1
>From 06a75aff1f695f128f49664af027eeadab255caa Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Wed, 7 Oct 2015 16:50:04 +0200
Subject: [PATCH 3/9] HSA: change linkage for a function that cannot be
emitted.
gcc/ChangeLog:
2015-10-07 Martin Liska <mliska@suse.cz>
* hsa-brig.c (emit_function_declaration): Return pointer to
BrigDirectiveExecutable structure.
(hsa_brig_emit_function): Save the structure for all emitted
declarations.
(hsa_output_brig): Fix function linkage, if needed.
* hsa-gen.c (HSA_SORRY_AT{V}): Put all failed functions
to a set.
* hsa.h (hsa_failed_functions): New global variable.
(hsa_fail_cfun): New function.
* hsa.c (hsa_failed_functions): Define the variable.
(hsa_fail_cfun): New function.
---
gcc/hsa-brig.c | 30 +++++++++++++++++++++++-------
gcc/hsa-gen.c | 4 ++--
gcc/hsa.c | 17 +++++++++++++++++
gcc/hsa.h | 3 +++
4 files changed, 45 insertions(+), 9 deletions(-)
diff --git a/gcc/hsa-brig.c b/gcc/hsa-brig.c
index 53c9d32..29ddbba 100644
--- a/gcc/hsa-brig.c
+++ b/gcc/hsa-brig.c
@@ -132,7 +132,7 @@ static bool brig_initialized = false;
static hash_map<tree, BrigCodeOffset32_t> *function_offsets;
/* Set of emitted function declarations. */
-static hash_set <tree> *emitted_declarations;
+static hash_map <tree, BrigDirectiveExecutable *> *emitted_declarations;
/* List of sbr instructions. */
static vec <hsa_insn_sbr *> *switch_instructions;
@@ -1110,15 +1110,18 @@ emit_queued_operands (void)
/* Emit directives describing the function that is used for
a function declaration. */
-static void
+
+static BrigDirectiveExecutable *
emit_function_declaration (tree decl)
{
hsa_function_representation *f = hsa_generate_function_declaration (decl);
- emit_function_directives (f, true);
+ BrigDirectiveExecutable *e = emit_function_directives (f, true);
emit_queued_operands ();
delete f;
+
+ return e;
}
/* Emit an HSA memory instruction and all necessary directives, schedule
@@ -1881,17 +1884,17 @@ hsa_brig_emit_function (void)
function_offsets = new hash_map<tree, BrigCodeOffset32_t> ();
if (!emitted_declarations)
- emitted_declarations = new hash_set<tree> ();
+ emitted_declarations = new hash_map <tree, BrigDirectiveExecutable *> ();
for (unsigned i = 0; i < hsa_cfun->called_functions.length (); i++)
{
tree called = hsa_cfun->called_functions[i];
/* If the function has no definition, emit a declaration. */
- if (!emitted_declarations->contains (called))
+ if (!emitted_declarations->get (called))
{
- emit_function_declaration (called);
- emitted_declarations->add (called);
+ BrigDirectiveExecutable *e = emit_function_declaration (called);
+ emitted_declarations->put (called, e);
}
}
@@ -2280,6 +2283,19 @@ hsa_output_brig (void)
code_ref->ref = htole32 (*func_offset);
}
+ /* Iterate all function declarations and if we meet a function that should
+ have module linkage and we are unable to emit HSAIL for the function,
+ then change the linkage to program linkage. Doing so, we will emit
+ a valid BRIG image. */
+ if (hsa_failed_functions != NULL && emitted_declarations != NULL)
+ for (hash_map <tree, BrigDirectiveExecutable *>::iterator it =
+ emitted_declarations->begin (); it != emitted_declarations->end ();
+ ++it)
+ {
+ if (hsa_failed_functions->contains ((*it).first))
+ (*it).second->linkage = BRIG_LINKAGE_PROGRAM;
+ }
+
saved_section = in_section;
switch_to_section (get_section (BRIG_ELF_SECTION_NAME, SECTION_NOTYPE, NULL));
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 3dd7613..ae03361 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -84,7 +84,7 @@ along with GCC; see the file COPYING3. If not see
#define HSA_SORRY_ATV(location, message, ...) \
do \
{ \
- hsa_cfun->seen_error = true; \
+ hsa_fail_cfun (); \
if (warning_at (EXPR_LOCATION (hsa_cfun->decl), OPT_Whsa, \
HSA_SORRY_MSG)) \
inform (location, message, ##__VA_ARGS__); \
@@ -96,7 +96,7 @@ along with GCC; see the file COPYING3. If not see
#define HSA_SORRY_AT(location, message) \
do \
{ \
- hsa_cfun->seen_error = true; \
+ hsa_fail_cfun (); \
if (warning_at (EXPR_LOCATION (hsa_cfun->decl), OPT_Whsa, \
HSA_SORRY_MSG)) \
inform (location, message); \
diff --git a/gcc/hsa.c b/gcc/hsa.c
index c938248..5aa40fc 100644
--- a/gcc/hsa.c
+++ b/gcc/hsa.c
@@ -107,6 +107,9 @@ hsa_summary_t *hsa_summaries = NULL;
/* HSA number of threads. */
hsa_symbol *hsa_num_threads = NULL;
+/* HSA function that cannot be expanded to HSAIL. */
+hash_set <tree> *hsa_failed_functions = NULL;
+
/* True if compilation unit-wide data are already allocated and initialized. */
static bool compilation_unit_data_initialized;
@@ -140,6 +143,9 @@ hsa_deinit_compilation_unit_data (void)
{
if (compilation_unit_data_initialized)
delete hsa_global_variable_symbols;
+
+ if (hsa_failed_functions)
+ delete hsa_failed_functions;
}
/* Return true if we are generating large HSA machine model. */
@@ -653,4 +659,15 @@ hsa_seen_error (void)
return hsa_cfun->seen_error;
}
+/* Mark current HSA function as failed. */
+
+void
+hsa_fail_cfun (void)
+{
+ if (hsa_failed_functions == NULL)
+ hsa_failed_functions = new hash_set <tree> ();
+ hsa_failed_functions->add (hsa_cfun->decl);
+ hsa_cfun->seen_error = true;
+}
+
#include "gt-hsa.h"
diff --git a/gcc/hsa.h b/gcc/hsa.h
index b400b39..4861e00 100644
--- a/gcc/hsa.h
+++ b/gcc/hsa.h
@@ -1043,6 +1043,8 @@ extern hash_map <tree, vec <char *> *> *hsa_decl_kernel_dependencies;
extern hsa_summary_t *hsa_summaries;
extern hsa_symbol *hsa_num_threads;
extern unsigned hsa_kernel_calls_counter;
+extern hash_set <tree> *hsa_failed_functions;
+
bool hsa_callable_function_p (tree fndecl);
void hsa_init_compilation_unit_data (void);
void hsa_deinit_compilation_unit_data (void);
@@ -1068,6 +1070,7 @@ const char *hsa_get_declaration_name (tree decl);
void hsa_register_kernel (cgraph_node *host);
void hsa_register_kernel (cgraph_node *gpu, cgraph_node *host);
bool hsa_seen_error (void);
+void hsa_fail_cfun (void);
/* In hsa-gen.c. */
void hsa_build_append_simple_mov (hsa_op_reg *, hsa_op_base *, hsa_bb *);
--
2.5.1
>From 9ae73202db1fd631d7ac4b56263cf5c91b256be6 Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Wed, 7 Oct 2015 17:07:32 +0200
Subject: [PATCH 4/9] HSA: surrender HSA emission if an error is seen.
gcc/ChangeLog:
2015-10-07 Martin Liska <mliska@suse.cz>
* hsa-gen.c (gen_hsa_insns_for_operation_assignment): Surrender
if an error is seen.
---
gcc/hsa-gen.c | 3 +++
1 file changed, 3 insertions(+)
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index ae03361..f831611 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -2910,6 +2910,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb,
hsa_op_with_type *op2 = rhs2 != NULL_TREE ?
hsa_reg_or_immed_for_gimple_op (rhs2, hbb, ssa_map) : NULL;
+ if (hsa_seen_error ())
+ return;
+
switch (rhs_class)
{
case GIMPLE_TERNARY_RHS:
--
2.5.1
>From d6cfc72c388553316619d9d51b25c24539ca34d3 Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Wed, 7 Oct 2015 22:14:19 +0200
Subject: [PATCH 5/9] HSA: preserve HSA function names and visibility
gcc/ChangeLog:
2015-10-08 Martin Liska <mliska@suse.cz>
* hsa-brig.c (hsa_output_kernel_mapping): Change type of kernel
dependencies container.
* hsa-gen.c (has_host_function_p): New function.
(hsa_generate_function_declaration): Emit host implementation of
a function declaration.
(generate_hsa): Likewise.
(hsa_get_gpu_function): Remove unused function.
* hsa.c (hsa_add_kernel_dependency): Change type of a function
argument.
* hsa.h: Likewise.
* ipa-hsa.c (process_hsa_functions): Copy visibility from
a cloned function.
---
gcc/hsa-brig.c | 7 ++++---
gcc/hsa-gen.c | 39 +++++++++++++++++++--------------------
gcc/hsa.c | 12 ++++++------
gcc/hsa.h | 4 ++--
gcc/ipa-hsa.c | 2 ++
5 files changed, 33 insertions(+), 31 deletions(-)
diff --git a/gcc/hsa-brig.c b/gcc/hsa-brig.c
index 29ddbba..b6eabfd 100644
--- a/gcc/hsa-brig.c
+++ b/gcc/hsa-brig.c
@@ -2057,10 +2057,11 @@ hsa_output_kernel_mapping (tree brig_decl)
vec<constructor_elt, va_gc> *kernel_dependencies_vec = NULL;
if (hsa_decl_kernel_dependencies)
{
- vec<char *> **slot = hsa_decl_kernel_dependencies->get (kernel);
+ vec<const char *> **slot;
+ slot = hsa_decl_kernel_dependencies->get (kernel);
if (slot)
{
- vec <char *> *dependencies = *slot;
+ vec <const char *> *dependencies = *slot;
count = dependencies->length ();
kernel_dependencies_vector_type = build_array_type
@@ -2069,7 +2070,7 @@ hsa_output_kernel_mapping (tree brig_decl)
for (unsigned j = 0; j < count; j++)
{
- char *d = (*dependencies)[j];
+ const char *d = (*dependencies)[j];
len = strlen (d);
tree dependency_name = build_string (len, d);
TREE_TYPE (dependency_name) = build_array_type
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index f831611..cf36882 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -713,19 +713,6 @@ get_symbol_for_decl (tree decl)
return sym;
}
-/* For a given function declaration, return a GPU function
- of the function. */
-
-static tree
-hsa_get_gpu_function (tree decl)
-{
- hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (decl));
- gcc_assert (s->kind != HSA_NONE);
- gcc_assert (!s->gpu_implementation_p);
-
- return s->binded_function->decl;
-}
-
/* For a given HSA function declaration, return a host
function declaration. */
@@ -739,6 +726,15 @@ hsa_get_host_function (tree decl)
return s->binded_function->decl;
}
+/* Return true if function DECL has a host equivalent function. */
+
+static bool
+has_host_function_p (tree decl)
+{
+ hsa_function_summary *s = hsa_summaries->get (cgraph_node::get_create (decl));
+ return s != NULL && s->kind != HSA_NONE && s->gpu_implementation_p;
+}
+
/* Create a spill symbol of type TYPE. */
hsa_symbol *
@@ -4322,10 +4318,9 @@ specialop:
called = TREE_OPERAND (called, 0);
gcc_checking_assert (TREE_CODE (called) == FUNCTION_DECL);
- const char *name = hsa_get_declaration_name
- (hsa_get_gpu_function (called));
- hsa_add_kernel_dependency (hsa_cfun->decl,
- hsa_brig_function_name (name));
+ char *name = xstrdup (hsa_get_declaration_name (called));
+ hsa_add_kernel_dependency
+ (hsa_cfun->decl, hsa_brig_function_name (name));
gen_hsa_insns_for_kernel_call (hbb, as_a <gcall *> (stmt));
break;
@@ -4832,9 +4827,12 @@ hsa_generate_function_declaration (tree decl)
{
hsa_function_representation *fun = XCNEW (hsa_function_representation);
+ tree host_decl = has_host_function_p (decl) ? hsa_get_host_function (decl) :
+ decl;
+
fun->declaration_p = true;
fun->decl = decl;
- fun->name = xstrdup (hsa_get_declaration_name (decl));
+ fun->name = xstrdup (hsa_get_declaration_name (host_decl));
hsa_sanitize_name (fun->name);
gen_function_decl_parameters (fun, decl);
@@ -5118,6 +5116,7 @@ static void
generate_hsa (bool kernel)
{
vec <hsa_op_reg_p> ssa_map = vNULL;
+ tree host_decl = NULL_TREE;
if (hsa_num_threads == NULL)
emit_hsa_module_variables ();
@@ -5130,9 +5129,9 @@ generate_hsa (bool kernel)
hsa_cfun->decl = cfun->decl;
hsa_cfun->kern_p = kernel;
+ host_decl = hsa_get_host_function (current_function_decl);
ssa_map.safe_grow_cleared (SSANAMES (cfun)->length ());
- hsa_cfun->name
- = xstrdup (hsa_get_declaration_name (current_function_decl));
+ hsa_cfun->name = xstrdup (hsa_get_declaration_name (host_decl));
hsa_sanitize_name (hsa_cfun->name);
gen_function_def_parameters (hsa_cfun, &ssa_map);
diff --git a/gcc/hsa.c b/gcc/hsa.c
index 5aa40fc..7da7b6c 100644
--- a/gcc/hsa.c
+++ b/gcc/hsa.c
@@ -96,7 +96,7 @@ static GTY (()) vec<hsa_decl_kernel_map_element, va_gc> *hsa_decl_kernel_mapping
/* Mapping between decls and corresponding HSA kernels
called by the function. */
-hash_map <tree, vec <char *> *> *hsa_decl_kernel_dependencies;
+hash_map <tree, vec <const char *> *> *hsa_decl_kernel_dependencies;
/* Hash function to lookup a symbol for a decl. */
hash_table <hsa_free_symbol_hasher> *hsa_global_variable_symbols;
@@ -564,16 +564,16 @@ hsa_free_decl_kernel_mapping (void)
/* Add new kernel dependency. */
void
-hsa_add_kernel_dependency (tree caller, char *called_function)
+hsa_add_kernel_dependency (tree caller, const char *called_function)
{
if (hsa_decl_kernel_dependencies == NULL)
- hsa_decl_kernel_dependencies = new hash_map<tree, vec<char *> *> ();
+ hsa_decl_kernel_dependencies = new hash_map<tree, vec<const char *> *> ();
- vec <char *> *s = NULL;
- vec <char *> **slot = hsa_decl_kernel_dependencies->get (caller);
+ vec <const char *> *s = NULL;
+ vec <const char *> **slot = hsa_decl_kernel_dependencies->get (caller);
if (slot == NULL)
{
- s = new vec <char *> ();
+ s = new vec <const char *> ();
hsa_decl_kernel_dependencies->put (caller, s);
}
else
diff --git a/gcc/hsa.h b/gcc/hsa.h
index 4861e00..98d70e0 100644
--- a/gcc/hsa.h
+++ b/gcc/hsa.h
@@ -1039,7 +1039,7 @@ hsa_summary_t::link_functions (cgraph_node *gpu, cgraph_node *host,
/* in hsa.c */
extern struct hsa_function_representation *hsa_cfun;
extern hash_table <hsa_free_symbol_hasher> *hsa_global_variable_symbols;
-extern hash_map <tree, vec <char *> *> *hsa_decl_kernel_dependencies;
+extern hash_map <tree, vec <const char *> *> *hsa_decl_kernel_dependencies;
extern hsa_summary_t *hsa_summaries;
extern hsa_symbol *hsa_num_threads;
extern unsigned hsa_kernel_calls_counter;
@@ -1063,7 +1063,7 @@ tree hsa_get_decl_kernel_mapping_decl (unsigned i);
char *hsa_get_decl_kernel_mapping_name (unsigned i);
unsigned hsa_get_decl_kernel_mapping_omp_size (unsigned i);
void hsa_free_decl_kernel_mapping (void);
-void hsa_add_kernel_dependency (tree caller, char *called_function);
+void hsa_add_kernel_dependency (tree caller, const char *called_function);
void hsa_sanitize_name (char *p);
char *hsa_brig_function_name (const char *p);
const char *hsa_get_declaration_name (tree decl);
diff --git a/gcc/ipa-hsa.c b/gcc/ipa-hsa.c
index d87a17b..0c072da 100644
--- a/gcc/ipa-hsa.c
+++ b/gcc/ipa-hsa.c
@@ -101,6 +101,7 @@ process_hsa_functions (void)
{
cgraph_node *clone = node->create_virtual_clone
(vec <cgraph_edge *> (), NULL, NULL, "hsa");
+ TREE_PUBLIC (clone->decl) = TREE_PUBLIC (node->decl);
clone->force_output = true;
hsa_summaries->link_functions (clone, node, s->kind);
@@ -114,6 +115,7 @@ process_hsa_functions (void)
{
cgraph_node *clone = node->create_virtual_clone
(vec <cgraph_edge *> (), NULL, NULL, "hsa");
+ TREE_PUBLIC (clone->decl) = TREE_PUBLIC (node->decl);
if (!cgraph_local_p (node))
clone->force_output = true;
--
2.5.1
>From 55b82750b576588883d41407bf96e3cd1adc4c62 Mon Sep 17 00:00:00 2001
From: marxin <mliska@suse.cz>
Date: Thu, 8 Oct 2015 10:06:25 +0200
Subject: [PATCH 6/9] HSA: simplify LTO partitioning and improve kernel
dependencies resolution.
gcc/lto/ChangeLog:
2015-10-08 Martin Liska <mliska@suse.cz>
* lto-partition.c (add_symbol_to_partition_1): Simplify a much
partitioning of HSA symbols.
libgomp/ChangeLog:
2015-10-08 Martin Liska <mliska@suse.cz>
* plugin/plugin-hsa.c (get_kernel_for_agent): New function.
(init_single_kernel): Use it.
(create_kernel_dispatch_recursive): Dtto.
---
gcc/lto/lto-partition.c | 30 ------------------------------
libgomp/plugin/plugin-hsa.c | 30 ++++++++++++++++--------------
2 files changed, 16 insertions(+), 44 deletions(-)
diff --git a/gcc/lto/lto-partition.c b/gcc/lto/lto-partition.c
index 01a60b2..9e0b95d 100644
--- a/gcc/lto/lto-partition.c
+++ b/gcc/lto/lto-partition.c
@@ -196,36 +196,6 @@ add_symbol_to_partition_1 (ltrans_partition part, symtab_node *node)
"adding an HSA function (host/gpu) to the "
"partition: %s\n",
s->binded_function->name ());
-
- ipa_ref *ref;
-
- /* Add all parents nodes that have HSA type. */
- for (unsigned i = 0; node->iterate_referring (i, ref); i++)
- {
- cgraph_node *r = dyn_cast <cgraph_node *> (ref->referring);
- if (r && hsa_summaries->get (r)->kind != HSA_NONE)
- {
- add_symbol_to_partition_1 (part, r);
- if (symtab->dump_file)
- fprintf (symtab->dump_file,
- "adding an HSA referring node: %s\n",
- r->name ());
- }
- }
-
- /* Add all children nodes that have HSA type. */
- for (unsigned i = 0; node->iterate_reference (i, ref); i++)
- {
- cgraph_node *r = dyn_cast <cgraph_node *> (ref->referred);
- if (r && hsa_summaries->get (r)->kind != HSA_NONE)
- {
- add_symbol_to_partition_1 (part, r);
- if (symtab->dump_file)
- fprintf (symtab->dump_file,
- "adding an HSA referred symbol: %s\n",
- r->name ());
- }
- }
}
}
}
diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c
index f1c0427..ed3573f 100644
--- a/libgomp/plugin/plugin-hsa.c
+++ b/libgomp/plugin/plugin-hsa.c
@@ -265,14 +265,21 @@ struct hsa_context_info
static struct hsa_context_info hsa_context;
-/* Find kernel in MODULE by name provided in KERNEL_NAME. */
+/* Find kernel for an AGENT by name provided in KERNEL_NAME. */
static struct kernel_info *
-get_kernel_in_module (struct module_info *module, const char *kernel_name)
+get_kernel_for_agent (struct agent_info *agent, const char *kernel_name)
{
- for (unsigned i = 0; i < module->kernel_count; i++)
- if (strcmp (module->kernels[i].name, kernel_name) == 0)
- return &module->kernels[i];
+ struct module_info *module = agent->first_module;
+
+ while (module)
+ {
+ for (unsigned i = 0; i < module->kernel_count; i++)
+ if (strcmp (module->kernels[i].name, kernel_name) == 0)
+ return &module->kernels[i];
+
+ module = module->next;
+ }
return NULL;
}
@@ -841,12 +848,10 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size)
if (kernel->omp_data_size > *max_omp_data_size)
*max_omp_data_size = kernel->omp_data_size;
- /* FIXME: do not consider all kernels to live in a same module. */
- struct module_info *module = kernel->agent->first_module;
for (unsigned i = 0; i < kernel->dependencies_count; i++)
{
- struct kernel_info *dependency = get_kernel_in_module
- (module, kernel->dependencies[i]);
+ struct kernel_info *dependency = get_kernel_for_agent
+ (agent, kernel->dependencies[i]);
if (dependency == NULL)
{
@@ -917,9 +922,6 @@ static struct hsa_kernel_dispatch *
create_kernel_dispatch_recursive (struct kernel_info *kernel,
unsigned omp_data_size)
{
- // TODO: find correct module
- struct module_info *module = kernel->agent->first_module;
-
struct hsa_kernel_dispatch *shadow = create_kernel_dispatch (kernel,
omp_data_size);
shadow->omp_num_threads = 64;
@@ -927,8 +929,8 @@ create_kernel_dispatch_recursive (struct kernel_info *kernel,
for (unsigned i = 0; i < kernel->dependencies_count; i++)
{
- struct kernel_info *dependency = get_kernel_in_module
- (module, kernel->dependencies[i]);
+ struct kernel_info *dependency = get_kernel_for_agent
+ (kernel->agent, kernel->dependencies[i]);
shadow->children_dispatches[i] = create_kernel_dispatch_recursive
(dependency, omp_data_size);
}
--
2.5.1