[AMDGPU] Add GFX12 llvm.amdgcn.s.wait.*cnt intrinsics

[jayfoad]

No description provided.

[llvmbot]

@llvm/pr-subscribers-llvm-ir

@llvm/pr-subscribers-backend-amdgpu

Author: Jay Foad (jayfoad)

Changes

---

Full diff: https://github.com/llvm/llvm-project/pull/78723.diff

3 Files Affected:

• (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+11)
• (modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (+14-7)
• (added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.gfx12.ll (+94)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 9302e590a6fc93..21765cdd13a153 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -311,6 +311,17 @@ def int_amdgcn_iglp_opt : ClangBuiltin<"__builtin_amdgcn_iglp_opt">,
 def int_amdgcn_s_waitcnt : ClangBuiltin<"__builtin_amdgcn_s_waitcnt">,
   Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
 
+// GFX12 intrinsics
+class AMDGPUWaitIntrinsic :
+  Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
+def int_amdgcn_s_wait_bvhcnt         : AMDGPUWaitIntrinsic;
+def int_amdgcn_s_wait_dscnt          : AMDGPUWaitIntrinsic;
+def int_amdgcn_s_wait_expcnt         : AMDGPUWaitIntrinsic;
+def int_amdgcn_s_wait_kmcnt          : AMDGPUWaitIntrinsic;
+def int_amdgcn_s_wait_loadcnt        : AMDGPUWaitIntrinsic;
+def int_amdgcn_s_wait_samplecnt      : AMDGPUWaitIntrinsic;
+def int_amdgcn_s_wait_storecnt       : AMDGPUWaitIntrinsic;
+
 def int_amdgcn_div_scale : DefaultAttrsIntrinsic<
   // 1st parameter: Numerator
   // 2nd parameter: Denominator
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index eae4800ade0dce..d9c6b1a35e272c 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1720,23 +1720,30 @@ let SubtargetPredicate = HasVGPRSingleUseHintInsts in {
 
 let SubtargetPredicate = isGFX12Plus, hasSideEffects = 1 in {
   def S_WAIT_LOADCNT :
-    SOPP_Pseudo<"s_wait_loadcnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_loadcnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_loadcnt timm:$simm16)]>;
   def S_WAIT_LOADCNT_DSCNT :
     SOPP_Pseudo<"s_wait_loadcnt_dscnt", (ins s16imm:$simm16), "$simm16">;
   def S_WAIT_STORECNT :
-    SOPP_Pseudo<"s_wait_storecnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_storecnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_storecnt timm:$simm16)]>;
   def S_WAIT_STORECNT_DSCNT :
     SOPP_Pseudo<"s_wait_storecnt_dscnt", (ins s16imm:$simm16), "$simm16">;
   def S_WAIT_SAMPLECNT :
-    SOPP_Pseudo<"s_wait_samplecnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_samplecnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_samplecnt timm:$simm16)]>;
   def S_WAIT_BVHCNT :
-    SOPP_Pseudo<"s_wait_bvhcnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_bvhcnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_bvhcnt timm:$simm16)]>;
   def S_WAIT_EXPCNT :
-    SOPP_Pseudo<"s_wait_expcnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_expcnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_expcnt timm:$simm16)]>;
   def S_WAIT_DSCNT :
-    SOPP_Pseudo<"s_wait_dscnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_dscnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_dscnt timm:$simm16)]>;
   def S_WAIT_KMCNT :
-    SOPP_Pseudo<"s_wait_kmcnt", (ins s16imm:$simm16), "$simm16">;
+    SOPP_Pseudo<"s_wait_kmcnt", (ins s16imm:$simm16), "$simm16",
+                [(int_amdgcn_s_wait_kmcnt timm:$simm16)]>;
 } // End SubtargetPredicate = isGFX12Plus, hasSideEffects = 1
 
 //===----------------------------------------------------------------------===//
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.gfx12.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.gfx12.ll
new file mode 100644
index 00000000000000..f03dbb9eb16457
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.gfx12.ll
@@ -0,0 +1,94 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -global-isel=0 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck %s -check-prefix=GFX12
+; RUN: llc -global-isel=1 -march=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck %s -check-prefix=GFX12
+
+define amdgpu_ps void @test_bvhcnt() {
+; GFX12-LABEL: test_bvhcnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_bvhcnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.bvhcnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_dscnt() {
+; GFX12-LABEL: test_dscnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_dscnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.dscnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_expcnt() {
+; GFX12-LABEL: test_expcnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_expcnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.expcnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_kmcnt() {
+; GFX12-LABEL: test_kmcnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_kmcnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.kmcnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_loadcnt() {
+; GFX12-LABEL: test_loadcnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_loadcnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.loadcnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_loadcnt_dscnt() {
+; GFX12-LABEL: test_loadcnt_dscnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_loadcnt_dscnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.loadcnt(i16 0)
+  call void @llvm.amdgcn.s.wait.dscnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_samplecnt() {
+; GFX12-LABEL: test_samplecnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_samplecnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.samplecnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_storecnt() {
+; GFX12-LABEL: test_storecnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_storecnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.storecnt(i16 0)
+  ret void
+}
+
+define amdgpu_ps void @test_storecnt_dscnt() {
+; GFX12-LABEL: test_storecnt_dscnt:
+; GFX12:       ; %bb.0:
+; GFX12-NEXT:    s_wait_storecnt_dscnt 0x0
+; GFX12-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.wait.storecnt(i16 0)
+  call void @llvm.amdgcn.s.wait.dscnt(i16 0)
+  ret void
+}
+
+declare void @llvm.amdgcn.s.wait.bvhcnt(i16)
+declare void @llvm.amdgcn.s.wait.dscnt(i16)
+declare void @llvm.amdgcn.s.wait.expcnt(i16)
+declare void @llvm.amdgcn.s.wait.kmcnt(i16)
+declare void @llvm.amdgcn.s.wait.loadcnt(i16)
+declare void @llvm.amdgcn.s.wait.samplecnt(i16)
+declare void @llvm.amdgcn.s.wait.storecnt(i16)

[jayfoad]

This is by analogy with the existing amdgcn.s.waitcnt intrinsic. TBH I don't know if there is a good use case for the new intrinsics.