From patchwork Mon Aug 15 16:48:35 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Martin Jambor X-Patchwork-Id: 659346 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 3sChK61Rtxz9t39 for ; Tue, 16 Aug 2016 02:48:56 +1000 (AEST) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=uQ2Q430W; dkim-atps=neutral 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=afy3lp2+4QX7R6UlNnMWjPYuE53CAJd5uCWbSGcF7Pzr3ty5ZF3Mq +JZlNMFe2nXJeLBQA0psSUYDY+b/gbeqIzSa38+EbyHHSDfVdMksKbIDbm00nNB4 GH20lmYWuD/5EsCXEotq7Ah5tnk+nJRA7hWW3gZt1i7F0HerNCZhGo= 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=auMpGpqIU8QIrV3tMbEjXDY7DxQ=; b=uQ2Q430WvU9qL0v63FD1 JFY4ePxR8QkQMMlzTXj1XJjntpRYBEla1yF4FUoRkP/uBdu+AysVrbTxG9Eki/bU XRU2nSEo7GHRc85lD0yWqS+rCp/m3AqusYmm+aRKyKchbsEnX4PfGT1aa9by+6D5 sXLCv2VnZubwx9S1HwnM748= Received: (qmail 94600 invoked by alias); 15 Aug 2016 16:48:48 -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 94564 invoked by uid 89); 15 Aug 2016 16:48:47 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=BAYES_00, SPF_PASS autolearn=ham version=3.3.2 spammy=__builtin_abort, m_type 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, 15 Aug 2016 16:48:37 +0000 Received: from relay1.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 64E2FAD33 for ; Mon, 15 Aug 2016 16:48:35 +0000 (UTC) Date: Mon, 15 Aug 2016 18:48:35 +0200 From: Martin Jambor To: GCC Patches Subject: [hsa-branch] Fix issues with firstbit and popcount source types Message-ID: <20160815164834.47upouxs7twqtgp6@virgil.suse.cz> Mail-Followup-To: GCC Patches MIME-Version: 1.0 Content-Disposition: inline User-Agent: Mutt/1.6.2-neo (2016-06-11) X-IsSubscribed: yes Hi, this patch addresses a regression caused by my patch that avoided useless register copies but in a few cases caused us to generate instruction types that did not make the finalizer happy. Fixed thusly. I am going to commit to the branch now and will queue it for trunk for later. Thanks, Martin 2016-08-12 Martin Jambor * hsa-gen.c (gen_hsa_unary_operation): Make sure the function does not use bittype source type for firstbit and lastbit operations. (gen_hsa_popcount_to_dest): Make sure the function uses a bittype source type. libgomp/ * testsuite/libgomp.hsa.c/bits-insns.c: New test. --- gcc/hsa-gen.c | 13 +++-- libgomp/testsuite/libgomp.hsa.c/bits-insns.c | 73 ++++++++++++++++++++++++++++ 2 files changed, 82 insertions(+), 4 deletions(-) create mode 100644 libgomp/testsuite/libgomp.hsa.c/bits-insns.c diff --git a/gcc/hsa-gen.c b/gcc/hsa-gen.c index baa20b9..c946b2f 100644 --- a/gcc/hsa-gen.c +++ b/gcc/hsa-gen.c @@ -2957,8 +2957,12 @@ gen_hsa_unary_operation (BrigOpcode opcode, hsa_op_reg *dest, 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_srctype (2, opcode, BRIG_TYPE_U32, op1->m_type, NULL, - op1); + { + BrigType16_t srctype = hsa_type_integer_p (op1->m_type) ? op1->m_type + : hsa_unsigned_type_for_type (op1->m_type); + insn = new hsa_insn_srctype (2, opcode, BRIG_TYPE_U32, srctype, NULL, + op1); + } else { insn = new hsa_insn_basic (2, opcode, dest->m_type, dest, op1); @@ -4250,12 +4254,13 @@ gen_hsa_popcount_to_dest (hsa_op_reg *dest, hsa_op_with_type *arg, hsa_bb *hbb) if (hsa_type_bit_size (arg->m_type) < 32) arg = arg->get_in_type (BRIG_TYPE_B32, hbb); + BrigType16_t srctype = hsa_bittype_for_type (arg->m_type); if (!hsa_btype_p (arg->m_type)) - arg = arg->get_in_type (hsa_bittype_for_type (arg->m_type), hbb); + arg = arg->get_in_type (srctype, hbb); hsa_insn_srctype *popcount = new hsa_insn_srctype (2, BRIG_OPCODE_POPCOUNT, BRIG_TYPE_U32, - arg->m_type, NULL, arg); + srctype, NULL, arg); hbb->append_insn (popcount); popcount->set_output_in_type (dest, 0, hbb); } diff --git a/libgomp/testsuite/libgomp.hsa.c/bits-insns.c b/libgomp/testsuite/libgomp.hsa.c/bits-insns.c new file mode 100644 index 0000000..21cac72 --- /dev/null +++ b/libgomp/testsuite/libgomp.hsa.c/bits-insns.c @@ -0,0 +1,73 @@ +#include + +#define N 12 + +int main() +{ + unsigned int arguments[N] = {0u, 1u, 2u, 3u, 111u, 333u, 444u, 0x80000000u, 0x0000ffffu, 0xf0000000u, 0xff000000u, 0xffffffffu}; + int clrsb[N] = {}; + int clz[N] = {}; + int ctz[N] = {}; + int ffs[N] = {}; + int parity[N] = {}; + int popcount[N] = {}; + + int ref_clrsb[N] = {}; + int ref_clz[N] = {}; + int ref_ctz[N] = {}; + int ref_ffs[N] = {}; + int ref_parity[N] = {}; + int ref_popcount[N] = {}; + + for (unsigned i = 0; i < N; i++) + { + ref_clrsb[i] = __builtin_clrsb (arguments[i]); + ref_clz[i] = __builtin_clz (arguments[i]); + ref_ctz[i] = __builtin_ctz (arguments[i]); + ref_ffs[i] = __builtin_ffs (arguments[i]); + ref_parity[i] = __builtin_parity (arguments[i]); + ref_popcount[i] = __builtin_popcount (arguments[i]); + } + + #pragma omp target map(from:clz, ctz, ffs, parity, popcount) + { + for (unsigned i = 0; i < N; i++) + { + clrsb[i] = __builtin_clrsb (arguments[i]); + clz[i] = __builtin_clz (arguments[i]); + ctz[i] = __builtin_ctz (arguments[i]); + ffs[i] = __builtin_ffs (arguments[i]); + parity[i] = __builtin_parity (arguments[i]); + popcount[i] = __builtin_popcount (arguments[i]); + } + } + + for (unsigned i = 0; i < N; i++) + if (ref_clrsb[i] != clrsb[i]) + __builtin_abort (); + + /* CLZ of zero is undefined for zero. */ + for (unsigned i = 1; i < N; i++) + if (ref_clz[i] != clz[i]) + __builtin_abort (); + + /* Likewise for ctz */ + for (unsigned i = 1; i < N; i++) + if (ref_ctz[i] != ctz[i]) + __builtin_abort (); + + for (unsigned i = 0; i < N; i++) + if (ref_ffs[i] != ffs[i]) + __builtin_abort (); + + for (unsigned i = 0; i < N; i++) + if (ref_parity[i] != parity[i]) + __builtin_abort (); + + for (unsigned i = 0; i < N; i++) + if (ref_popcount[i] != popcount[i]) + __builtin_abort (); + + return 0; +} +