[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

1 files changed, 37 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.hsa.c/pr82416.c b/libgomp/testsuite/libgomp.hsa.c/pr82416.c
new file mode 100644
index 0000000..b89d421
--- /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;
+}