Skip to content

Commit 06b67c0

Browse files
mariusz-sikora-at-amdjph-13
mariusz-sikora-at-amd
authored and
jph-13
committed
[AMDGPU] Remove unused s_barrier_{init,join,leave} instructions (llvm#129548)
1 parent 2672268 commit 06b67c0

19 files changed

+3
-624
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -487,9 +487,6 @@ TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts")
487487
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vv*i", "n", "gfx12-insts")
488488
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts")
489489
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts")
490-
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vv*i", "n", "gfx12-insts")
491-
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vv*", "n", "gfx12-insts")
492-
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "vIs", "n", "gfx12-insts")
493490
TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts")
494491
TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*", "n", "gfx12-insts")
495492
TARGET_BUILTIN(__builtin_amdgcn_s_prefetch_data, "vvC*Ui", "nc", "gfx12-insts")

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12-param-err.cl

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -23,13 +23,6 @@ kernel void builtins_amdgcn_s_barrier_signal_isfirst_err(global int* in, global
2323
*out = *in;
2424
}
2525

26-
kernel void builtins_amdgcn_s_barrier_leave_err(global int* in, global int* out, int barrier) {
27-
28-
__builtin_amdgcn_s_barrier_signal(-1);
29-
__builtin_amdgcn_s_barrier_leave(barrier); // expected-error {{'__builtin_amdgcn_s_barrier_leave' must be a constant integer}}
30-
*out = *in;
31-
}
32-
3326
void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, unsigned int off)
3427
{
3528
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, off, 31); // expected-error {{'__builtin_amdgcn_s_buffer_prefetch_data' must be a constant integer}}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl

Lines changed: 0 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -139,50 +139,6 @@ void test_s_barrier_signal_isfirst(int* a, int* b, int *c)
139139
__builtin_amdgcn_s_barrier_wait(1);
140140
}
141141

142-
// CHECK-LABEL: @test_s_barrier_init(
143-
// CHECK-NEXT: entry:
144-
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
145-
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
146-
// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr
147-
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
148-
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8
149-
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
150-
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8
151-
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
152-
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
153-
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) [[TMP1]], i32 [[TMP2]])
154-
// CHECK-NEXT: ret void
155-
//
156-
void test_s_barrier_init(void *bar, int a)
157-
{
158-
__builtin_amdgcn_s_barrier_init(bar, a);
159-
}
160-
161-
// CHECK-LABEL: @test_s_barrier_join(
162-
// CHECK-NEXT: entry:
163-
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
164-
// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr
165-
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8
166-
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8
167-
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
168-
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) [[TMP1]])
169-
// CHECK-NEXT: ret void
170-
//
171-
void test_s_barrier_join(void *bar)
172-
{
173-
__builtin_amdgcn_s_barrier_join(bar);
174-
}
175-
176-
// CHECK-LABEL: @test_s_barrier_leave(
177-
// CHECK-NEXT: entry:
178-
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.leave(i16 1)
179-
// CHECK-NEXT: ret void
180-
//
181-
void test_s_barrier_leave()
182-
{
183-
__builtin_amdgcn_s_barrier_leave(1);
184-
}
185-
186142
// CHECK-LABEL: @test_s_get_barrier_state(
187143
// CHECK-NEXT: entry:
188144
// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 0 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -272,28 +272,11 @@ def int_amdgcn_s_barrier_signal_isfirst : ClangBuiltin<"__builtin_amdgcn_s_barri
272272
Intrinsic<[llvm_i1_ty], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,
273273
IntrWillReturn, IntrNoCallback, IntrNoFree]>;
274274

275-
// void @llvm.amdgcn.s.barrier.init(ptr addrspace(3) %barrier, i32 %memberCnt)
276-
// The %barrier and %memberCnt argument must be uniform, otherwise behavior is undefined.
277-
def int_amdgcn_s_barrier_init : ClangBuiltin<"__builtin_amdgcn_s_barrier_init">,
278-
Intrinsic<[], [local_ptr_ty, llvm_i32_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent,
279-
IntrWillReturn, IntrNoCallback, IntrNoFree]>;
280-
281-
// void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) %barrier)
282-
// The %barrier argument must be uniform, otherwise behavior is undefined.
283-
def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">,
284-
Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
285-
IntrNoCallback, IntrNoFree]>;
286-
287275
// void @llvm.amdgcn.s.barrier.wait(i16 %barrierType)
288276
def int_amdgcn_s_barrier_wait : ClangBuiltin<"__builtin_amdgcn_s_barrier_wait">,
289277
Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,
290278
IntrWillReturn, IntrNoCallback, IntrNoFree]>;
291279

292-
// void @llvm.amdgcn.s.barrier.leave(i16 %barrierType)
293-
def int_amdgcn_s_barrier_leave : ClangBuiltin<"__builtin_amdgcn_s_barrier_leave">,
294-
Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,
295-
IntrWillReturn, IntrNoCallback, IntrNoFree]>;
296-
297280
// uint32_t @llvm.amdgcn.s.get.barrier.state(i32 %barrierId)
298281
// The %barrierType argument must be uniform, otherwise behavior is undefined.
299282
def int_amdgcn_s_get_barrier_state : ClangBuiltin<"__builtin_amdgcn_s_get_barrier_state">,

‎llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

Lines changed: 1 addition & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -2322,10 +2322,8 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS(
23222322
break;
23232323
case Intrinsic::amdgcn_ds_bvh_stack_rtn:
23242324
return selectDSBvhStackIntrinsic(I);
2325-
case Intrinsic::amdgcn_s_barrier_init:
23262325
case Intrinsic::amdgcn_s_barrier_signal_var:
23272326
return selectNamedBarrierInit(I, IntrinsicID);
2328-
case Intrinsic::amdgcn_s_barrier_join:
23292327
case Intrinsic::amdgcn_s_get_named_barrier_state:
23302328
return selectNamedBarrierInst(I, IntrinsicID);
23312329
case Intrinsic::amdgcn_s_get_barrier_state:
@@ -5928,17 +5926,13 @@ unsigned getNamedBarrierOp(bool HasInlineConst, Intrinsic::ID IntrID) {
59285926
switch (IntrID) {
59295927
default:
59305928
llvm_unreachable("not a named barrier op");
5931-
case Intrinsic::amdgcn_s_barrier_join:
5932-
return AMDGPU::S_BARRIER_JOIN_IMM;
59335929
case Intrinsic::amdgcn_s_get_named_barrier_state:
59345930
return AMDGPU::S_GET_BARRIER_STATE_IMM;
59355931
};
59365932
} else {
59375933
switch (IntrID) {
59385934
default:
59395935
llvm_unreachable("not a named barrier op");
5940-
case Intrinsic::amdgcn_s_barrier_join:
5941-
return AMDGPU::S_BARRIER_JOIN_M0;
59425936
case Intrinsic::amdgcn_s_get_named_barrier_state:
59435937
return AMDGPU::S_GET_BARRIER_STATE_M0;
59445938
};
@@ -5989,11 +5983,8 @@ bool AMDGPUInstructionSelector::selectNamedBarrierInit(
59895983
BuildMI(*MBB, &I, DL, TII.get(AMDGPU::COPY), AMDGPU::M0).addReg(TmpReg4);
59905984
constrainSelectedInstRegOperands(*CopyMIB, TII, TRI, RBI);
59915985

5992-
unsigned Opc = IntrID == Intrinsic::amdgcn_s_barrier_init
5993-
? AMDGPU::S_BARRIER_INIT_M0
5994-
: AMDGPU::S_BARRIER_SIGNAL_M0;
59955986
MachineInstrBuilder MIB;
5996-
MIB = BuildMI(*MBB, &I, DL, TII.get(Opc));
5987+
MIB = BuildMI(*MBB, &I, DL, TII.get(AMDGPU::S_BARRIER_SIGNAL_M0));
59975988

59985989
I.eraseFromParent();
59995990
return true;

‎llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.h

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -156,7 +156,6 @@ class AMDGPUInstructionSelector final : public InstructionSelector {
156156
bool selectNamedBarrierInst(MachineInstr &I, Intrinsic::ID IID) const;
157157
bool selectSBarrierSignalIsfirst(MachineInstr &I, Intrinsic::ID IID) const;
158158
bool selectSGetBarrierState(MachineInstr &I, Intrinsic::ID IID) const;
159-
bool selectSBarrierLeave(MachineInstr &I) const;
160159

161160
std::pair<Register, unsigned> selectVOP3ModsImpl(Register Src,
162161
bool IsCanonicalizing = true,

‎llvm/lib/Target/AMDGPU/AMDGPUMemoryUtils.cpp

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -359,10 +359,7 @@ bool isReallyAClobber(const Value *Ptr, MemoryDef *Def, AAResults *AA) {
359359
case Intrinsic::amdgcn_s_barrier_signal:
360360
case Intrinsic::amdgcn_s_barrier_signal_var:
361361
case Intrinsic::amdgcn_s_barrier_signal_isfirst:
362-
case Intrinsic::amdgcn_s_barrier_init:
363-
case Intrinsic::amdgcn_s_barrier_join:
364362
case Intrinsic::amdgcn_s_barrier_wait:
365-
case Intrinsic::amdgcn_s_barrier_leave:
366363
case Intrinsic::amdgcn_s_get_barrier_state:
367364
case Intrinsic::amdgcn_wave_barrier:
368365
case Intrinsic::amdgcn_sched_barrier:

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -3303,10 +3303,6 @@ void AMDGPURegisterBankInfo::applyMappingImpl(
33033303
assert(OpdMapper.getVRegs(1).empty());
33043304
constrainOpWithReadfirstlane(B, MI, 1);
33053305
return;
3306-
case Intrinsic::amdgcn_s_barrier_join:
3307-
constrainOpWithReadfirstlane(B, MI, 1);
3308-
return;
3309-
case Intrinsic::amdgcn_s_barrier_init:
33103306
case Intrinsic::amdgcn_s_barrier_signal_var:
33113307
constrainOpWithReadfirstlane(B, MI, 1);
33123308
constrainOpWithReadfirstlane(B, MI, 2);
@@ -5272,10 +5268,6 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
52725268
case Intrinsic::amdgcn_s_sleep_var:
52735269
OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
52745270
break;
5275-
case Intrinsic::amdgcn_s_barrier_join:
5276-
OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
5277-
break;
5278-
case Intrinsic::amdgcn_s_barrier_init:
52795271
case Intrinsic::amdgcn_s_barrier_signal_var:
52805272
OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
52815273
OpdsMapping[2] = getSGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);

‎llvm/lib/Target/AMDGPU/SIISelLowering.cpp

Lines changed: 2 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -10180,17 +10180,13 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
1018010180
return SDValue(DAG.getMachineNode(AMDGPU::SI_END_CF, DL, MVT::Other,
1018110181
Op->getOperand(2), Chain),
1018210182
0);
10183-
case Intrinsic::amdgcn_s_barrier_init:
1018410183
case Intrinsic::amdgcn_s_barrier_signal_var: {
1018510184
// these two intrinsics have two operands: barrier pointer and member count
1018610185
SDValue Chain = Op->getOperand(0);
1018710186
SmallVector<SDValue, 2> Ops;
1018810187
SDValue BarOp = Op->getOperand(2);
1018910188
SDValue CntOp = Op->getOperand(3);
1019010189
SDValue M0Val;
10191-
unsigned Opc = IntrinsicID == Intrinsic::amdgcn_s_barrier_init
10192-
? AMDGPU::S_BARRIER_INIT_M0
10193-
: AMDGPU::S_BARRIER_SIGNAL_M0;
1019410190
// extract the BarrierID from bits 4-9 of BarOp
1019510191
SDValue BarID;
1019610192
BarID = DAG.getNode(ISD::SRL, DL, MVT::i32, BarOp,
@@ -10214,40 +10210,8 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
1021410210

1021510211
Ops.push_back(copyToM0(DAG, Chain, DL, M0Val).getValue(0));
1021610212

10217-
auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
10218-
return SDValue(NewMI, 0);
10219-
}
10220-
case Intrinsic::amdgcn_s_barrier_join: {
10221-
// these three intrinsics have one operand: barrier pointer
10222-
SDValue Chain = Op->getOperand(0);
10223-
SmallVector<SDValue, 2> Ops;
10224-
SDValue BarOp = Op->getOperand(2);
10225-
unsigned Opc;
10226-
10227-
if (isa<ConstantSDNode>(BarOp)) {
10228-
uint64_t BarVal = cast<ConstantSDNode>(BarOp)->getZExtValue();
10229-
Opc = AMDGPU::S_BARRIER_JOIN_IMM;
10230-
10231-
// extract the BarrierID from bits 4-9 of the immediate
10232-
unsigned BarID = (BarVal >> 4) & 0x3F;
10233-
SDValue K = DAG.getTargetConstant(BarID, DL, MVT::i32);
10234-
Ops.push_back(K);
10235-
Ops.push_back(Chain);
10236-
} else {
10237-
Opc = AMDGPU::S_BARRIER_JOIN_M0;
10238-
10239-
// extract the BarrierID from bits 4-9 of BarOp, copy to M0[5:0]
10240-
SDValue M0Val;
10241-
M0Val = DAG.getNode(ISD::SRL, DL, MVT::i32, BarOp,
10242-
DAG.getShiftAmountConstant(4, MVT::i32, DL));
10243-
M0Val =
10244-
SDValue(DAG.getMachineNode(AMDGPU::S_AND_B32, DL, MVT::i32, M0Val,
10245-
DAG.getTargetConstant(0x3F, DL, MVT::i32)),
10246-
0);
10247-
Ops.push_back(copyToM0(DAG, Chain, DL, M0Val).getValue(0));
10248-
}
10249-
10250-
auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
10213+
auto *NewMI = DAG.getMachineNode(AMDGPU::S_BARRIER_SIGNAL_M0, DL,
10214+
Op->getVTList(), Ops);
1025110215
return SDValue(NewMI, 0);
1025210216
}
1025310217
case Intrinsic::amdgcn_s_prefetch_data: {

‎llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2078,7 +2078,6 @@ void SIInsertWaitcnts::updateEventWaitcntAfter(MachineInstr &Inst,
20782078
case AMDGPU::S_MEMREALTIME:
20792079
case AMDGPU::S_BARRIER_SIGNAL_ISFIRST_M0:
20802080
case AMDGPU::S_BARRIER_SIGNAL_ISFIRST_IMM:
2081-
case AMDGPU::S_BARRIER_LEAVE:
20822081
case AMDGPU::S_GET_BARRIER_STATE_M0:
20832082
case AMDGPU::S_GET_BARRIER_STATE_IMM:
20842083
ScoreBrackets->updateByEvent(TII, TRI, MRI, SMEM_ACCESS, Inst);

‎llvm/lib/Target/AMDGPU/SIInstrInfo.h

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -959,11 +959,6 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
959959

960960
bool isBarrier(unsigned Opcode) const {
961961
return isBarrierStart(Opcode) || Opcode == AMDGPU::S_BARRIER_WAIT ||
962-
Opcode == AMDGPU::S_BARRIER_INIT_M0 ||
963-
Opcode == AMDGPU::S_BARRIER_INIT_IMM ||
964-
Opcode == AMDGPU::S_BARRIER_JOIN_IMM ||
965-
Opcode == AMDGPU::S_BARRIER_LEAVE ||
966-
Opcode == AMDGPU::S_BARRIER_LEAVE_IMM ||
967962
Opcode == AMDGPU::DS_GWS_INIT || Opcode == AMDGPU::DS_GWS_BARRIER;
968963
}
969964

‎llvm/lib/Target/AMDGPU/SOPInstructions.td

Lines changed: 0 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -470,24 +470,6 @@ def S_BARRIER_SIGNAL_ISFIRST_M0 : SOP1_Pseudo <"s_barrier_signal_isfirst m0", (o
470470
let isConvergent = 1;
471471
}
472472

473-
def S_BARRIER_INIT_M0 : SOP1_Pseudo <"s_barrier_init m0", (outs), (ins),
474-
"", []>{
475-
let SchedRW = [WriteBarrier];
476-
let isConvergent = 1;
477-
}
478-
479-
def S_BARRIER_INIT_IMM : SOP1_Pseudo <"s_barrier_init", (outs),
480-
(ins SplitBarrier:$src0), "$src0", []>{
481-
let SchedRW = [WriteBarrier];
482-
let isConvergent = 1;
483-
}
484-
485-
def S_BARRIER_JOIN_M0 : SOP1_Pseudo <"s_barrier_join m0", (outs), (ins),
486-
"", []>{
487-
let SchedRW = [WriteBarrier];
488-
let isConvergent = 1;
489-
}
490-
491473
} // End Uses = [M0]
492474

493475
def S_BARRIER_SIGNAL_IMM : SOP1_Pseudo <"s_barrier_signal", (outs),
@@ -503,12 +485,6 @@ def S_BARRIER_SIGNAL_ISFIRST_IMM : SOP1_Pseudo <"s_barrier_signal_isfirst", (out
503485
let isConvergent = 1;
504486
}
505487

506-
def S_BARRIER_JOIN_IMM : SOP1_Pseudo <"s_barrier_join", (outs),
507-
(ins SplitBarrier:$src0), "$src0", []>{
508-
let SchedRW = [WriteBarrier];
509-
let isConvergent = 1;
510-
}
511-
512488
} // End has_sdst = 0
513489

514490
def S_GET_BARRIER_STATE_IMM : SOP1_Pseudo <"s_get_barrier_state", (outs SSrc_b32:$sdst),
@@ -1594,17 +1570,6 @@ def S_BARRIER_WAIT : SOPP_Pseudo <"s_barrier_wait", (ins i16imm:$simm16), "$simm
15941570
let isConvergent = 1;
15951571
}
15961572

1597-
def S_BARRIER_LEAVE : SOPP_Pseudo <"s_barrier_leave", (ins)> {
1598-
let SchedRW = [WriteBarrier];
1599-
let simm16 = 0;
1600-
let fixed_imm = 1;
1601-
let isConvergent = 1;
1602-
let Defs = [SCC];
1603-
}
1604-
1605-
def S_BARRIER_LEAVE_IMM : SOPP_Pseudo <"s_barrier_leave",
1606-
(ins i16imm:$simm16), "$simm16", [(int_amdgcn_s_barrier_leave timm:$simm16)]>;
1607-
16081573
def S_WAKEUP : SOPP_Pseudo <"s_wakeup", (ins) > {
16091574
let SubtargetPredicate = isGFX8Plus;
16101575
let simm16 = 0;
@@ -2080,13 +2045,9 @@ defm S_SENDMSG_RTN_B64 : SOP1_Real_gfx11_gfx12<0x04d>;
20802045
defm S_BARRIER_SIGNAL_M0 : SOP1_M0_Real_gfx12<0x04e>;
20812046
defm S_BARRIER_SIGNAL_ISFIRST_M0 : SOP1_M0_Real_gfx12<0x04f>;
20822047
defm S_GET_BARRIER_STATE_M0 : SOP1_M0_Real_gfx12<0x050>;
2083-
defm S_BARRIER_INIT_M0 : SOP1_M0_Real_gfx12<0x051>;
2084-
defm S_BARRIER_JOIN_M0 : SOP1_M0_Real_gfx12<0x052>;
20852048
defm S_BARRIER_SIGNAL_IMM : SOP1_IMM_Real_gfx12<0x04e>;
20862049
defm S_BARRIER_SIGNAL_ISFIRST_IMM : SOP1_IMM_Real_gfx12<0x04f>;
20872050
defm S_GET_BARRIER_STATE_IMM : SOP1_IMM_Real_gfx12<0x050>;
2088-
defm S_BARRIER_INIT_IMM : SOP1_IMM_Real_gfx12<0x051>;
2089-
defm S_BARRIER_JOIN_IMM : SOP1_IMM_Real_gfx12<0x052>;
20902051
defm S_SLEEP_VAR : SOP1_IMM_Real_gfx12<0x058>;
20912052

20922053
//===----------------------------------------------------------------------===//
@@ -2563,7 +2524,6 @@ multiclass SOPP_Real_32_gfx12<bits<7> op, string name = !tolower(NAME)> {
25632524
}
25642525

25652526
defm S_BARRIER_WAIT : SOPP_Real_32_gfx12<0x014>;
2566-
defm S_BARRIER_LEAVE : SOPP_Real_32_gfx12<0x015>;
25672527
defm S_WAIT_LOADCNT : SOPP_Real_32_gfx12<0x040>;
25682528
defm S_WAIT_STORECNT : SOPP_Real_32_gfx12<0x041>;
25692529
defm S_WAIT_SAMPLECNT : SOPP_Real_32_gfx12<0x042>;

0 commit comments

Comments
 (0)