llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-mc Author: Mirko Brkušanin (mbrkusanin) <details> <summary>Changes</summary> --- Full diff: https://github.com/llvm/llvm-project/pull/122277.diff 10 Files Affected: - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (-1) - (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (-6) - (modified) llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp (-5) - (modified) llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp (-1) - (modified) llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp (-2) - (modified) llvm/lib/Target/AMDGPU/SIISelLowering.cpp (+5-22) - (modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (-12) - (modified) llvm/test/CodeGen/AMDGPU/s-barrier.ll (-9) - (modified) llvm/test/MC/AMDGPU/gfx12_asm_sop1.s (-9) - (modified) llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_sop1.txt (-9) ``````````diff diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 14c1746716cdd6..1b29a8e359c205 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -489,7 +489,6 @@ TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vv*i", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vv*", "n", "gfx12-insts") -TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vv*", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "vIs", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts") TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*", "n", "gfx12-insts") diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 92418b9104ad14..b930d6983e2251 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -284,12 +284,6 @@ def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">, Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>; -// void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %barrier) -// The %barrier argument must be uniform, otherwise behavior is undefined. -def int_amdgcn_s_wakeup_barrier : ClangBuiltin<"__builtin_amdgcn_s_wakeup_barrier">, - Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn, - IntrNoCallback, IntrNoFree]>; - // void @llvm.amdgcn.s.barrier.wait(i16 %barrierType) def int_amdgcn_s_barrier_wait : ClangBuiltin<"__builtin_amdgcn_s_barrier_wait">, Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent, diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp index 041b9b4d66f63f..50e0faef9e7c27 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp @@ -2239,7 +2239,6 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS( case Intrinsic::amdgcn_s_barrier_signal_var: return selectNamedBarrierInit(I, IntrinsicID); case Intrinsic::amdgcn_s_barrier_join: - case Intrinsic::amdgcn_s_wakeup_barrier: case Intrinsic::amdgcn_s_get_named_barrier_state: return selectNamedBarrierInst(I, IntrinsicID); case Intrinsic::amdgcn_s_get_barrier_state: @@ -5839,8 +5838,6 @@ unsigned getNamedBarrierOp(bool HasInlineConst, Intrinsic::ID IntrID) { llvm_unreachable("not a named barrier op"); case Intrinsic::amdgcn_s_barrier_join: return AMDGPU::S_BARRIER_JOIN_IMM; - case Intrinsic::amdgcn_s_wakeup_barrier: - return AMDGPU::S_WAKEUP_BARRIER_IMM; case Intrinsic::amdgcn_s_get_named_barrier_state: return AMDGPU::S_GET_BARRIER_STATE_IMM; }; @@ -5850,8 +5847,6 @@ unsigned getNamedBarrierOp(bool HasInlineConst, Intrinsic::ID IntrID) { llvm_unreachable("not a named barrier op"); case Intrinsic::amdgcn_s_barrier_join: return AMDGPU::S_BARRIER_JOIN_M0; - case Intrinsic::amdgcn_s_wakeup_barrier: - return AMDGPU::S_WAKEUP_BARRIER_M0; case Intrinsic::amdgcn_s_get_named_barrier_state: return AMDGPU::S_GET_BARRIER_STATE_M0; }; diff --git a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp index 2df068d8fb007b..0406ba9c68ccd3 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp @@ -326,7 +326,6 @@ bool isReallyAClobber(const Value *Ptr, MemoryDef *Def, AAResults *AA) { case Intrinsic::amdgcn_s_barrier_wait: case Intrinsic::amdgcn_s_barrier_leave: case Intrinsic::amdgcn_s_get_barrier_state: - case Intrinsic::amdgcn_s_wakeup_barrier: case Intrinsic::amdgcn_wave_barrier: case Intrinsic::amdgcn_sched_barrier: case Intrinsic::amdgcn_sched_group_barrier: diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp index 08e23cbf34e42b..224c368cff4a1f 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -3304,7 +3304,6 @@ void AMDGPURegisterBankInfo::applyMappingImpl( constrainOpWithReadfirstlane(B, MI, 1); return; case Intrinsic::amdgcn_s_barrier_join: - case Intrinsic::amdgcn_s_wakeup_barrier: constrainOpWithReadfirstlane(B, MI, 1); return; case Intrinsic::amdgcn_s_barrier_init: @@ -5272,7 +5271,6 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI); break; case Intrinsic::amdgcn_s_barrier_join: - case Intrinsic::amdgcn_s_wakeup_barrier: OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI); break; case Intrinsic::amdgcn_s_barrier_init: diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index 0ac84f4e1f02af..7e07c2234a132e 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -10076,8 +10076,7 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops); return SDValue(NewMI, 0); } - case Intrinsic::amdgcn_s_barrier_join: - case Intrinsic::amdgcn_s_wakeup_barrier: { + case Intrinsic::amdgcn_s_barrier_join: { // these three intrinsics have one operand: barrier pointer SDValue Chain = Op->getOperand(0); SmallVector<SDValue, 2> Ops; @@ -10086,32 +10085,16 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, if (isa<ConstantSDNode>(BarOp)) { uint64_t BarVal = cast<ConstantSDNode>(BarOp)->getZExtValue(); - switch (IntrinsicID) { - default: - return SDValue(); - case Intrinsic::amdgcn_s_barrier_join: - Opc = AMDGPU::S_BARRIER_JOIN_IMM; - break; - case Intrinsic::amdgcn_s_wakeup_barrier: - Opc = AMDGPU::S_WAKEUP_BARRIER_IMM; - break; - } + Opc = AMDGPU::S_BARRIER_JOIN_IMM; + // extract the BarrierID from bits 4-9 of the immediate unsigned BarID = (BarVal >> 4) & 0x3F; SDValue K = DAG.getTargetConstant(BarID, DL, MVT::i32); Ops.push_back(K); Ops.push_back(Chain); } else { - switch (IntrinsicID) { - default: - return SDValue(); - case Intrinsic::amdgcn_s_barrier_join: - Opc = AMDGPU::S_BARRIER_JOIN_M0; - break; - case Intrinsic::amdgcn_s_wakeup_barrier: - Opc = AMDGPU::S_WAKEUP_BARRIER_M0; - break; - } + Opc = AMDGPU::S_BARRIER_JOIN_M0; + // extract the BarrierID from bits 4-9 of BarOp, copy to M0[5:0] SDValue M0Val; M0Val = DAG.getNode(ISD::SRL, DL, MVT::i32, BarOp, diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td index 46ac2a4992c456..da186f7058d18a 100644 --- a/llvm/lib/Target/AMDGPU/SOPInstructions.td +++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td @@ -488,11 +488,6 @@ def S_BARRIER_JOIN_M0 : SOP1_Pseudo <"s_barrier_join m0", (outs), (ins), let isConvergent = 1; } -def S_WAKEUP_BARRIER_M0 : SOP1_Pseudo <"s_wakeup_barrier m0", (outs), (ins), - "", []>{ - let SchedRW = [WriteBarrier]; - let isConvergent = 1; -} } // End Uses = [M0] def S_BARRIER_SIGNAL_IMM : SOP1_Pseudo <"s_barrier_signal", (outs), @@ -514,11 +509,6 @@ def S_BARRIER_JOIN_IMM : SOP1_Pseudo <"s_barrier_join", (outs), let isConvergent = 1; } -def S_WAKEUP_BARRIER_IMM : SOP1_Pseudo <"s_wakeup_barrier", (outs), - (ins SplitBarrier:$src0), "$src0", []>{ - let SchedRW = [WriteBarrier]; - let isConvergent = 1; -} } // End has_sdst = 0 def S_GET_BARRIER_STATE_IMM : SOP1_Pseudo <"s_get_barrier_state", (outs SSrc_b32:$sdst), @@ -2092,13 +2082,11 @@ defm S_BARRIER_SIGNAL_ISFIRST_M0 : SOP1_M0_Real_gfx12<0x04f>; defm S_GET_BARRIER_STATE_M0 : SOP1_M0_Real_gfx12<0x050>; defm S_BARRIER_INIT_M0 : SOP1_M0_Real_gfx12<0x051>; defm S_BARRIER_JOIN_M0 : SOP1_M0_Real_gfx12<0x052>; -defm S_WAKEUP_BARRIER_M0 : SOP1_M0_Real_gfx12<0x057>; defm S_BARRIER_SIGNAL_IMM : SOP1_IMM_Real_gfx12<0x04e>; defm S_BARRIER_SIGNAL_ISFIRST_IMM : SOP1_IMM_Real_gfx12<0x04f>; defm S_GET_BARRIER_STATE_IMM : SOP1_IMM_Real_gfx12<0x050>; defm S_BARRIER_INIT_IMM : SOP1_IMM_Real_gfx12<0x051>; defm S_BARRIER_JOIN_IMM : SOP1_IMM_Real_gfx12<0x052>; -defm S_WAKEUP_BARRIER_IMM : SOP1_IMM_Real_gfx12<0x057>; defm S_SLEEP_VAR : SOP1_IMM_Real_gfx12<0x058>; //===----------------------------------------------------------------------===// diff --git a/llvm/test/CodeGen/AMDGPU/s-barrier.ll b/llvm/test/CodeGen/AMDGPU/s-barrier.ll index a2624e5f61307a..83a077f7f74db1 100644 --- a/llvm/test/CodeGen/AMDGPU/s-barrier.ll +++ b/llvm/test/CodeGen/AMDGPU/s-barrier.ll @@ -112,10 +112,6 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX12-SDAG-NEXT: s_mov_b32 m0, 2 ; GFX12-SDAG-NEXT: s_barrier_wait 1 ; GFX12-SDAG-NEXT: s_barrier_leave -; GFX12-SDAG-NEXT: s_wakeup_barrier m0 -; GFX12-SDAG-NEXT: s_mov_b32 m0, s2 -; GFX12-SDAG-NEXT: s_wakeup_barrier m0 -; GFX12-SDAG-NEXT: s_mov_b32 m0, 2 ; GFX12-SDAG-NEXT: s_get_barrier_state s3, m0 ; GFX12-SDAG-NEXT: s_mov_b32 m0, s2 ; GFX12-SDAG-NEXT: s_get_barrier_state s2, m0 @@ -176,8 +172,6 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) ; GFX12-GISEL-NEXT: s_barrier_join m0 ; GFX12-GISEL-NEXT: s_barrier_wait 1 ; GFX12-GISEL-NEXT: s_barrier_leave -; GFX12-GISEL-NEXT: s_wakeup_barrier 2 -; GFX12-GISEL-NEXT: s_wakeup_barrier m0 ; GFX12-GISEL-NEXT: s_get_barrier_state s0, 2 ; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0 ; GFX12-GISEL-NEXT: s_get_barrier_state s0, m0 @@ -218,8 +212,6 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in) call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) %in) call void @llvm.amdgcn.s.barrier.wait(i16 1) call void @llvm.amdgcn.s.barrier.leave(i16 1) - call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) @bar) - call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %in) %state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar) %state2 = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) %in) call void @llvm.amdgcn.s.barrier() @@ -295,7 +287,6 @@ declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1 declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1 declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1 declare void @llvm.amdgcn.s.barrier.leave(i16) #1 -declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1 declare i32 @llvm.amdgcn.s.get.barrier.state(i32) #1 declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1 diff --git a/llvm/test/MC/AMDGPU/gfx12_asm_sop1.s b/llvm/test/MC/AMDGPU/gfx12_asm_sop1.s index 939320e9ef2dce..d93ea2e82c1d28 100644 --- a/llvm/test/MC/AMDGPU/gfx12_asm_sop1.s +++ b/llvm/test/MC/AMDGPU/gfx12_asm_sop1.s @@ -726,15 +726,6 @@ s_barrier_join -2 s_barrier_join m0 // GFX12: encoding: [0x7d,0x52,0x80,0xbe] -s_wakeup_barrier 1 -// GFX12: encoding: [0x81,0x57,0x80,0xbe] - -s_wakeup_barrier -1 -// GFX12: encoding: [0xc1,0x57,0x80,0xbe] - -s_wakeup_barrier m0 -// GFX12: encoding: [0x7d,0x57,0x80,0xbe] - s_get_barrier_state s3, -1 // GFX12: encoding: [0xc1,0x50,0x83,0xbe] diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_sop1.txt b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_sop1.txt index f8c235f77b5f56..2cb6da42213e32 100644 --- a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_sop1.txt +++ b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_sop1.txt @@ -726,15 +726,6 @@ # GFX12: s_barrier_join m0 ; encoding: [0x7d,0x52,0x80,0xbe] 0x7d,0x52,0x80,0xbe -# GFX12: s_wakeup_barrier 1 ; encoding: [0x81,0x57,0x80,0xbe] -0x81,0x57,0x80,0xbe - -# GFX12: s_wakeup_barrier -1 ; encoding: [0xc1,0x57,0x80,0xbe] -0xc1,0x57,0x80,0xbe - -# GFX12: s_wakeup_barrier m0 ; encoding: [0x7d,0x57,0x80,0xbe] -0x7d,0x57,0x80,0xbe - # GFX12: s_get_barrier_state s3, -1 ; encoding: [0xc1,0x50,0x83,0xbe] 0xc1,0x50,0x83,0xbe `````````` </details> https://github.com/llvm/llvm-project/pull/122277 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
