Author: Artem Belevich Date: 2024-08-08T13:28:05-07:00 New Revision: 08500dc4838cff0b99fbbcb2b571664745a7b9e6
URL: https://github.com/llvm/llvm-project/commit/08500dc4838cff0b99fbbcb2b571664745a7b9e6 DIFF: https://github.com/llvm/llvm-project/commit/08500dc4838cff0b99fbbcb2b571664745a7b9e6.diff LOG: Revert "[NVPTX] support switch statement with brx.idx (#102400)" This reverts commit ba976971898d74df38d155c55e008c898120d1e4. Added: Modified: llvm/include/llvm/CodeGen/TargetLowering.h llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp llvm/lib/Target/NVPTX/NVPTXISelLowering.h llvm/lib/Target/NVPTX/NVPTXInstrInfo.td Removed: llvm/test/CodeGen/NVPTX/jump-table.ll ################################################################################ diff --git a/llvm/include/llvm/CodeGen/TargetLowering.h b/llvm/include/llvm/CodeGen/TargetLowering.h index 5b2214fa66c40b..9ccdbab008aec8 100644 --- a/llvm/include/llvm/CodeGen/TargetLowering.h +++ b/llvm/include/llvm/CodeGen/TargetLowering.h @@ -3843,10 +3843,6 @@ class TargetLowering : public TargetLoweringBase { /// returned value is a member of the MachineJumpTableInfo::JTEntryKind enum. virtual unsigned getJumpTableEncoding() const; - virtual MVT getJumpTableRegTy(const DataLayout &DL) const { - return getPointerTy(DL); - } - virtual const MCExpr * LowerCustomJumpTableEntry(const MachineJumpTableInfo * /*MJTI*/, const MachineBasicBlock * /*MBB*/, unsigned /*uid*/, diff --git a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp index 37ba62911ec70b..1f4436fb3a4966 100644 --- a/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp @@ -2977,7 +2977,7 @@ void SelectionDAGBuilder::visitJumpTable(SwitchCG::JumpTable &JT) { // Emit the code for the jump table assert(JT.SL && "Should set SDLoc for SelectionDAG!"); assert(JT.Reg != -1U && "Should lower JT Header first!"); - EVT PTy = DAG.getTargetLoweringInfo().getJumpTableRegTy(DAG.getDataLayout()); + EVT PTy = DAG.getTargetLoweringInfo().getPointerTy(DAG.getDataLayout()); SDValue Index = DAG.getCopyFromReg(getControlRoot(), *JT.SL, JT.Reg, PTy); SDValue Table = DAG.getJumpTable(JT.JTI, PTy); SDValue BrJumpTable = DAG.getNode(ISD::BR_JT, *JT.SL, MVT::Other, @@ -3005,13 +3005,12 @@ void SelectionDAGBuilder::visitJumpTableHeader(SwitchCG::JumpTable &JT, // This value may be smaller or larger than the target's pointer type, and // therefore require extension or truncating. const TargetLowering &TLI = DAG.getTargetLoweringInfo(); - SwitchOp = - DAG.getZExtOrTrunc(Sub, dl, TLI.getJumpTableRegTy(DAG.getDataLayout())); + SwitchOp = DAG.getZExtOrTrunc(Sub, dl, TLI.getPointerTy(DAG.getDataLayout())); unsigned JumpTableReg = - FuncInfo.CreateReg(TLI.getJumpTableRegTy(DAG.getDataLayout())); - SDValue CopyTo = - DAG.getCopyToReg(getControlRoot(), dl, JumpTableReg, SwitchOp); + FuncInfo.CreateReg(TLI.getPointerTy(DAG.getDataLayout())); + SDValue CopyTo = DAG.getCopyToReg(getControlRoot(), dl, + JumpTableReg, SwitchOp); JT.Reg = JumpTableReg; if (!JTH.FallthroughUnreachable) { diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index bf647c88f00e28..516fc7339a4bf3 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -25,7 +25,6 @@ #include "llvm/CodeGen/Analysis.h" #include "llvm/CodeGen/ISDOpcodes.h" #include "llvm/CodeGen/MachineFunction.h" -#include "llvm/CodeGen/MachineJumpTableInfo.h" #include "llvm/CodeGen/MachineMemOperand.h" #include "llvm/CodeGen/SelectionDAG.h" #include "llvm/CodeGen/SelectionDAGNodes.h" @@ -583,7 +582,9 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM, setOperationAction(ISD::ROTR, MVT::i8, Expand); setOperationAction(ISD::BSWAP, MVT::i16, Expand); - setOperationAction(ISD::BR_JT, MVT::Other, Custom); + // Indirect branch is not supported. + // This also disables Jump Table creation. + setOperationAction(ISD::BR_JT, MVT::Other, Expand); setOperationAction(ISD::BRIND, MVT::Other, Expand); setOperationAction(ISD::GlobalAddress, MVT::i32, Custom); @@ -944,9 +945,6 @@ const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const { MAKE_CASE(NVPTXISD::Dummy) MAKE_CASE(NVPTXISD::MUL_WIDE_SIGNED) MAKE_CASE(NVPTXISD::MUL_WIDE_UNSIGNED) - MAKE_CASE(NVPTXISD::BrxEnd) - MAKE_CASE(NVPTXISD::BrxItem) - MAKE_CASE(NVPTXISD::BrxStart) MAKE_CASE(NVPTXISD::Tex1DFloatS32) MAKE_CASE(NVPTXISD::Tex1DFloatFloat) MAKE_CASE(NVPTXISD::Tex1DFloatFloatLevel) @@ -2787,8 +2785,6 @@ NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const { return LowerFP_ROUND(Op, DAG); case ISD::FP_EXTEND: return LowerFP_EXTEND(Op, DAG); - case ISD::BR_JT: - return LowerBR_JT(Op, DAG); case ISD::VAARG: return LowerVAARG(Op, DAG); case ISD::VASTART: @@ -2814,41 +2810,6 @@ NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const { } } -SDValue NVPTXTargetLowering::LowerBR_JT(SDValue Op, SelectionDAG &DAG) const { - SDLoc DL(Op); - SDValue Chain = Op.getOperand(0); - const auto *JT = cast<JumpTableSDNode>(Op.getOperand(1)); - SDValue Index = Op.getOperand(2); - - unsigned JId = JT->getIndex(); - MachineJumpTableInfo *MJTI = DAG.getMachineFunction().getJumpTableInfo(); - ArrayRef<MachineBasicBlock *> MBBs = MJTI->getJumpTables()[JId].MBBs; - - SDValue IdV = DAG.getConstant(JId, DL, MVT::i32); - - // Generate BrxStart node - SDVTList VTs = DAG.getVTList(MVT::Other, MVT::Glue); - Chain = DAG.getNode(NVPTXISD::BrxStart, DL, VTs, Chain, IdV); - - // Generate BrxItem nodes - assert(!MBBs.empty()); - for (MachineBasicBlock *MBB : MBBs.drop_back()) - Chain = DAG.getNode(NVPTXISD::BrxItem, DL, VTs, Chain.getValue(0), - DAG.getBasicBlock(MBB), Chain.getValue(1)); - - // Generate BrxEnd nodes - SDValue EndOps[] = {Chain.getValue(0), DAG.getBasicBlock(MBBs.back()), Index, - IdV, Chain.getValue(1)}; - SDValue BrxEnd = DAG.getNode(NVPTXISD::BrxEnd, DL, VTs, EndOps); - - return BrxEnd; -} - -// This will prevent AsmPrinter from trying to print the jump tables itself. -unsigned NVPTXTargetLowering::getJumpTableEncoding() const { - return MachineJumpTableInfo::EK_Inline; -} - // This function is almost a copy of SelectionDAG::expandVAArg(). // The only diff is that this one produces loads from local address space. SDValue NVPTXTargetLowering::LowerVAARG(SDValue Op, SelectionDAG &DAG) const { diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h index 32e6b044b0de1f..63262961b363ed 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h @@ -62,9 +62,6 @@ enum NodeType : unsigned { BFI, PRMT, DYNAMIC_STACKALLOC, - BrxStart, - BrxItem, - BrxEnd, Dummy, LoadV2 = ISD::FIRST_TARGET_MEMORY_OPCODE, @@ -583,11 +580,6 @@ class NVPTXTargetLowering : public TargetLowering { return true; } - // The default is the same as pointer type, but brx.idx only accepts i32 - MVT getJumpTableRegTy(const DataLayout &) const override { return MVT::i32; } - - unsigned getJumpTableEncoding() const override; - bool enableAggressiveFMAFusion(EVT VT) const override { return true; } // The default is to transform llvm.ctlz(x, false) (where false indicates that @@ -645,8 +637,6 @@ class NVPTXTargetLowering : public TargetLowering { SDValue LowerSelect(SDValue Op, SelectionDAG &DAG) const; - SDValue LowerBR_JT(SDValue Op, SelectionDAG &DAG) const; - SDValue LowerVAARG(SDValue Op, SelectionDAG &DAG) const; SDValue LowerVASTART(SDValue Op, SelectionDAG &DAG) const; diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td index 904574b2e1d660..6a096fa5acea7c 100644 --- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -3880,44 +3880,6 @@ def DYNAMIC_STACKALLOC64 : [(set Int64Regs:$ptr, (dyn_alloca Int64Regs:$size, (i32 timm:$align)))]>, Requires<[hasPTX<73>, hasSM<52>]>; - -// -// BRX -// - -def SDTBrxStartProfile : SDTypeProfile<0, 1, [SDTCisInt<0>]>; -def SDTBrxItemProfile : SDTypeProfile<0, 1, [SDTCisVT<0, OtherVT>]>; -def SDTBrxEndProfile : SDTypeProfile<0, 3, [SDTCisVT<0, OtherVT>, SDTCisInt<1>, SDTCisInt<2>]>; - -def brx_start : - SDNode<"NVPTXISD::BrxStart", SDTBrxStartProfile, - [SDNPHasChain, SDNPOutGlue, SDNPSideEffect]>; -def brx_item : - SDNode<"NVPTXISD::BrxItem", SDTBrxItemProfile, - [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>; -def brx_end : - SDNode<"NVPTXISD::BrxEnd", SDTBrxEndProfile, - [SDNPHasChain, SDNPInGlue, SDNPSideEffect]>; - -let isTerminator = 1, isBranch = 1, isIndirectBranch = 1 in { - - def BRX_START : - NVPTXInst<(outs), (ins i32imm:$id), - "$$L_brx_$id: .branchtargets", - [(brx_start (i32 imm:$id))]>; - - def BRX_ITEM : - NVPTXInst<(outs), (ins brtarget:$target), - "\t$target,", - [(brx_item bb:$target)]>; - - def BRX_END : - NVPTXInst<(outs), (ins brtarget:$target, Int32Regs:$val, i32imm:$id), - "\t$target;\n\tbrx.idx \t$val, $$L_brx_$id;", - [(brx_end bb:$target, (i32 Int32Regs:$val), (i32 imm:$id))]>; -} - - include "NVPTXIntrinsics.td" //----------------------------------- diff --git a/llvm/test/CodeGen/NVPTX/jump-table.ll b/llvm/test/CodeGen/NVPTX/jump-table.ll deleted file mode 100644 index 867e171a5840ae..00000000000000 --- a/llvm/test/CodeGen/NVPTX/jump-table.ll +++ /dev/null @@ -1,69 +0,0 @@ -; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 -; RUN: llc < %s | FileCheck %s -; RUN: %if ptxas %{ llc < %s | %ptxas-verify %} - -target triple = "nvptx64-nvidia-cuda" - -@out = addrspace(1) global i32 0, align 4 - -define void @foo(i32 %i) { -; CHECK-LABEL: foo( -; CHECK: { -; CHECK-NEXT: .reg .pred %p<2>; -; CHECK-NEXT: .reg .b32 %r<7>; -; CHECK-EMPTY: -; CHECK-NEXT: // %bb.0: // %entry -; CHECK-NEXT: ld.param.u32 %r2, [foo_param_0]; -; CHECK-NEXT: setp.gt.u32 %p1, %r2, 3; -; CHECK-NEXT: @%p1 bra $L__BB0_6; -; CHECK-NEXT: // %bb.1: // %entry -; CHECK-NEXT: $L_brx_0: .branchtargets -; CHECK-NEXT: $L__BB0_2, -; CHECK-NEXT: $L__BB0_3, -; CHECK-NEXT: $L__BB0_4, -; CHECK-NEXT: $L__BB0_5; -; CHECK-NEXT: brx.idx %r2, $L_brx_0; -; CHECK-NEXT: $L__BB0_2: // %case0 -; CHECK-NEXT: mov.b32 %r6, 0; -; CHECK-NEXT: st.global.u32 [out], %r6; -; CHECK-NEXT: bra.uni $L__BB0_6; -; CHECK-NEXT: $L__BB0_4: // %case2 -; CHECK-NEXT: mov.b32 %r4, 2; -; CHECK-NEXT: st.global.u32 [out], %r4; -; CHECK-NEXT: bra.uni $L__BB0_6; -; CHECK-NEXT: $L__BB0_5: // %case3 -; CHECK-NEXT: mov.b32 %r3, 3; -; CHECK-NEXT: st.global.u32 [out], %r3; -; CHECK-NEXT: bra.uni $L__BB0_6; -; CHECK-NEXT: $L__BB0_3: // %case1 -; CHECK-NEXT: mov.b32 %r5, 1; -; CHECK-NEXT: st.global.u32 [out], %r5; -; CHECK-NEXT: $L__BB0_6: // %end -; CHECK-NEXT: ret; -entry: - switch i32 %i, label %end [ - i32 0, label %case0 - i32 1, label %case1 - i32 2, label %case2 - i32 3, label %case3 - ] - -case0: - store i32 0, ptr addrspace(1) @out, align 4 - br label %end - -case1: - store i32 1, ptr addrspace(1) @out, align 4 - br label %end - -case2: - store i32 2, ptr addrspace(1) @out, align 4 - br label %end - -case3: - store i32 3, ptr addrspace(1) @out, align 4 - br label %end - -end: - ret void -} _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
