This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[HSA, PR 82416] Do not extend operands to at least 32 bits
- From: Martin Jambor <mjambor at suse dot cz>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>
- Date: Mon, 9 Oct 2017 11:49:56 +0200
- Subject: [HSA, PR 82416] Do not extend operands to at least 32 bits
- Authentication-results: sourceware.org; auth=none
Hi,
Pekka came up with a nice testcase demonstrating that passing true as
min32int to hsa_type_for_scalar_tree_type in reg_for_gimple_ssa was
just wrong. So this patch changes that to false and adds all the
necessary conversions when dealing with instructions that operate on
32bit data or larger.
Tested on an HSA APU, tested by Pekka, bootstrapped on an x86_64-linux
with HSA generation enabled, committed to trunk a few moments ago.
Thanks,
Martin
2017-10-09 Martin Jambor <mjambor@suse.cz>
PR hsa/82416
gcc/
* hsa-common.h (hsa_op_with_type): New method extend_int_to_32bit.
* hsa-gen.c (hsa_extend_inttype_to_32bit): New function.
(hsa_type_for_scalar_tree_type): Use it. Always force min32int for
COMPLEX types.
(hsa_fixup_mov_insn_type): New function.
(hsa_op_with_type::get_in_type): Use it.
(hsa_build_append_simple_mov): Likewise. Allow sub-32bit
immediates in an assert.
(hsa_op_with_type::extend_int_to_32bit): New method.
(gen_hsa_insns_for_bitfield): Fixup instruction and intermediary
types. Convert to dest type if necessary.
(gen_hsa_insns_for_bitfield_load): Fixup load type if necessary.
(reg_for_gimple_ssa): Pass false as min32int to
hsa_type_for_scalar_tree_type.
(gen_hsa_addr): Fixup type when creating addresable temporary.
(gen_hsa_cmp_insn_from_gimple): Extend operands if necessary.
(gen_hsa_unary_operation): Extend operands and convert to dest type if
necessary. Call hsa_fixup_mov_insn_type.
(gen_hsa_binary_operation): Changed operand types to hsa_op_with_type,
extend operands and convert to dest type if necessary.
(gen_hsa_insns_for_operation_assignment): Extend operands and convert
to dest type if necessary.
(set_output_in_type): Call hsa_fixup_mov_insn_type. Just ude dest
if conversion nt necessary and size matches.
(gen_hsa_insns_for_load): Call hsa_fixup_mov_insn_type, convert
to dest type if necessary.
(gen_hsa_insns_for_store): Call hsa_fixup_mov_insn_type.
(gen_hsa_insns_for_switch_stmt): Likewise. Also extend operands if
necessary.
(gen_hsa_clrsb): Likewise.
(gen_hsa_ffs): Likewise.
(gen_hsa_divmod): Extend operands and convert to dest type if
necessary.
(gen_hsa_atomic_for_builtin): Change type of op to hsa_op_with_type.
libgomp/
* testsuite/libgomp.hsa.c/pr82416.c: New test.
---
gcc/hsa-common.h | 3 +
gcc/hsa-gen.c | 218 +++++++++++++++++++++---------
libgomp/testsuite/libgomp.hsa.c/pr82416.c | 37 +++++
3 files changed, 197 insertions(+), 61 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.hsa.c/pr82416.c
diff --git a/gcc/hsa-common.h b/gcc/hsa-common.h
index 810624e4e1c..3075163a020 100644
--- a/gcc/hsa-common.h
+++ b/gcc/hsa-common.h
@@ -157,6 +157,9 @@ public:
/* Convert an operand to a destination type DTYPE and attach insns
to HBB if needed. */
hsa_op_with_type *get_in_type (BrigType16_t dtype, hsa_bb *hbb);
+ /* If this operand has integer type smaller than 32 bits, extend it to 32
+ bits, adding instructions to HBB if needed. */
+ hsa_op_with_type *extend_int_to_32bit (hsa_bb *hbb);
protected:
hsa_op_with_type (BrigKind16_t k, BrigType16_t t);
diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c
index 6e054c0ce82..b5a8c73731a 100644
--- a/gcc/hsa-gen.c
+++ b/gcc/hsa-gen.c
@@ -564,6 +564,19 @@ get_integer_type_by_bytes (unsigned size, bool sign)
return 0;
}
+/* If T points to an integral type smaller than 32 bits, change it to a 32bit
+ equivalent and return the result. Otherwise just return the result. */
+
+static BrigType16_t
+hsa_extend_inttype_to_32bit (BrigType16_t t)
+{
+ if (t == BRIG_TYPE_U8 || t == BRIG_TYPE_U16)
+ return BRIG_TYPE_U32;
+ else if (t == BRIG_TYPE_S8 || t == BRIG_TYPE_S16)
+ return BRIG_TYPE_S32;
+ return t;
+}
+
/* Return HSA type for tree TYPE, which has to fit into BrigType16_t. Pointers
are assumed to use flat addressing. If min32int is true, always expand
integer types to one that has at least 32 bits. */
@@ -580,8 +593,13 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
if (POINTER_TYPE_P (type))
return hsa_get_segment_addr_type (BRIG_SEGMENT_FLAT);
- if (TREE_CODE (type) == VECTOR_TYPE || TREE_CODE (type) == COMPLEX_TYPE)
+ if (TREE_CODE (type) == VECTOR_TYPE)
base = TREE_TYPE (type);
+ else if (TREE_CODE (type) == COMPLEX_TYPE)
+ {
+ base = TREE_TYPE (type);
+ min32int = true;
+ }
else
base = type;
@@ -652,14 +670,9 @@ hsa_type_for_scalar_tree_type (const_tree type, bool min32int)
}
if (min32int)
- {
- /* Registers/immediate operands can only be 32bit or more except for
- f16. */
- if (res == BRIG_TYPE_U8 || res == BRIG_TYPE_U16)
- res = BRIG_TYPE_U32;
- else if (res == BRIG_TYPE_S8 || res == BRIG_TYPE_S16)
- res = BRIG_TYPE_S32;
- }
+ /* Registers/immediate operands can only be 32bit or more except for
+ f16. */
+ res = hsa_extend_inttype_to_32bit (res);
if (TREE_CODE (type) == COMPLEX_TYPE)
{
@@ -1009,6 +1022,16 @@ hsa_get_string_cst_symbol (tree string_cst)
return sym;
}
+/* Make the type of a MOV instruction larger if mandated by HSAIL rules. */
+
+static void
+hsa_fixup_mov_insn_type (hsa_insn_basic *insn)
+{
+ insn->m_type = hsa_extend_inttype_to_32bit (insn->m_type);
+ if (insn->m_type == BRIG_TYPE_B8 || insn->m_type == BRIG_TYPE_B16)
+ insn->m_type = BRIG_TYPE_B32;
+}
+
/* Constructor of the ancestor of all operands. K is BRIG kind that identified
what the operator is. */
@@ -1050,9 +1073,11 @@ hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
else
{
dest = new hsa_op_reg (m_type);
- hbb->append_insn (new hsa_insn_basic (2, BRIG_OPCODE_MOV,
- dest->m_type, dest, this));
+ hsa_insn_basic *mov = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
+ dest->m_type, dest, this);
+ hsa_fixup_mov_insn_type (mov);
+ hbb->append_insn (mov);
/* We cannot simply for instance: 'mov_u32 $_3, 48 (s32)' because
type of the operand must be same as type of the instruction. */
dest->m_type = dtype;
@@ -1061,6 +1086,20 @@ hsa_op_with_type::get_in_type (BrigType16_t dtype, hsa_bb *hbb)
return dest;
}
+/* If this operand has integer type smaller than 32 bits, extend it to 32 bits,
+ adding instructions to HBB if needed. */
+
+hsa_op_with_type *
+hsa_op_with_type::extend_int_to_32bit (hsa_bb *hbb)
+{
+ if (m_type == BRIG_TYPE_U8 || m_type == BRIG_TYPE_U16)
+ return get_in_type (BRIG_TYPE_U32, hbb);
+ else if (m_type == BRIG_TYPE_S8 || m_type == BRIG_TYPE_S16)
+ return get_in_type (BRIG_TYPE_S32, hbb);
+ else
+ return this;
+}
+
/* Constructor of class representing HSA immediate values. TREE_VAL is the
tree representation of the immediate value. If min32int is true,
always expand integer types to one that has at least 32 bits. */
@@ -1292,7 +1331,7 @@ hsa_function_representation::reg_for_gimple_ssa (tree ssa)
return m_ssa_map[SSA_NAME_VERSION (ssa)];
hreg = new hsa_op_reg (hsa_type_for_scalar_tree_type (TREE_TYPE (ssa),
- true));
+ false));
hreg->m_gimple_ssa = ssa;
m_ssa_map[SSA_NAME_VERSION (ssa)] = hreg;
@@ -1799,7 +1838,7 @@ gen_address_calculation (tree exp, hsa_bb *hbb, BrigType16_t addrtype)
case INTEGER_CST:
{
- hsa_op_immed *imm = new hsa_op_immed (exp);
+ hsa_op_immed *imm = new hsa_op_immed (exp);
if (addrtype != imm->m_type)
imm->m_type = addrtype;
return imm;
@@ -1957,8 +1996,10 @@ gen_hsa_addr (tree ref, hsa_bb *hbb, HOST_WIDE_INT *output_bitsize = NULL,
case SSA_NAME:
{
addrtype = hsa_get_segment_addr_type (BRIG_SEGMENT_PRIVATE);
- symbol = hsa_cfun->create_hsa_temporary (flat_addrtype);
- hsa_op_reg *r = hsa_cfun->reg_for_gimple_ssa (ref);
+ hsa_op_with_type *r = hsa_cfun->reg_for_gimple_ssa (ref);
+ if (r->m_type == BRIG_TYPE_B1)
+ r = r->get_in_type (BRIG_TYPE_U32, hbb);
+ symbol = hsa_cfun->create_hsa_temporary (r->m_type);
hbb->append_insn (new hsa_insn_mem (BRIG_OPCODE_ST, r->m_type,
r, new hsa_op_address (symbol)));
@@ -2247,13 +2288,18 @@ hsa_build_append_simple_mov (hsa_op_reg *dest, hsa_op_base *src, hsa_bb *hbb)
rules like when dealing with memory. */
BrigType16_t tp = mem_type_for_type (dest->m_type);
hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, tp, dest, src);
+ hsa_fixup_mov_insn_type (insn);
+ unsigned dest_size = hsa_type_bit_size (dest->m_type);
if (hsa_op_reg *sreg = dyn_cast <hsa_op_reg *> (src))
- gcc_assert (hsa_type_bit_size (dest->m_type)
- == hsa_type_bit_size (sreg->m_type));
+ gcc_assert (dest_size == hsa_type_bit_size (sreg->m_type));
else
- gcc_assert (hsa_type_bit_size (dest->m_type)
- == hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type));
-
+ {
+ unsigned imm_size
+ = hsa_type_bit_size (as_a <hsa_op_immed *> (src)->m_type);
+ gcc_assert ((dest_size == imm_size)
+ /* Eventually < 32bit registers will be promoted to 32bit. */
+ || (dest_size < 32 && imm_size == 32));
+ }
hbb->append_insn (insn);
}
@@ -2268,13 +2314,15 @@ gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
HOST_WIDE_INT bitsize, HOST_WIDE_INT bitpos,
hsa_bb *hbb)
{
- unsigned type_bitsize = hsa_type_bit_size (dest->m_type);
+ unsigned type_bitsize
+ = hsa_type_bit_size (hsa_extend_inttype_to_32bit (dest->m_type));
unsigned left_shift = type_bitsize - (bitsize + bitpos);
unsigned right_shift = left_shift + bitpos;
if (left_shift)
{
- hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
+ hsa_op_reg *value_reg_2
+ = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type));
hsa_op_immed *c = new hsa_op_immed (left_shift, BRIG_TYPE_U32);
hsa_insn_basic *lshift
@@ -2288,7 +2336,8 @@ gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
if (right_shift)
{
- hsa_op_reg *value_reg_2 = new hsa_op_reg (dest->m_type);
+ hsa_op_reg *value_reg_2
+ = new hsa_op_reg (hsa_extend_inttype_to_32bit (dest->m_type));
hsa_op_immed *c = new hsa_op_immed (right_shift, BRIG_TYPE_U32);
hsa_insn_basic *rshift
@@ -2301,8 +2350,10 @@ gen_hsa_insns_for_bitfield (hsa_op_reg *dest, hsa_op_reg *value_reg,
}
hsa_insn_basic *assignment
- = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, dest, value_reg);
+ = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type, NULL, value_reg);
+ hsa_fixup_mov_insn_type (assignment);
hbb->append_insn (assignment);
+ assignment->set_output_in_type (dest, 0, hbb);
}
@@ -2318,8 +2369,10 @@ gen_hsa_insns_for_bitfield_load (hsa_op_reg *dest, hsa_op_address *addr,
hsa_bb *hbb, BrigAlignment8_t align)
{
hsa_op_reg *value_reg = new hsa_op_reg (dest->m_type);
- hsa_insn_mem *mem = new hsa_insn_mem (BRIG_OPCODE_LD, dest->m_type, value_reg,
- addr);
+ hsa_insn_mem *mem
+ = new hsa_insn_mem (BRIG_OPCODE_LD,
+ hsa_extend_inttype_to_32bit (dest->m_type),
+ value_reg, addr);
mem->set_align (align);
hbb->append_insn (mem);
gen_hsa_insns_for_bitfield (dest, value_reg, bitsize, bitpos, hbb);
@@ -2446,9 +2499,10 @@ gen_hsa_insns_for_load (hsa_op_reg *dest, tree rhs, tree type, hsa_bb *hbb)
real_reg : imag_reg;
hsa_insn_basic *insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV,
- dest->m_type, dest, source);
-
+ dest->m_type, NULL, source);
+ hsa_fixup_mov_insn_type (insn);
hbb->append_insn (insn);
+ insn->set_output_in_type (dest, 0, hbb);
}
else if (TREE_CODE (rhs) == BIT_FIELD_REF
&& TREE_CODE (TREE_OPERAND (rhs, 0)) == SSA_NAME)
@@ -2584,6 +2638,7 @@ gen_hsa_insns_for_store (tree lhs, hsa_op_base *src, hsa_bb *hbb)
hsa_insn_basic *basic = new hsa_insn_basic (2, BRIG_OPCODE_MOV, mem_type,
new_value_reg, src);
+ hsa_fixup_mov_insn_type (basic);
hbb->append_insn (basic);
if (bitpos)
@@ -2954,8 +3009,10 @@ gen_hsa_cmp_insn_from_gimple (enum tree_code code, tree lhs, tree rhs,
? (BrigType16_t) BRIG_TYPE_B1 : dest->m_type;
hsa_insn_cmp *cmp = new hsa_insn_cmp (compare, dest_type);
- cmp->set_op (1, hsa_reg_or_immed_for_gimple_op (lhs, hbb));
- cmp->set_op (2, hsa_reg_or_immed_for_gimple_op (rhs, hbb));
+ hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (lhs, hbb);
+ cmp->set_op (1, op1->extend_int_to_32bit (hbb));
+ hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs, hbb);
+ cmp->set_op (2, op2->extend_int_to_32bit (hbb));
hbb->append_insn (cmp);
cmp->set_output_in_type (dest, 0, hbb);
@@ -2973,8 +3030,14 @@ gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
hsa_insn_basic *insn;
if (opcode == BRIG_OPCODE_MOV && hsa_needs_cvt (dest->m_type, op1->m_type))
- insn = new hsa_insn_cvt (dest, op1);
- else if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
+ {
+ insn = new hsa_insn_cvt (dest, op1);
+ hbb->append_insn (insn);
+ return;
+ }
+
+ op1 = op1->extend_int_to_32bit (hbb);
+ if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
{
BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type
: hsa_unsigned_type_for_type (op1->m_type);
@@ -2983,9 +3046,12 @@ gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
}
else
{
- insn = new hsa_insn_basic (2, opcode, dest->m_type, dest, op1);
+ BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
+ insn = new hsa_insn_basic (2, opcode, optype, NULL, op1);
- if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
+ if (opcode == BRIG_OPCODE_MOV)
+ hsa_fixup_mov_insn_type (insn);
+ else if (opcode == BRIG_OPCODE_ABS || opcode == BRIG_OPCODE_NEG)
{
/* ABS and NEG only exist in _s form :-/ */
if (insn->m_type == BRIG_TYPE_U32)
@@ -2996,9 +3062,7 @@ gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
}
hbb->append_insn (insn);
-
- if (opcode == BRIG_OPCODE_FIRSTBIT || opcode == BRIG_OPCODE_LASTBIT)
- insn->set_output_in_type (dest, 0, hbb);
+ insn->set_output_in_type (dest, 0, hbb);
}
/* Generate a binary instruction with OPCODE and append it to a basic block
@@ -3007,10 +3071,15 @@ gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest,
static void
gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
- hsa_op_base *op1, hsa_op_base *op2, hsa_bb *hbb)
+ hsa_op_with_type *op1, hsa_op_with_type *op2,
+ hsa_bb *hbb)
{
gcc_checking_assert (dest);
+ BrigType16_t optype = hsa_extend_inttype_to_32bit (dest->m_type);
+ op1 = op1->extend_int_to_32bit (hbb);
+ op2 = op2->extend_int_to_32bit (hbb);
+
if ((opcode == BRIG_OPCODE_SHL || opcode == BRIG_OPCODE_SHR)
&& is_a <hsa_op_immed *> (op2))
{
@@ -3026,9 +3095,10 @@ gen_hsa_binary_operation (int opcode, hsa_op_reg *dest,
i->set_type (hsa_unsigned_type_for_type (i->m_type));
}
- hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, dest->m_type, dest,
+ hsa_insn_basic *insn = new hsa_insn_basic (3, opcode, optype, NULL,
op1, op2);
hbb->append_insn (insn);
+ insn->set_output_in_type (dest, 0, hbb);
}
/* Generate HSA instructions for a single assignment. HBB is the basic block
@@ -3150,6 +3220,7 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
else if (TREE_CODE (rhs2) == SSA_NAME)
{
hsa_op_reg *s = hsa_cfun->reg_for_gimple_ssa (rhs2);
+ s = as_a <hsa_op_reg *> (s->extend_int_to_32bit (hbb));
hsa_op_reg *d = new hsa_op_reg (s->m_type);
hsa_op_immed *size_imm = new hsa_op_immed (bitsize, BRIG_TYPE_U32);
@@ -3253,8 +3324,11 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
hsa_op_with_type *op2 = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
+ op2 = op2->extend_int_to_32bit (hbb);
+ op3 = op3->extend_int_to_32bit (hbb);
- BrigType16_t utype = hsa_unsigned_type_for_type (dest->m_type);
+ BrigType16_t type = hsa_extend_inttype_to_32bit (dest->m_type);
+ BrigType16_t utype = hsa_unsigned_type_for_type (type);
if (is_a <hsa_op_immed *> (op2))
op2->m_type = utype;
if (is_a <hsa_op_immed *> (op3))
@@ -3262,10 +3336,11 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
hsa_insn_basic *insn
= new hsa_insn_basic (4, BRIG_OPCODE_CMOV,
- hsa_bittype_for_type (dest->m_type),
- dest, ctrl, op2, op3);
+ hsa_bittype_for_type (type),
+ NULL, ctrl, op2, op3);
hbb->append_insn (insn);
+ insn->set_output_in_type (dest, 0, hbb);
return;
}
case COMPLEX_EXPR:
@@ -3273,7 +3348,9 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
hsa_op_reg *dest
= hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
hsa_op_with_type *rhs1_reg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
+ rhs1_reg = rhs1_reg->extend_int_to_32bit (hbb);
hsa_op_with_type *rhs2_reg = hsa_reg_or_immed_for_gimple_op (rhs2, hbb);
+ rhs2_reg = rhs2_reg->extend_int_to_32bit (hbb);
if (hsa_seen_error ())
return;
@@ -3298,11 +3375,10 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
}
- hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (gimple_assign_lhs (assign));
-
+ hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
hsa_op_with_type *op1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
- hsa_op_with_type *op2 = rhs2 != NULL_TREE ?
- hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL;
+ hsa_op_with_type *op2
+ = rhs2 ? hsa_reg_or_immed_for_gimple_op (rhs2, hbb) : NULL;
if (hsa_seen_error ())
return;
@@ -3312,6 +3388,7 @@ gen_hsa_insns_for_operation_assignment (gimple *assign, hsa_bb *hbb)
case GIMPLE_TERNARY_RHS:
{
hsa_op_with_type *op3 = hsa_reg_or_immed_for_gimple_op (rhs3, hbb);
+ op3 = op3->extend_int_to_32bit (hbb);
hsa_insn_basic *insn = new hsa_insn_basic (4, opcode, dest->m_type, dest,
op1, op2, op3);
hbb->append_insn (insn);
@@ -3407,14 +3484,15 @@ gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
tree highest = get_switch_high (s);
hsa_op_reg *index = hsa_cfun->reg_for_gimple_ssa (index_tree);
+ index = as_a <hsa_op_reg *> (index->extend_int_to_32bit (hbb));
hsa_op_reg *cmp1_reg = new hsa_op_reg (BRIG_TYPE_B1);
- hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest);
+ hsa_op_immed *cmp1_immed = new hsa_op_immed (lowest, true);
hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_GE, cmp1_reg->m_type,
cmp1_reg, index, cmp1_immed));
hsa_op_reg *cmp2_reg = new hsa_op_reg (BRIG_TYPE_B1);
- hsa_op_immed *cmp2_immed = new hsa_op_immed (highest);
+ hsa_op_immed *cmp2_immed = new hsa_op_immed (highest, true);
hbb->append_insn (new hsa_insn_cmp (BRIG_COMPARE_LE, cmp2_reg->m_type,
cmp2_reg, index, cmp2_immed));
@@ -3444,7 +3522,7 @@ gen_hsa_insns_for_switch_stmt (gswitch *s, hsa_bb *hbb)
hsa_op_reg *sub_index = new hsa_op_reg (index->m_type);
hbb->append_insn (new hsa_insn_basic (3, BRIG_OPCODE_SUB, sub_index->m_type,
sub_index, index,
- new hsa_op_immed (lowest)));
+ new hsa_op_immed (lowest, true)));
hsa_op_base *tmp = sub_index->get_in_type (BRIG_TYPE_U64, hbb);
sub_index = as_a <hsa_op_reg *> (tmp);
@@ -3760,7 +3838,6 @@ void
hsa_insn_basic::set_output_in_type (hsa_op_reg *dest, unsigned op_index,
hsa_bb *hbb)
{
- hsa_insn_basic *insn;
gcc_checking_assert (op_output_p (op_index));
if (dest->m_type == m_type)
@@ -3769,15 +3846,28 @@ hsa_insn_basic::set_output_in_type (hsa_op_reg *dest, unsigned op_index,
return;
}
- hsa_op_reg *tmp = new hsa_op_reg (m_type);
- set_op (op_index, tmp);
-
+ hsa_insn_basic *insn;
+ hsa_op_reg *tmp;
if (hsa_needs_cvt (dest->m_type, m_type))
- insn = new hsa_insn_cvt (dest, tmp);
+ {
+ tmp = new hsa_op_reg (m_type);
+ insn = new hsa_insn_cvt (dest, tmp);
+ }
+ else if (hsa_type_bit_size (dest->m_type) == hsa_type_bit_size (m_type))
+ {
+ /* When output, HSA registers do not really have types, only sizes, so if
+ the sizes match, we can use the register directly. */
+ set_op (op_index, dest);
+ return;
+ }
else
- insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type,
- dest, tmp->get_in_type (dest->m_type, hbb));
-
+ {
+ tmp = new hsa_op_reg (m_type);
+ insn = new hsa_insn_basic (2, BRIG_OPCODE_MOV, dest->m_type,
+ dest, tmp->get_in_type (dest->m_type, hbb));
+ hsa_fixup_mov_insn_type (insn);
+ }
+ set_op (op_index, tmp);
hbb->append_insn (insn);
}
@@ -4200,6 +4290,7 @@ gen_hsa_clrsb (gcall *call, hsa_bb *hbb)
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
tree rhs1 = gimple_call_arg (call, 0);
hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
+ arg->extend_int_to_32bit (hbb);
BrigType16_t bittype = hsa_bittype_for_type (arg->m_type);
unsigned bitsize = tree_to_uhwi (TYPE_SIZE (TREE_TYPE (rhs1)));
@@ -4272,6 +4363,7 @@ gen_hsa_ffs (gcall *call, hsa_bb *hbb)
tree rhs1 = gimple_call_arg (call, 0);
hsa_op_with_type *arg = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
+ arg = arg->extend_int_to_32bit (hbb);
hsa_op_reg *tmp = new hsa_op_reg (BRIG_TYPE_U32);
hsa_insn_srctype *insn = new hsa_insn_srctype (2, BRIG_OPCODE_LASTBIT,
@@ -4361,7 +4453,9 @@ gen_hsa_divmod (gcall *call, hsa_bb *hbb)
tree rhs1 = gimple_call_arg (call, 1);
hsa_op_with_type *arg0 = hsa_reg_or_immed_for_gimple_op (rhs0, hbb);
+ arg0 = arg0->extend_int_to_32bit (hbb);
hsa_op_with_type *arg1 = hsa_reg_or_immed_for_gimple_op (rhs1, hbb);
+ arg1 = arg1->extend_int_to_32bit (hbb);
hsa_op_reg *dest0 = new hsa_op_reg (arg0->m_type);
hsa_op_reg *dest1 = new hsa_op_reg (arg1->m_type);
@@ -4374,11 +4468,13 @@ gen_hsa_divmod (gcall *call, hsa_bb *hbb)
hbb->append_insn (insn);
hsa_op_reg *dest = hsa_cfun->reg_for_gimple_ssa (lhs);
+ BrigType16_t dst_type = hsa_extend_inttype_to_32bit (dest->m_type);
BrigType16_t src_type = hsa_bittype_for_type (dest0->m_type);
- insn = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dest->m_type,
- src_type, dest, dest0, dest1);
+ insn = new hsa_insn_packed (3, BRIG_OPCODE_COMBINE, dst_type,
+ src_type, NULL, dest0, dest1);
hbb->append_insn (insn);
+ insn->set_output_in_type (dest, 0, hbb);
}
/* Set VALUE to a shadow kernel debug argument and append a new instruction
@@ -4936,8 +5032,8 @@ gen_hsa_atomic_for_builtin (bool ret_orig, enum BrigAtomicOperation acode,
tgt = addr;
}
- hsa_op_base *op = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1),
- hbb);
+ hsa_op_with_type *op
+ = hsa_reg_or_immed_for_gimple_op (gimple_call_arg (stmt, 1), hbb);
if (lhs)
{
atominsn->set_op (0, dest);
diff --git a/libgomp/testsuite/libgomp.hsa.c/pr82416.c b/libgomp/testsuite/libgomp.hsa.c/pr82416.c
new file mode 100644
index 00000000000..b89d421e8f3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.hsa.c/pr82416.c
@@ -0,0 +1,37 @@
+char __attribute__ ((noipa))
+toup (char X)
+{
+ if (X >= 97 && X <= 122)
+ return X - 32;
+ else
+ return X;
+}
+
+char __attribute__ ((noipa))
+target_toup (char X)
+{
+ char r;
+#pragma omp target map(to:X) map(from:r)
+ {
+ if (X >= 97 && X <= 122)
+ r = X - 32;
+ else
+ r = X;
+ }
+ return r;
+}
+
+int main (int argc, char **argv)
+{
+ char a = 'a';
+ if (toup (a) != target_toup (a))
+ __builtin_abort ();
+ a = 'Z';
+ if (toup (a) != target_toup (a))
+ __builtin_abort ();
+ a = 5;
+ if (toup (a) != target_toup (a))
+ __builtin_abort ();
+
+ return 0;
+}
--
2.14.1