From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 7A4473858425 for ; Mon, 4 Sep 2023 16:24:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 7A4473858425 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.02,226,1688457600"; d="scan'208,223";a="18170527" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 04 Sep 2023 08:24:54 -0800 IronPort-SDR: HDGceH5Vqt6l1y6lyTofkaTqVRMuS+yssJVSQjy6mgmNDVQckWmILMbAHjxUT8UasgD/nYZFyV 26RI36T5grfWSRpdrMttApP3+JdCurAxuVI9BUzaTXENBmx2Ad63+yWlDrXx1r9De+5VRp+YD7 z3wkbA+vDWz4RXS1U7s4Y51lOJ2iZki+J/tCTog/A5GiQZnYiy2pRdrzCVUwPEwySzTKtIdT0l kZ9pkcuCAQlZdKiAO6hQ2eOSeUPsSRBRntyWbOkP0NVNzNqs/9SBxvjj1E67XQDnj9yqJs6lnC PiQ= From: Thomas Schwinge To: , Roger Sayle , "Tom de Vries" Subject: [WIP] nvptx: Also allow immediate input operand to 'bitrev2' User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Mon, 4 Sep 2023 18:24:42 +0200 Message-ID: <87bkehx5x1.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="=-=-=" X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-11.8 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,KAM_SHORT,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: --=-=-= Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Hi! I'm working towards reviewing some (of Roger's) GCC/nvptx patches, and therefore learning some more GCC/nvptx, and generally RTL etc., and the conventions around it. Please bear with me asking "obvious" questions. For the PTX bit reverse instruction, GCC/nvptx currently ("forever") defines: (define_insn "bitrev2" [(set (match_operand:SDIM 0 "nvptx_register_operand" "=3DR") (unspec:SDIM [(match_operand:SDIM 1 "nvptx_register_operand" "R"= )] UNSPEC_BITREV))] "" "%.\\tbrev.b%T0\\t%0, %1;") ..., with: (define_predicate "nvptx_register_operand" (match_code "reg") { return register_operand (op, mode); }) (define_constraint "R" "A pseudo register." (match_code "reg")) That is, only a register input operand is permitted, not an immediate. However, I don't see such a restriction in the manual, . If I change that 'define_insn': - (unspec:SDIM [(match_operand:SDIM 1 "nvptx_register_operand" "R= ")] + (unspec:SDIM [(match_operand:SDIM 1 "nvptx_nonmemory_operand" "= Ri")] ..., with (existing): (define_predicate "nvptx_nonmemory_operand" (match_code "reg,const_int,const_double") { return (REG_P (op) ? register_operand (op, mode) : immediate_operand (op, mode)); }) ..., then a simple code: return __builtin_nvptx_brev(0xe6a2c480) !=3D 0x01234567; ... for '-O1' subsequently changes: [...] .reg .u32 %r22; -.reg .u32 %r25; .reg .pred %r29; -mov.u32 %r25,-425540480; -brev.b32 %r22,%r25; +brev.b32 %r22,-425540480; setp.ne.u32 %r29,%r22,19088743; [...] (I understand that, in the end, that's probably equivalent, assuming that the later PTX -> SASS compiler does the same optimization, but I find it easier to read: one less '.reg' to keep track of textually/mentally.) Does that make sense to you, too? I'd then extend my attached "[WIP] nvptx: Also allow immediate input operand to 'bitrev2'" with a test case. (Similarly then, a number of other GCC/nvptx 'define_insn's to be reviewed/revised, later on.) Relatedly, I see that a lot of GCC/nvptx' two input operands instructions ('add3', etc.) similarly do allow for their second input operand to be an immediate in addition to a register. I suppose that only allowing for the second input operand to be an immediate is sufficient/desirable: reduces load on the matching; we shouldn't ever end up with 'IMM + IMM', for example: should've optimized that before. As I learn from , "for commutative [...] operators, a constant is always made the second operand". (Confirmed to ICE if swapping that around for 'add3'; so, that's all as expected.) Gr=C3=BC=C3=9Fe Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstra=C3=9Fe 201= , 80634 M=C3=BCnchen; Gesellschaft mit beschr=C3=A4nkter Haftung; Gesch=C3= =A4ftsf=C3=BChrer: Thomas Heurung, Frank Th=C3=BCrauf; Sitz der Gesellschaf= t: M=C3=BCnchen; Registergericht M=C3=BCnchen, HRB 106955 --=-=-= Content-Type: text/x-diff Content-Disposition: inline; filename="0001-WIP-nvptx-Also-allow-immediate-input-operand-to-bitr.patch" >From 4a5138eb61a026ad6bc5470a648ebc596af1b1ed Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Mon, 4 Sep 2023 16:48:53 +0200 Subject: [PATCH] [WIP] nvptx: Also allow immediate input operand to 'bitrev2' --- gcc/config/nvptx/nvptx.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md index 1bb93045403..e1c822f2ea8 100644 --- a/gcc/config/nvptx/nvptx.md +++ b/gcc/config/nvptx/nvptx.md @@ -636,7 +636,7 @@ (define_insn "bitrev2" [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R") - (unspec:SDIM [(match_operand:SDIM 1 "nvptx_register_operand" "R")] + (unspec:SDIM [(match_operand:SDIM 1 "nvptx_nonmemory_operand" "Ri")] UNSPEC_BITREV))] "" "%.\\tbrev.b%T0\\t%0, %1;") -- 2.34.1 --=-=-=--