From patchwork Mon Sep 4 16:24:42 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1829704 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4RfYrn6t14z1yhH for ; Tue, 5 Sep 2023 02:25:12 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9F764385802F for ; Mon, 4 Sep 2023 16:25:10 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org 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 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 X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" 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" "=R") (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) != 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üße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 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