Skip to content

Commit 3def49c

Browse files
mbrkusaninmbrkusanin
authored
[AMDGPU] Remove s_wakeup_barrier instruction (#122277)
1 parent 1ef2580 commit 3def49c

File tree

11 files changed

+5
-91
lines changed

11 files changed

+5
-91
lines changed

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -489,7 +489,6 @@ TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts")
489489
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts")
490490
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vv*i", "n", "gfx12-insts")
491491
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vv*", "n", "gfx12-insts")
492-
TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vv*", "n", "gfx12-insts")
493492
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "vIs", "n", "gfx12-insts")
494493
TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts")
495494
TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*", "n", "gfx12-insts")

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

Lines changed: 0 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -173,21 +173,6 @@ void test_s_barrier_join(void *bar)
173173
__builtin_amdgcn_s_barrier_join(bar);
174174
}
175175

176-
// CHECK-LABEL: @test_s_wakeup_barrier(
177-
// CHECK-NEXT: entry:
178-
// CHECK-NEXT: [[BAR_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
179-
// CHECK-NEXT: [[BAR_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[BAR_ADDR]] to ptr
180-
// CHECK-NEXT: store ptr [[BAR:%.*]], ptr [[BAR_ADDR_ASCAST]], align 8
181-
// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[BAR_ADDR_ASCAST]], align 8
182-
// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr [[TMP0]] to ptr addrspace(3)
183-
// CHECK-NEXT: call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) [[TMP1]])
184-
// CHECK-NEXT: ret void
185-
//
186-
void test_s_wakeup_barrier(void *bar)
187-
{
188-
__builtin_amdgcn_s_wakeup_barrier(bar);
189-
}
190-
191176
// CHECK-LABEL: @test_s_barrier_leave(
192177
// CHECK-NEXT: entry:
193178
// CHECK-NEXT: call void @llvm.amdgcn.s.barrier.leave(i16 1)

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

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -284,12 +284,6 @@ def int_amdgcn_s_barrier_join : ClangBuiltin<"__builtin_amdgcn_s_barrier_join">,
284284
Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
285285
IntrNoCallback, IntrNoFree]>;
286286

287-
// void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %barrier)
288-
// The %barrier argument must be uniform, otherwise behavior is undefined.
289-
def int_amdgcn_s_wakeup_barrier : ClangBuiltin<"__builtin_amdgcn_s_wakeup_barrier">,
290-
Intrinsic<[], [local_ptr_ty], [IntrNoMem, IntrHasSideEffects, IntrConvergent, IntrWillReturn,
291-
IntrNoCallback, IntrNoFree]>;
292-
293287
// void @llvm.amdgcn.s.barrier.wait(i16 %barrierType)
294288
def int_amdgcn_s_barrier_wait : ClangBuiltin<"__builtin_amdgcn_s_barrier_wait">,
295289
Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrConvergent,

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

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2239,7 +2239,6 @@ bool AMDGPUInstructionSelector::selectG_INTRINSIC_W_SIDE_EFFECTS(
22392239
case Intrinsic::amdgcn_s_barrier_signal_var:
22402240
return selectNamedBarrierInit(I, IntrinsicID);
22412241
case Intrinsic::amdgcn_s_barrier_join:
2242-
case Intrinsic::amdgcn_s_wakeup_barrier:
22432242
case Intrinsic::amdgcn_s_get_named_barrier_state:
22442243
return selectNamedBarrierInst(I, IntrinsicID);
22452244
case Intrinsic::amdgcn_s_get_barrier_state:
@@ -5838,8 +5837,6 @@ unsigned getNamedBarrierOp(bool HasInlineConst, Intrinsic::ID IntrID) {
58385837
llvm_unreachable("not a named barrier op");
58395838
case Intrinsic::amdgcn_s_barrier_join:
58405839
return AMDGPU::S_BARRIER_JOIN_IMM;
5841-
case Intrinsic::amdgcn_s_wakeup_barrier:
5842-
return AMDGPU::S_WAKEUP_BARRIER_IMM;
58435840
case Intrinsic::amdgcn_s_get_named_barrier_state:
58445841
return AMDGPU::S_GET_BARRIER_STATE_IMM;
58455842
};
@@ -5849,8 +5846,6 @@ unsigned getNamedBarrierOp(bool HasInlineConst, Intrinsic::ID IntrID) {
58495846
llvm_unreachable("not a named barrier op");
58505847
case Intrinsic::amdgcn_s_barrier_join:
58515848
return AMDGPU::S_BARRIER_JOIN_M0;
5852-
case Intrinsic::amdgcn_s_wakeup_barrier:
5853-
return AMDGPU::S_WAKEUP_BARRIER_M0;
58545849
case Intrinsic::amdgcn_s_get_named_barrier_state:
58555850
return AMDGPU::S_GET_BARRIER_STATE_M0;
58565851
};

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -326,7 +326,6 @@ bool isReallyAClobber(const Value *Ptr, MemoryDef *Def, AAResults *AA) {
326326
case Intrinsic::amdgcn_s_barrier_wait:
327327
case Intrinsic::amdgcn_s_barrier_leave:
328328
case Intrinsic::amdgcn_s_get_barrier_state:
329-
case Intrinsic::amdgcn_s_wakeup_barrier:
330329
case Intrinsic::amdgcn_wave_barrier:
331330
case Intrinsic::amdgcn_sched_barrier:
332331
case Intrinsic::amdgcn_sched_group_barrier:

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

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3304,7 +3304,6 @@ void AMDGPURegisterBankInfo::applyMappingImpl(
33043304
constrainOpWithReadfirstlane(B, MI, 1);
33053305
return;
33063306
case Intrinsic::amdgcn_s_barrier_join:
3307-
case Intrinsic::amdgcn_s_wakeup_barrier:
33083307
constrainOpWithReadfirstlane(B, MI, 1);
33093308
return;
33103309
case Intrinsic::amdgcn_s_barrier_init:
@@ -5272,7 +5271,6 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
52725271
OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
52735272
break;
52745273
case Intrinsic::amdgcn_s_barrier_join:
5275-
case Intrinsic::amdgcn_s_wakeup_barrier:
52765274
OpdsMapping[1] = getSGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
52775275
break;
52785276
case Intrinsic::amdgcn_s_barrier_init:

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

Lines changed: 5 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -10107,8 +10107,7 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
1010710107
auto *NewMI = DAG.getMachineNode(Opc, DL, Op->getVTList(), Ops);
1010810108
return SDValue(NewMI, 0);
1010910109
}
10110-
case Intrinsic::amdgcn_s_barrier_join:
10111-
case Intrinsic::amdgcn_s_wakeup_barrier: {
10110+
case Intrinsic::amdgcn_s_barrier_join: {
1011210111
// these three intrinsics have one operand: barrier pointer
1011310112
SDValue Chain = Op->getOperand(0);
1011410113
SmallVector<SDValue, 2> Ops;
@@ -10117,32 +10116,16 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
1011710116

1011810117
if (isa<ConstantSDNode>(BarOp)) {
1011910118
uint64_t BarVal = cast<ConstantSDNode>(BarOp)->getZExtValue();
10120-
switch (IntrinsicID) {
10121-
default:
10122-
return SDValue();
10123-
case Intrinsic::amdgcn_s_barrier_join:
10124-
Opc = AMDGPU::S_BARRIER_JOIN_IMM;
10125-
break;
10126-
case Intrinsic::amdgcn_s_wakeup_barrier:
10127-
Opc = AMDGPU::S_WAKEUP_BARRIER_IMM;
10128-
break;
10129-
}
10119+
Opc = AMDGPU::S_BARRIER_JOIN_IMM;
10120+
1013010121
// extract the BarrierID from bits 4-9 of the immediate
1013110122
unsigned BarID = (BarVal >> 4) & 0x3F;
1013210123
SDValue K = DAG.getTargetConstant(BarID, DL, MVT::i32);
1013310124
Ops.push_back(K);
1013410125
Ops.push_back(Chain);
1013510126
} else {
10136-
switch (IntrinsicID) {
10137-
default:
10138-
return SDValue();
10139-
case Intrinsic::amdgcn_s_barrier_join:
10140-
Opc = AMDGPU::S_BARRIER_JOIN_M0;
10141-
break;
10142-
case Intrinsic::amdgcn_s_wakeup_barrier:
10143-
Opc = AMDGPU::S_WAKEUP_BARRIER_M0;
10144-
break;
10145-
}
10127+
Opc = AMDGPU::S_BARRIER_JOIN_M0;
10128+
1014610129
// extract the BarrierID from bits 4-9 of BarOp, copy to M0[5:0]
1014710130
SDValue M0Val;
1014810131
M0Val = DAG.getNode(ISD::SRL, DL, MVT::i32, BarOp,

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

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -488,11 +488,6 @@ def S_BARRIER_JOIN_M0 : SOP1_Pseudo <"s_barrier_join m0", (outs), (ins),
488488
let isConvergent = 1;
489489
}
490490

491-
def S_WAKEUP_BARRIER_M0 : SOP1_Pseudo <"s_wakeup_barrier m0", (outs), (ins),
492-
"", []>{
493-
let SchedRW = [WriteBarrier];
494-
let isConvergent = 1;
495-
}
496491
} // End Uses = [M0]
497492

498493
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),
514509
let isConvergent = 1;
515510
}
516511

517-
def S_WAKEUP_BARRIER_IMM : SOP1_Pseudo <"s_wakeup_barrier", (outs),
518-
(ins SplitBarrier:$src0), "$src0", []>{
519-
let SchedRW = [WriteBarrier];
520-
let isConvergent = 1;
521-
}
522512
} // End has_sdst = 0
523513

524514
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>;
20922082
defm S_GET_BARRIER_STATE_M0 : SOP1_M0_Real_gfx12<0x050>;
20932083
defm S_BARRIER_INIT_M0 : SOP1_M0_Real_gfx12<0x051>;
20942084
defm S_BARRIER_JOIN_M0 : SOP1_M0_Real_gfx12<0x052>;
2095-
defm S_WAKEUP_BARRIER_M0 : SOP1_M0_Real_gfx12<0x057>;
20962085
defm S_BARRIER_SIGNAL_IMM : SOP1_IMM_Real_gfx12<0x04e>;
20972086
defm S_BARRIER_SIGNAL_ISFIRST_IMM : SOP1_IMM_Real_gfx12<0x04f>;
20982087
defm S_GET_BARRIER_STATE_IMM : SOP1_IMM_Real_gfx12<0x050>;
20992088
defm S_BARRIER_INIT_IMM : SOP1_IMM_Real_gfx12<0x051>;
21002089
defm S_BARRIER_JOIN_IMM : SOP1_IMM_Real_gfx12<0x052>;
2101-
defm S_WAKEUP_BARRIER_IMM : SOP1_IMM_Real_gfx12<0x057>;
21022090
defm S_SLEEP_VAR : SOP1_IMM_Real_gfx12<0x058>;
21032091

21042092
//===----------------------------------------------------------------------===//

‎llvm/test/CodeGen/AMDGPU/s-barrier.ll

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -112,10 +112,6 @@ define amdgpu_kernel void @kernel1(ptr addrspace(1) %out, ptr addrspace(3) %in)
112112
; GFX12-SDAG-NEXT: s_mov_b32 m0, 2
113113
; GFX12-SDAG-NEXT: s_barrier_wait 1
114114
; GFX12-SDAG-NEXT: s_barrier_leave
115-
; GFX12-SDAG-NEXT: s_wakeup_barrier m0
116-
; GFX12-SDAG-NEXT: s_mov_b32 m0, s2
117-
; GFX12-SDAG-NEXT: s_wakeup_barrier m0
118-
; GFX12-SDAG-NEXT: s_mov_b32 m0, 2
119115
; GFX12-SDAG-NEXT: s_get_barrier_state s3, m0
120116
; GFX12-SDAG-NEXT: s_mov_b32 m0, s2
121117
; 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)
176172
; GFX12-GISEL-NEXT: s_barrier_join m0
177173
; GFX12-GISEL-NEXT: s_barrier_wait 1
178174
; GFX12-GISEL-NEXT: s_barrier_leave
179-
; GFX12-GISEL-NEXT: s_wakeup_barrier 2
180-
; GFX12-GISEL-NEXT: s_wakeup_barrier m0
181175
; GFX12-GISEL-NEXT: s_get_barrier_state s0, 2
182176
; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
183177
; 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)
218212
call void @llvm.amdgcn.s.barrier.join(ptr addrspace(3) %in)
219213
call void @llvm.amdgcn.s.barrier.wait(i16 1)
220214
call void @llvm.amdgcn.s.barrier.leave(i16 1)
221-
call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) @bar)
222-
call void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3) %in)
223215
%state = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) @bar)
224216
%state2 = call i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3) %in)
225217
call void @llvm.amdgcn.s.barrier()
@@ -295,7 +287,6 @@ declare i1 @llvm.amdgcn.s.barrier.signal.isfirst(i32) #1
295287
declare void @llvm.amdgcn.s.barrier.init(ptr addrspace(3), i32) #1
296288
declare void @llvm.amdgcn.s.barrier.join(ptr addrspace(3)) #1
297289
declare void @llvm.amdgcn.s.barrier.leave(i16) #1
298-
declare void @llvm.amdgcn.s.wakeup.barrier(ptr addrspace(3)) #1
299290
declare i32 @llvm.amdgcn.s.get.barrier.state(i32) #1
300291
declare i32 @llvm.amdgcn.s.get.named.barrier.state(ptr addrspace(3)) #1
301292

‎llvm/test/MC/AMDGPU/gfx12_asm_sop1.s

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -726,15 +726,6 @@ s_barrier_join -2
726726
s_barrier_join m0
727727
// GFX12: encoding: [0x7d,0x52,0x80,0xbe]
728728

729-
s_wakeup_barrier 1
730-
// GFX12: encoding: [0x81,0x57,0x80,0xbe]
731-
732-
s_wakeup_barrier -1
733-
// GFX12: encoding: [0xc1,0x57,0x80,0xbe]
734-
735-
s_wakeup_barrier m0
736-
// GFX12: encoding: [0x7d,0x57,0x80,0xbe]
737-
738729
s_get_barrier_state s3, -1
739730
// GFX12: encoding: [0xc1,0x50,0x83,0xbe]
740731

0 commit comments

Comments
 (0)