[PR 82416] Do not extend operands to at least 32 bits

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.

From-SVN: r253538

gcc/ChangeLog

+2017-10-09  Martin Jambor  <mjambor@suse.cz>
+
+	PR hsa/82416
+	* 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.
+
 2017-10-08  Segher Boessenkool  <segher@kernel.crashing.org>
 
 	* config/rs6000/rs6000.md (conditional branch): Clean up formatting.

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);

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);

libgomp/ChangeLog

+2017-10-09  Martin Jambor  <mjambor@suse.cz>
+
+	PR hsa/82416
+	* testsuite/libgomp.hsa.c/pr82416.c: New test.
+
 2017-10-07  Tom de Vries  <tom@codesourcery.com>
 
 	* testsuite/libgomp.oacc-fortran/firstprivate-1.f90 (firstprivate):

libgomp/testsuite/libgomp.hsa.c/pr82416.c

+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;
+}