Skip to content

Commit 63d7ca9

Browse files
jayfoadjayfoad
authored
[AMDGPU] Add GFX12 llvm.amdgcn.s.wait.*cnt intrinsics (#78723)
1 parent 920bb54 commit 63d7ca9

File tree

3 files changed

+119
-7
lines changed

3 files changed

+119
-7
lines changed

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

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -311,6 +311,17 @@ def int_amdgcn_iglp_opt : ClangBuiltin<"__builtin_amdgcn_iglp_opt">,
311311
def int_amdgcn_s_waitcnt : ClangBuiltin<"__builtin_amdgcn_s_waitcnt">,
312312
Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
313313

314+
// GFX12 intrinsics
315+
class AMDGPUWaitIntrinsic :
316+
Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
317+
def int_amdgcn_s_wait_bvhcnt : AMDGPUWaitIntrinsic;
318+
def int_amdgcn_s_wait_dscnt : AMDGPUWaitIntrinsic;
319+
def int_amdgcn_s_wait_expcnt : AMDGPUWaitIntrinsic;
320+
def int_amdgcn_s_wait_kmcnt : AMDGPUWaitIntrinsic;
321+
def int_amdgcn_s_wait_loadcnt : AMDGPUWaitIntrinsic;
322+
def int_amdgcn_s_wait_samplecnt : AMDGPUWaitIntrinsic;
323+
def int_amdgcn_s_wait_storecnt : AMDGPUWaitIntrinsic;
324+
314325
def int_amdgcn_div_scale : DefaultAttrsIntrinsic<
315326
// 1st parameter: Numerator
316327
// 2nd parameter: Denominator

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

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1720,23 +1720,30 @@ let SubtargetPredicate = HasVGPRSingleUseHintInsts in {
17201720

17211721
let SubtargetPredicate = isGFX12Plus, hasSideEffects = 1 in {
17221722
def S_WAIT_LOADCNT :
1723-
SOPP_Pseudo<"s_wait_loadcnt", (ins s16imm:$simm16), "$simm16">;
1723+
SOPP_Pseudo<"s_wait_loadcnt", (ins s16imm:$simm16), "$simm16",
1724+
[(int_amdgcn_s_wait_loadcnt timm:$simm16)]>;
17241725
def S_WAIT_LOADCNT_DSCNT :
17251726
SOPP_Pseudo<"s_wait_loadcnt_dscnt", (ins s16imm:$simm16), "$simm16">;
17261727
def S_WAIT_STORECNT :
1727-
SOPP_Pseudo<"s_wait_storecnt", (ins s16imm:$simm16), "$simm16">;
1728+
SOPP_Pseudo<"s_wait_storecnt", (ins s16imm:$simm16), "$simm16",
1729+
[(int_amdgcn_s_wait_storecnt timm:$simm16)]>;
17281730
def S_WAIT_STORECNT_DSCNT :
17291731
SOPP_Pseudo<"s_wait_storecnt_dscnt", (ins s16imm:$simm16), "$simm16">;
17301732
def S_WAIT_SAMPLECNT :
1731-
SOPP_Pseudo<"s_wait_samplecnt", (ins s16imm:$simm16), "$simm16">;
1733+
SOPP_Pseudo<"s_wait_samplecnt", (ins s16imm:$simm16), "$simm16",
1734+
[(int_amdgcn_s_wait_samplecnt timm:$simm16)]>;
17321735
def S_WAIT_BVHCNT :
1733-
SOPP_Pseudo<"s_wait_bvhcnt", (ins s16imm:$simm16), "$simm16">;
1736+
SOPP_Pseudo<"s_wait_bvhcnt", (ins s16imm:$simm16), "$simm16",
1737+
[(int_amdgcn_s_wait_bvhcnt timm:$simm16)]>;
17341738
def S_WAIT_EXPCNT :
1735-
SOPP_Pseudo<"s_wait_expcnt", (ins s16imm:$simm16), "$simm16">;
1739+
SOPP_Pseudo<"s_wait_expcnt", (ins s16imm:$simm16), "$simm16",
1740+
[(int_amdgcn_s_wait_expcnt timm:$simm16)]>;
17361741
def S_WAIT_DSCNT :
1737-
SOPP_Pseudo<"s_wait_dscnt", (ins s16imm:$simm16), "$simm16">;
1742+
SOPP_Pseudo<"s_wait_dscnt", (ins s16imm:$simm16), "$simm16",
1743+
[(int_amdgcn_s_wait_dscnt timm:$simm16)]>;
17381744
def S_WAIT_KMCNT :
1739-
SOPP_Pseudo<"s_wait_kmcnt", (ins s16imm:$simm16), "$simm16">;
1745+
SOPP_Pseudo<"s_wait_kmcnt", (ins s16imm:$simm16), "$simm16",
1746+
[(int_amdgcn_s_wait_kmcnt timm:$simm16)]>;
17401747
} // End SubtargetPredicate = isGFX12Plus, hasSideEffects = 1
17411748

17421749
//===----------------------------------------------------------------------===//

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.gfx12.ll

Lines changed: 94 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,94 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
2+
; RUN: llc -global-isel=0 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck %s -check-prefix=GFX12
3+
; RUN: llc -global-isel=1 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck %s -check-prefix=GFX12
4+
5+
define amdgpu_ps void @test_bvhcnt() {
6+
; GFX12-LABEL: test_bvhcnt:
7+
; GFX12: ; %bb.0:
8+
; GFX12-NEXT: s_wait_bvhcnt 0x0
9+
; GFX12-NEXT: s_endpgm
10+
call void @llvm.amdgcn.s.wait.bvhcnt(i16 0)
11+
ret void
12+
}
13+
14+
define amdgpu_ps void @test_dscnt() {
15+
; GFX12-LABEL: test_dscnt:
16+
; GFX12: ; %bb.0:
17+
; GFX12-NEXT: s_wait_dscnt 0x0
18+
; GFX12-NEXT: s_endpgm
19+
call void @llvm.amdgcn.s.wait.dscnt(i16 0)
20+
ret void
21+
}
22+
23+
define amdgpu_ps void @test_expcnt() {
24+
; GFX12-LABEL: test_expcnt:
25+
; GFX12: ; %bb.0:
26+
; GFX12-NEXT: s_wait_expcnt 0x0
27+
; GFX12-NEXT: s_endpgm
28+
call void @llvm.amdgcn.s.wait.expcnt(i16 0)
29+
ret void
30+
}
31+
32+
define amdgpu_ps void @test_kmcnt() {
33+
; GFX12-LABEL: test_kmcnt:
34+
; GFX12: ; %bb.0:
35+
; GFX12-NEXT: s_wait_kmcnt 0x0
36+
; GFX12-NEXT: s_endpgm
37+
call void @llvm.amdgcn.s.wait.kmcnt(i16 0)
38+
ret void
39+
}
40+
41+
define amdgpu_ps void @test_loadcnt() {
42+
; GFX12-LABEL: test_loadcnt:
43+
; GFX12: ; %bb.0:
44+
; GFX12-NEXT: s_wait_loadcnt 0x0
45+
; GFX12-NEXT: s_endpgm
46+
call void @llvm.amdgcn.s.wait.loadcnt(i16 0)
47+
ret void
48+
}
49+
50+
define amdgpu_ps void @test_loadcnt_dscnt() {
51+
; GFX12-LABEL: test_loadcnt_dscnt:
52+
; GFX12: ; %bb.0:
53+
; GFX12-NEXT: s_wait_loadcnt_dscnt 0x0
54+
; GFX12-NEXT: s_endpgm
55+
call void @llvm.amdgcn.s.wait.loadcnt(i16 0)
56+
call void @llvm.amdgcn.s.wait.dscnt(i16 0)
57+
ret void
58+
}
59+
60+
define amdgpu_ps void @test_samplecnt() {
61+
; GFX12-LABEL: test_samplecnt:
62+
; GFX12: ; %bb.0:
63+
; GFX12-NEXT: s_wait_samplecnt 0x0
64+
; GFX12-NEXT: s_endpgm
65+
call void @llvm.amdgcn.s.wait.samplecnt(i16 0)
66+
ret void
67+
}
68+
69+
define amdgpu_ps void @test_storecnt() {
70+
; GFX12-LABEL: test_storecnt:
71+
; GFX12: ; %bb.0:
72+
; GFX12-NEXT: s_wait_storecnt 0x0
73+
; GFX12-NEXT: s_endpgm
74+
call void @llvm.amdgcn.s.wait.storecnt(i16 0)
75+
ret void
76+
}
77+
78+
define amdgpu_ps void @test_storecnt_dscnt() {
79+
; GFX12-LABEL: test_storecnt_dscnt:
80+
; GFX12: ; %bb.0:
81+
; GFX12-NEXT: s_wait_storecnt_dscnt 0x0
82+
; GFX12-NEXT: s_endpgm
83+
call void @llvm.amdgcn.s.wait.storecnt(i16 0)
84+
call void @llvm.amdgcn.s.wait.dscnt(i16 0)
85+
ret void
86+
}
87+
88+
declare void @llvm.amdgcn.s.wait.bvhcnt(i16)
89+
declare void @llvm.amdgcn.s.wait.dscnt(i16)
90+
declare void @llvm.amdgcn.s.wait.expcnt(i16)
91+
declare void @llvm.amdgcn.s.wait.kmcnt(i16)
92+
declare void @llvm.amdgcn.s.wait.loadcnt(i16)
93+
declare void @llvm.amdgcn.s.wait.samplecnt(i16)
94+
declare void @llvm.amdgcn.s.wait.storecnt(i16)

0 commit comments

Comments
 (0)