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: 137454 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:ab0a:0:b0:3f2:4152:657d with SMTP id m10csp1191427vqo; Mon, 4 Sep 2023 09:25:29 -0700 (PDT) X-Google-Smtp-Source: AGHT+IGz5yyGUS75ZmZJQbjnbiAVTY7Wh3FgbrQpGvR2TlmRDtlDeRMMJnq8Ws9lLVksxeSKoQFp X-Received: by 2002:aa7:c602:0:b0:51e:ed6:df38 with SMTP id h2-20020aa7c602000000b0051e0ed6df38mr7632412edq.13.1693844729369; Mon, 04 Sep 2023 09:25:29 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1693844729; cv=none; d=google.com; s=arc-20160816; b=LOA/1CHX05wX6mIEc3WEHlNPxFO9BfNxZPScdcaCFGbB1o861rf3ks4ifwpgCTKNy8 V96/XIAg3VWrzfI16pinKMAY48c/r07jEFcLqj7lNrwvzfx6LsPCBX8HlcFFqxymL7cD 32tdlhDWwOvt0nhI153d/l1IOr1AhRmXJIxC9YxPP6yNUVrD0z2YgFAs8X3buGgBgv8B vL/bnlZgzL1CTFE/FaY4Wly6CjvxpVUXEks8cGPbwJSvTOdpSd2m5qC84GCJ3I3atL+m 7tk9gOd03OyuPeQ7zNgxM/c6dO56SHj7Z0rS+IpXi46dpmC+iBl0eM+edjddwg3r3PV/ xwWw== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:mime-version:message-id:date :user-agent:subject:to:from:ironport-sdr:dmarc-filter:delivered-to; bh=GibXWuM9T8f6jD53Z3v/KQeIzepAaKdRoI2VkiDTu8I=; fh=VIZAUtIjXWToTH0uKYvgjcOeFS47G8vG/jgGJD+v5ag=; b=ezbA64Z5buSvDTIrknaEq3qk66V+FDtna0Ch50tMp2UyBJwTkR6E/jHTRw29cukKBd 7FJCs5QxtuIHXgChu5/qVIJ+00csGq/rwTFCvvY9tbZ7pzDF9A+YDrOJZ2jUiEyScoEm fusMm9Zbo5AjN4tHOhtGT12pw9XKYz33nxatP0x6G1KrUQrfQZmsOLXFbPhNDDfGV9IJ RUQ+gop6xCXRvTsW/TuPAu3i6YP3qcFFD24z/ocxwyOwcDlOg72sWXbWLslXxflqsHKQ s7s7yaES4gohq6ftWspCnHiqBLMkqUAP3kErhemVTpFD4a2e41jli5/U1gzI0wQXPyXe GSRg== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id k15-20020aa7c38f000000b00529fd6decf9si6552635edq.11.2023.09.04.09.25.29 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 04 Sep 2023 09:25:29 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 4FA843857351 for ; Mon, 4 Sep 2023 16:25:23 +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+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1776124931281744741 X-GMAIL-MSGID: 1776124931281744741 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