[llvm] 63d7ca9 - [AMDGPU] Add GFX12 llvm.amdgcn.s.wait.*cnt intrinsics (#78723)
via llvm-commits
llvm-commits at lists.llvm.org
Sat Jan 20 03:44:47 PST 2024
Author: Jay Foad
Date: 2024-01-20T11:44:42Z
New Revision: 63d7ca924feea1948329ba80c4dc4fad56112be3
URL: https://github.com/llvm/llvm-project/commit/63d7ca924feea1948329ba80c4dc4fad56112be3
DIFF: https://github.com/llvm/llvm-project/commit/63d7ca924feea1948329ba80c4dc4fad56112be3.diff
LOG: [AMDGPU] Add GFX12 llvm.amdgcn.s.wait.*cnt intrinsics (#78723)
Added:
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.wait.gfx12.ll
Modified:
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/lib/Target/AMDGPU/SOPInstructions.td
Removed:
################################################################################
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)
More information about the llvm-commits
mailing list