Hi,
I have found out that the hsa branch ICEs when expanding the following
tot HSAIL:
int foo ()
{
#pragma omp target
{
int q[8];
__builtin_memset (&q[2], 0, sizeof (int) * 6);
}
}
The problem was that gen_hsa_addr was used for both values and memory
reference trees, which are really separate things, so I introduced a
special function for the former. Consequently, I had to fix a bit
bit-field access signaling in gen_hsa_addr.
Committed to the branch after the standard testing I do.
Martin
2015-09-10 Martin Jambor <mjam...@suse.cz>
* hsa-gen.c (get_address_from_value): New function.
(gen_hsa_ternary_atomic_for_builtin): Use it instead of
gen_hsa_addr.
(gen_hsa_insns_for_call): Likewise.
(gen_hsa_addr): Set bitsize to zero if the thing is not a bitfield
access.
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 7796895..6e39c78 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -1546,18 +1546,9 @@ gen_hsa_addr (tree ref, hsa_bb *hbb, vec <hsa_op_reg_p>
*ssa_map,
switch (TREE_CODE (ref))
{
case SSA_NAME:
- /* The SSA_NAME and ADDR_EXPR cases cannot occur in a valid gimple memory
- reference but we also use this function to generate addresses of
- instructions representing operands of atomic memory access builtins
- which are just addresses and not references. */
- gcc_assert (!reg);
- reg = hsa_reg_for_gimple_ssa_reqtype (ref, ssa_map, hbb, addrtype);
- break;
-
case ADDR_EXPR:
- ref = TREE_OPERAND (ref, 0);
- gcc_assert (DECL_P (ref));
- /* Fall-through. */
+ gcc_unreachable ();
+
case PARM_DECL:
case VAR_DECL:
case RESULT_DECL:
@@ -1636,6 +1627,12 @@ out:
/* Calculate remaining bitsize offset (if presented). */
bitpos %= BITS_PER_UNIT;
+ /* If bitsize is a power of two that is greater or equal to BITS_PER_UNIT, it
+ is not a reason to think this is a bit-field access. */
+ if (bitpos == 0
+ && (bitsize >= BITS_PER_UNIT)
+ && !(bitsize & (bitsize - 1)))
+ bitsize = 0;
if ((bitpos || bitsize) && (output_bitpos == NULL || output_bitsize == NULL))
sorry ("Support for HSA does not implement unhandled bit field reference "
@@ -3297,6 +3294,37 @@ gen_hsa_unaryop_for_builtin (int opcode, gimple stmt,
hsa_bb *hbb,
gen_hsa_unary_operation (opcode, dest, op, hbb);
}
+/* Generate HSA address corresponding to a value VAL (as opposed to a memory
+ reference tree), for example an SSA_NAME or an ADDR_EXPR. HBB is the HSA BB
+ to which the instruction should be added and SSA_MAP is used to map gimple
+ SSA names to HSA pseudoregisters. */
+
+static hsa_op_address *
+get_address_from_value (tree val, hsa_bb *hbb, vec <hsa_op_reg_p> *ssa_map)
+{
+ switch (TREE_CODE (val))
+ {
+ case SSA_NAME:
+ {
+ BrigType16_t addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
+ hsa_op_reg *reg = hsa_reg_for_gimple_ssa_reqtype (val, ssa_map,
+ hbb, addrtype);
+ return new hsa_op_address (NULL, reg, 0);
+ }
+ case ADDR_EXPR:
+ return gen_hsa_addr (TREE_OPERAND (val, 0), hbb, ssa_map);
+
+ case INTEGER_CST:
+ if (tree_fits_shwi_p (val))
+ return new hsa_op_address (NULL, NULL, tree_to_shwi (val));
+ /* Otherwise fall-through */
+
+ default:
+ sorry ("Support for HSA does not implement memory access to %E", val);
+ return new hsa_op_address (NULL, NULL, 0);
+ }
+}
+
/* Helper function to create an HSA atomic binary operation instruction out of
calls to atomic builtins. RET_ORIG is true if the built-in is the variant
that return s the value before applying operation, and false if it should
@@ -3337,7 +3365,7 @@ gen_hsa_ternary_atomic_for_builtin (bool ret_orig,
hsa_insn_atomic *atominsn = new hsa_insn_atomic (nops, opcode, acode,
bit_type);
hsa_op_address *addr;
- addr = gen_hsa_addr (gimple_call_arg (stmt, 0), hbb, ssa_map);
+ addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb, ssa_map);
hsa_op_base *op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1),
hbb, ssa_map);
@@ -3520,8 +3548,8 @@ specialop:
{
BrigType16_t mtype = mem_type_for_type (hsa_type_for_scalar_tree_type
(TREE_TYPE (lhs), false));
- hsa_op_address *addr = gen_hsa_addr (gimple_call_arg (stmt, 0),
- hbb, ssa_map);
+ hsa_op_address *addr;
+ addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb, ssa_map);
dest = hsa_reg_for_gimple_ssa (lhs, ssa_map);
hsa_insn_atomic *atominsn
= new hsa_insn_atomic (2, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_LD, mtype,
@@ -3657,7 +3685,7 @@ specialop:
hsa_insn_atomic *atominsn = new hsa_insn_atomic
(4, BRIG_OPCODE_ATOMIC, BRIG_ATOMIC_CAS, atype);
hsa_op_address *addr;
- addr = gen_hsa_addr (gimple_call_arg (stmt, 0), hbb, ssa_map);
+ addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb, ssa_map);
if (lhs != NULL)
dest = hsa_reg_for_gimple_ssa (lhs, ssa_map);
@@ -3709,8 +3737,8 @@ specialop:
tree dst = gimple_call_arg (stmt, 0);
tree src = gimple_call_arg (stmt, 1);
- hsa_op_address *dst_addr = gen_hsa_addr (dst, hbb, ssa_map);
- hsa_op_address *src_addr = gen_hsa_addr (src, hbb, ssa_map);
+ hsa_op_address *dst_addr = get_address_from_value (dst, hbb, ssa_map);
+ hsa_op_address *src_addr = get_address_from_value (src, hbb, ssa_map);
unsigned n = tree_to_uhwi (byte_size);
gen_hsa_memory_copy (hbb, dst_addr, src_addr, n);
@@ -3737,8 +3765,9 @@ specialop:
return;
}
- hsa_op_address *dst_addr = gen_hsa_addr (gimple_call_arg (stmt, 0),
- hbb, ssa_map);
+ hsa_op_address *dst_addr;
+ dst_addr = get_address_from_value (gimple_call_arg (stmt, 0), hbb,
+ ssa_map);
unsigned n = tree_to_uhwi (byte_size);
unsigned HOST_WIDE_INT constant = tree_to_uhwi
(fold_convert (unsigned_char_type_node, c));
