[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.



git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@253538 138bc75d-0d04-0410-961f-82ee72b054a4

2 files changed, 42 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index afa373ec7bb..1c5c5a9a912 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+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):
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;
+}