From patchwork Mon Oct 9 09:49:56 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 823156 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-463734-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="EBbDQxXw"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3y9b8938Z7z9tXx for ; Mon, 9 Oct 2017 20:50:16 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; q=dns; s= default; b=shu2cz9QXC2Bhp9AvGO28pEeCahbQ+mNCbdf4/erjVpZoMV3SeEoG sCNVSiqx7lKvm8QkgJ96Sf4R8nMxbnzjMJRuOVtKreurBxmtWerHy5xPSX037Jwn 8EP52M5TKEcCvq+/IBUNz/2ezzMcUkchyZtwK7Bnyesm38lrBXKj/I= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:date :from:to:subject:message-id:mime-version:content-type; s= default; bh=MdSbzKk54Qe9NznbEvsPp7S2RYU=; b=EBbDQxXwiTVelHziGB1N s4CmzKRljDMARYgllgUTE6aBVUQM5hmBpRwY9a3xpDHdZnjNzZIzwAHQj4phs1qP sjpSVHouxNR/9vNyz1AmrHyKlke+NrkGAY7Ixf+2LB+n+5wPuZpiAXDTAC6+IfKS 8YQUK0n00zAAjp84DdLIqqA= Received: (qmail 67253 invoked by alias); 9 Oct 2017 09:50:04 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 67222 invoked by uid 89); 9 Oct 2017 09:50:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.2 spammy=_3, ctrl, mandated X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Mon, 09 Oct 2017 09:49:59 +0000 Received: from relay2.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 1DA96ABFC for ; Mon, 9 Oct 2017 09:49:57 +0000 (UTC) Date: Mon, 9 Oct 2017 11:49:56 +0200 From: Martin Jambor To: GCC Patches Subject: [HSA, PR 82416] Do not extend operands to at least 32 bits Message-ID: <20171009094956.sxlnqpm5mw6bqlrj@virgil.suse.cz> Mail-Followup-To: GCC Patches MIME-Version: 1.0 Content-Disposition: inline User-Agent: NeoMutt/20170609 (1.8.3) X-IsSubscribed: yes 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 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 (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 (src)->m_type)); - + { + unsigned imm_size + = hsa_type_bit_size (as_a (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 (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 (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 (op2)) op2->m_type = utype; if (is_a (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 (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 (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; +}