[llvm-branch-commits] [llvm] [NFC][AMDGPU] Create tables for async-related intrinsics (PR #225488)
Zach Goldthorpe via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Oct 5 09:46:27 PDT 2026
https://github.com/zGoldthorpe updated https://github.com/llvm/llvm-project/pull/225488
>From 8b9957d78bfac1769104a63a06fdf9f525d24210 Mon Sep 17 00:00:00 2001
From: Zach Goldthorpe <Zach.Goldthorpe at amd.com>
Date: Tue, 22 Sep 2026 13:53:34 -0500
Subject: [PATCH 1/2] [NFC][AMDGPU] Create tables for async-related intrinsics
---
.../AMDGPU/AMDGPUInstructionSelector.cpp | 17 +--
.../Target/AMDGPU/AMDGPUSearchableTables.td | 101 ++++++++++++++++++
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 21 +---
.../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 73 ++++++-------
llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 3 +
5 files changed, 147 insertions(+), 68 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
index 140cf58e8fdd3..09544e4b12263 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp
@@ -3521,19 +3521,6 @@ bool AMDGPUInstructionSelector::selectG_INSERT_VECTOR_ELT(
return true;
}
-static bool isAsyncLDSDMA(Intrinsic::ID Intr) {
- switch (Intr) {
- case Intrinsic::amdgcn_raw_buffer_load_async_lds:
- case Intrinsic::amdgcn_raw_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_load_async_to_lds:
- case Intrinsic::amdgcn_global_load_async_lds:
- return true;
- }
- return false;
-}
-
bool AMDGPUInstructionSelector::selectBufferLoadLds(MachineInstr &MI) const {
if (!Subtarget->hasVMemToLDSLoad())
return false;
@@ -3629,7 +3616,7 @@ bool AMDGPUInstructionSelector::selectBufferLoadLds(MachineInstr &MI) const {
Aux & (IsGFX12Plus ? AMDGPU::CPol::SWZ : AMDGPU::CPol::SWZ_pregfx12)
? 1
: 0); // swz
- MIB.addImm(isAsyncLDSDMA(IntrinsicID));
+ MIB.addImm(AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID));
MachineMemOperand *LoadMMO = *MI.memoperands_begin();
// Don't set the offset value here because the pointer points to the base of
@@ -3824,7 +3811,7 @@ bool AMDGPUInstructionSelector::selectGlobalLoadLds(MachineInstr &MI) const{
unsigned Aux = MI.getOperand(5).getImm();
MIB.addImm(Aux & ~AMDGPU::CPol::VIRTUAL_BITS); // cpol
- MIB.addImm(isAsyncLDSDMA(IntrinsicID));
+ MIB.addImm(AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID));
MachineMemOperand *LoadMMO = *MI.memoperands_begin();
MachinePointerInfo LoadPtrI = LoadMMO->getPointerInfo();
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
index 81738d3fdc65a..c69b0b2871f35 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
@@ -429,3 +429,104 @@ def AMDGPUImageDMaskIntrinsicTable : GenericTable {
let PrimaryKeyName = "getAMDGPUImageDMaskIntrinsic";
let PrimaryKeyEarlyOut = 1;
}
+
+// DMA transfers between global memory and LDS.
+class LDSDMAIntrinsic<Intrinsic intr> {
+ Intrinsic Intr = intr;
+}
+
+def LDSDMAIntrinsics : GenericTable {
+ let FilterClass = "LDSDMAIntrinsic";
+ let Fields = ["Intr"];
+
+ let PrimaryKey = ["Intr"];
+ let PrimaryKeyName = "lookupLDSDMAIntrinsic";
+}
+
+def : LDSDMAIntrinsic<int_amdgcn_raw_buffer_load_lds>;
+def : LDSDMAIntrinsic<int_amdgcn_raw_ptr_buffer_load_lds>;
+def : LDSDMAIntrinsic<int_amdgcn_struct_buffer_load_lds>;
+def : LDSDMAIntrinsic<int_amdgcn_struct_ptr_buffer_load_lds>;
+def : LDSDMAIntrinsic<int_amdgcn_load_to_lds>;
+def : LDSDMAIntrinsic<int_amdgcn_global_load_lds>;
+
+// DMA transfers between global memory and LDS tracked with asyncmarks.
+class AsyncLDSDMAIntrinsic<Intrinsic intr> {
+ Intrinsic Intr = intr;
+}
+
+def AsyncLDSDMAIntrinsics : GenericTable {
+ let FilterClass = "AsyncLDSDMAIntrinsic";
+ let Fields = ["Intr"];
+
+ let PrimaryKey = ["Intr"];
+ let PrimaryKeyName = "lookupAsyncLDSDMAIntrinsic";
+}
+
+def : AsyncLDSDMAIntrinsic<int_amdgcn_raw_buffer_load_async_lds>;
+def : AsyncLDSDMAIntrinsic<int_amdgcn_raw_ptr_buffer_load_async_lds>;
+def : AsyncLDSDMAIntrinsic<int_amdgcn_struct_buffer_load_async_lds>;
+def : AsyncLDSDMAIntrinsic<int_amdgcn_struct_ptr_buffer_load_async_lds>;
+def : AsyncLDSDMAIntrinsic<int_amdgcn_load_async_to_lds>;
+def : AsyncLDSDMAIntrinsic<int_amdgcn_global_load_async_lds>;
+
+// Operations tracked by ASYNCcnt.
+class AsyncIntrinsic<Intrinsic intr> {
+ Intrinsic Intr = intr;
+}
+
+def AsyncIntrinsics : GenericTable {
+ let FilterClass = "AsyncIntrinsic";
+ let Fields = ["Intr"];
+
+ let PrimaryKey = ["Intr"];
+ let PrimaryKeyName = "lookupAsyncIntrinsic";
+}
+
+def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b8>;
+def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b32>;
+def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b64>;
+def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b128>;
+def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b8>;
+def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b32>;
+def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b64>;
+def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b128>;
+def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b8>;
+def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b32>;
+def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b64>;
+def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b128>;
+def : AsyncIntrinsic<int_amdgcn_ds_atomic_async_barrier_arrive_b64>;
+def : AsyncIntrinsic<int_amdgcn_s_wait_asynccnt>;
+
+// Operations tracked by TENSORcnt.
+class TensorIntrinsic<Intrinsic intr> {
+ Intrinsic Intr = intr;
+}
+
+def TensorIntrinsics : GenericTable {
+ let FilterClass = "TensorIntrinsic";
+ let Fields = ["Intr"];
+
+ let PrimaryKey = ["Intr"];
+ let PrimaryKeyName = "lookupTensorIntrinsic";
+}
+
+def : TensorIntrinsic<int_amdgcn_tensor_load_to_lds>;
+def : TensorIntrinsic<int_amdgcn_tensor_store_from_lds>;
+def : TensorIntrinsic<int_amdgcn_s_wait_tensorcnt>;
+
+// Asyncmark production and consumption.
+class AsyncMarkIntrinsic<Intrinsic intr> {
+ Intrinsic Intr = intr;
+}
+
+def AsyncMarkIntrinsics : GenericTable {
+ let FilterClass = "AsyncMarkIntrinsic";
+ let Fields = ["Intr"];
+
+ let PrimaryKey = ["Intr"];
+ let PrimaryKeyName = "lookupAsyncMarkIntrinsic";
+}
+
+def : AsyncMarkIntrinsic<int_amdgcn_asyncmark>;
+def : AsyncMarkIntrinsic<int_amdgcn_wait_asyncmark>;
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index f9a4ba4c1bd1d..4d04e98bf408a 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -12672,19 +12672,6 @@ SDValue SITargetLowering::handleD16VData(SDValue VData, SelectionDAG &DAG,
return VData;
}
-static bool isAsyncLDSDMA(Intrinsic::ID Intr) {
- switch (Intr) {
- case Intrinsic::amdgcn_raw_buffer_load_async_lds:
- case Intrinsic::amdgcn_raw_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_load_async_to_lds:
- case Intrinsic::amdgcn_global_load_async_lds:
- return true;
- }
- return false;
-}
-
SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
SelectionDAG &DAG) const {
SDLoc DL(Op);
@@ -12982,8 +12969,8 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
? 1
: 0,
DL, MVT::i8)); // swz
- Ops.push_back(
- DAG.getTargetConstant(isAsyncLDSDMA(IntrinsicID), DL, MVT::i8));
+ Ops.push_back(DAG.getTargetConstant(
+ AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID), DL, MVT::i8));
Ops.push_back(M0Val.getValue(0)); // Chain
Ops.push_back(M0Val.getValue(1)); // Glue
@@ -13068,8 +13055,8 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op,
unsigned Aux = Op.getConstantOperandVal(6);
Ops.push_back(DAG.getTargetConstant(Aux & ~AMDGPU::CPol::VIRTUAL_BITS, DL,
MVT::i32)); // CPol
- Ops.push_back(
- DAG.getTargetConstant(isAsyncLDSDMA(IntrinsicID), DL, MVT::i8));
+ Ops.push_back(DAG.getTargetConstant(
+ AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID), DL, MVT::i8));
Ops.push_back(M0Val.getValue(0)); // Chain
Ops.push_back(M0Val.getValue(1)); // Glue
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index 6e6c4ca8afa5f..a8e6822c21683 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -3386,8 +3386,38 @@ struct AlwaysUniform {
};
const AlwaysUniform *lookupAlwaysUniform(unsigned Intr);
+struct LDSDMAIntrinsic {
+ unsigned Intr;
+};
+const LDSDMAIntrinsic *lookupLDSDMAIntrinsic(unsigned Intr);
+
+struct AsyncLDSDMAIntrinsic {
+ unsigned Intr;
+};
+const AsyncLDSDMAIntrinsic *lookupAsyncLDSDMAIntrinsic(unsigned Intr);
+
+struct AsyncIntrinsic {
+ unsigned Intr;
+};
+const AsyncIntrinsic *lookupAsyncIntrinsic(unsigned Intr);
+
+struct TensorIntrinsic {
+ unsigned Intr;
+};
+const TensorIntrinsic *lookupTensorIntrinsic(unsigned Intr);
+
+struct AsyncMarkIntrinsic {
+ unsigned Intr;
+};
+const AsyncMarkIntrinsic *lookupAsyncMarkIntrinsic(unsigned Intr);
+
#define GET_SourcesOfDivergence_IMPL
#define GET_UniformIntrinsics_IMPL
+#define GET_LDSDMAIntrinsics_IMPL
+#define GET_AsyncLDSDMAIntrinsics_IMPL
+#define GET_AsyncIntrinsics_IMPL
+#define GET_TensorIntrinsics_IMPL
+#define GET_AsyncMarkIntrinsics_IMPL
#define GET_Gfx9BufferFormat_IMPL
#define GET_Gfx10BufferFormat_IMPL
#define GET_Gfx11PlusBufferFormat_IMPL
@@ -3404,43 +3434,14 @@ bool isIntrinsicAlwaysUniform(unsigned IntrID) {
return lookupAlwaysUniform(IntrID);
}
+bool isAsyncLDSDMAIntrinsic(unsigned IntrID) {
+ return lookupAsyncLDSDMAIntrinsic(IntrID);
+}
+
bool isAsyncIntrinsic(unsigned IntrID) {
- switch (IntrID) {
- case Intrinsic::amdgcn_raw_buffer_load_lds:
- case Intrinsic::amdgcn_raw_buffer_load_async_lds:
- case Intrinsic::amdgcn_raw_ptr_buffer_load_lds:
- case Intrinsic::amdgcn_raw_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_buffer_load_lds:
- case Intrinsic::amdgcn_struct_buffer_load_async_lds:
- case Intrinsic::amdgcn_struct_ptr_buffer_load_lds:
- case Intrinsic::amdgcn_struct_ptr_buffer_load_async_lds:
- case Intrinsic::amdgcn_load_to_lds:
- case Intrinsic::amdgcn_load_async_to_lds:
- case Intrinsic::amdgcn_global_load_lds:
- case Intrinsic::amdgcn_global_load_async_lds:
- case Intrinsic::amdgcn_cluster_load_async_to_lds_b8:
- case Intrinsic::amdgcn_cluster_load_async_to_lds_b32:
- case Intrinsic::amdgcn_cluster_load_async_to_lds_b64:
- case Intrinsic::amdgcn_cluster_load_async_to_lds_b128:
- case Intrinsic::amdgcn_global_load_async_to_lds_b8:
- case Intrinsic::amdgcn_global_load_async_to_lds_b32:
- case Intrinsic::amdgcn_global_load_async_to_lds_b64:
- case Intrinsic::amdgcn_global_load_async_to_lds_b128:
- case Intrinsic::amdgcn_global_store_async_from_lds_b8:
- case Intrinsic::amdgcn_global_store_async_from_lds_b32:
- case Intrinsic::amdgcn_global_store_async_from_lds_b64:
- case Intrinsic::amdgcn_global_store_async_from_lds_b128:
- case Intrinsic::amdgcn_tensor_load_to_lds:
- case Intrinsic::amdgcn_tensor_store_from_lds:
- case Intrinsic::amdgcn_asyncmark:
- case Intrinsic::amdgcn_wait_asyncmark:
- case Intrinsic::amdgcn_s_wait_asynccnt:
- case Intrinsic::amdgcn_s_wait_tensorcnt:
- case Intrinsic::amdgcn_ds_atomic_async_barrier_arrive_b64:
- return true;
- default:
- return false;
- }
+ return lookupLDSDMAIntrinsic(IntrID) || lookupAsyncLDSDMAIntrinsic(IntrID) ||
+ lookupAsyncIntrinsic(IntrID) || lookupTensorIntrinsic(IntrID) ||
+ lookupAsyncMarkIntrinsic(IntrID);
}
const GcnBufferFormatInfo *getGcnBufferFormatInfo(uint8_t BitsPerComp,
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
index 269594c9be2e8..405ff19c63da7 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
@@ -1778,6 +1778,9 @@ bool isIntrinsicSourceOfDivergence(unsigned IntrID);
/// \returns true if the intrinsic is uniform
bool isIntrinsicAlwaysUniform(unsigned IntrID);
+/// \returns true if the intrinsic is an LDS DMA
+bool isAsyncLDSDMAIntrinsic(unsigned IntrID);
+
/// \returns true if the intrinsic executes an asynchronous operation
bool isAsyncIntrinsic(unsigned IntrID);
>From 7b733c58a25623e1a6fb227b7d1b08f8281dd336 Mon Sep 17 00:00:00 2001
From: Zach Goldthorpe <Zach.Goldthorpe at amd.com>
Date: Wed, 23 Sep 2026 10:33:29 -0500
Subject: [PATCH 2/2] Identify async intrinsics with marker class
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 77 +++++++++++----
.../Target/AMDGPU/AMDGPUSearchableTables.td | 99 ++-----------------
.../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 30 +-----
3 files changed, 68 insertions(+), 138 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index e7829eb29ba34..4b7a06de8bcd5 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -40,6 +40,25 @@ class AMDGPURsrcIntrinsic<int rsrcarg, bit isimage = false> {
bit IsImage = isimage;
}
+// Marker classes for asynchronous operations.
+class AMDGPUAsyncIntrinsic {
+ Intrinsic Intr = !cast<Intrinsic>(NAME);
+}
+
+// DMA transfers between global memory and LDS.
+class AMDGPULDSDMAIntrinsic : AMDGPUAsyncIntrinsic;
+
+// DMA transfers between global memory and LDS tracked with asyncmarks.
+class AMDGPUAsyncLDSDMAIntrinsic : AMDGPUAsyncIntrinsic;
+
+// Operation tracked by ASYNCcnt.
+class AMDGPUAsyncCntIntrinsic : AMDGPUAsyncIntrinsic;
+
+// Operation tracked by TENSORcnt.
+class AMDGPUTensorCntIntrinsic : AMDGPUAsyncIntrinsic;
+
+class AMDGPUAsyncMarkIntrinsic : AMDGPUAsyncIntrinsic;
+
let TargetPrefix = "r600" in {
multiclass AMDGPUReadPreloadRegisterIntrinsic_xyz<
@@ -2086,8 +2105,10 @@ class AMDGPURawBufferLoadLDS : Intrinsic <
ImmArg<ArgIndex<6>>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
let TargetFeatures = "vmem-to-lds-load-insts" in {
- def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS;
- def int_amdgcn_raw_buffer_load_async_lds : AMDGPURawBufferLoadLDS;
+ def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS,
+ AMDGPULDSDMAIntrinsic;
+ def int_amdgcn_raw_buffer_load_async_lds : AMDGPURawBufferLoadLDS,
+ AMDGPUAsyncLDSDMAIntrinsic;
} // End TargetFeatures = "vmem-to-lds-load-insts"
class AMDGPURawPtrBufferLoadLDS :
@@ -2114,9 +2135,11 @@ class AMDGPURawPtrBufferLoadLDS :
let TargetFeatures = "vmem-to-lds-load-insts" in {
def int_amdgcn_raw_ptr_buffer_load_lds : AMDGPURawPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_lds">;
+ ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_lds">,
+ AMDGPULDSDMAIntrinsic;
def int_amdgcn_raw_ptr_buffer_load_async_lds : AMDGPURawPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_async_lds">;
+ ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_async_lds">,
+ AMDGPUAsyncLDSDMAIntrinsic;
} // End TargetFeatures = "vmem-to-lds-load-insts"
@@ -2140,8 +2163,10 @@ class AMDGPUStructBufferLoadLDS : Intrinsic <
ImmArg<ArgIndex<7>>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
let TargetFeatures = "vmem-to-lds-load-insts" in {
- def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS;
- def int_amdgcn_struct_buffer_load_async_lds : AMDGPUStructBufferLoadLDS;
+ def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS,
+ AMDGPULDSDMAIntrinsic;
+ def int_amdgcn_struct_buffer_load_async_lds : AMDGPUStructBufferLoadLDS,
+ AMDGPUAsyncLDSDMAIntrinsic;
} // End TargetFeatures = "vmem-to-lds-load-insts"
class AMDGPUStructPtrBufferLoadLDS :
@@ -2169,9 +2194,11 @@ class AMDGPUStructPtrBufferLoadLDS :
let TargetFeatures = "vmem-to-lds-load-insts" in {
def int_amdgcn_struct_ptr_buffer_load_lds : AMDGPUStructPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_lds">;
+ ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_lds">,
+ AMDGPULDSDMAIntrinsic;
def int_amdgcn_struct_ptr_buffer_load_async_lds : AMDGPUStructPtrBufferLoadLDS,
- ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_async_lds">;
+ ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_async_lds">,
+ AMDGPUAsyncLDSDMAIntrinsic;
} // End TargetFeatures = "vmem-to-lds-load-insts"
def int_amdgcn_s_buffer_prefetch_data : DefaultAttrsIntrinsic <
@@ -2897,8 +2924,8 @@ class AMDGPULoadToLDS :
NoCapture<ArgIndex<1>>, WriteOnly<ArgIndex<1>>,
ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>, IntrNoCallback, IntrNoFree],
"", [SDNPMemOperand]>;
-def int_amdgcn_load_to_lds : AMDGPULoadToLDS;
-def int_amdgcn_load_async_to_lds : AMDGPULoadToLDS;
+def int_amdgcn_load_to_lds : AMDGPULoadToLDS, AMDGPULDSDMAIntrinsic;
+def int_amdgcn_load_async_to_lds : AMDGPULoadToLDS, AMDGPUAsyncLDSDMAIntrinsic;
class AMDGPUGlobalLoadLDS : Intrinsic<
[],
@@ -2916,8 +2943,11 @@ class AMDGPUGlobalLoadLDS : Intrinsic<
ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>,
IntrNoCallback, IntrNoFree],
"", [SDNPMemOperand]>;
-def int_amdgcn_global_load_lds : AMDGPUGlobalLoadLDS, ClangBuiltin<"__builtin_amdgcn_global_load_lds">;
-def int_amdgcn_global_load_async_lds : AMDGPUGlobalLoadLDS, ClangBuiltin<"__builtin_amdgcn_global_load_async_lds">;
+def int_amdgcn_global_load_lds : AMDGPUGlobalLoadLDS,
+ ClangBuiltin<"__builtin_amdgcn_global_load_lds">, AMDGPULDSDMAIntrinsic;
+def int_amdgcn_global_load_async_lds : AMDGPUGlobalLoadLDS,
+ ClangBuiltin<"__builtin_amdgcn_global_load_async_lds">,
+ AMDGPUAsyncLDSDMAIntrinsic;
// This is IntrHasSideEffects because it reads from a volatile hardware register.
let TargetFeatures = "pops-exiting-wave-id" in
@@ -2927,14 +2957,16 @@ def int_amdgcn_pops_exiting_wave_id :
// Sets a marker in the stream of async requests. Modelled as InaccessibleMem.
def int_amdgcn_asyncmark : ClangBuiltin<"__builtin_amdgcn_asyncmark">,
Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrWillReturn,
- IntrNoCallback, IntrNoFree]>;
+ IntrNoCallback, IntrNoFree]>,
+ AMDGPUAsyncMarkIntrinsic;
// Waits until the Nth previous marker is completed, if it exists.
def int_amdgcn_wait_asyncmark :
ClangBuiltin<"__builtin_amdgcn_wait_asyncmark">,
Intrinsic<[], [llvm_i16_ty],
[ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects,
- IntrWillReturn, IntrNoCallback, IntrNoFree]>;
+ IntrWillReturn, IntrNoCallback, IntrNoFree]>,
+ AMDGPUAsyncMarkIntrinsic;
//===----------------------------------------------------------------------===//
// GFX10 Intrinsics
@@ -3846,16 +3878,19 @@ class AMDGPUWaitAsyncIntrinsic :
IntrNoFree]>;
def int_amdgcn_s_wait_asynccnt :
- ClangBuiltin<"__builtin_amdgcn_s_wait_asynccnt">, AMDGPUWaitAsyncIntrinsic;
+ ClangBuiltin<"__builtin_amdgcn_s_wait_asynccnt">, AMDGPUWaitAsyncIntrinsic,
+ AMDGPUAsyncCntIntrinsic;
def int_amdgcn_s_wait_tensorcnt :
- ClangBuiltin<"__builtin_amdgcn_s_wait_tensorcnt">, AMDGPUWaitAsyncIntrinsic;
+ ClangBuiltin<"__builtin_amdgcn_s_wait_tensorcnt">, AMDGPUWaitAsyncIntrinsic,
+ AMDGPUTensorCntIntrinsic;
def int_amdgcn_ds_atomic_async_barrier_arrive_b64 :
ClangBuiltin<"__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64">,
Intrinsic<[], [local_ptr_ty],
// Atomically updates LDS and also ASYNC_CNT which is modeled as InaccessibleMem.
[IntrConvergent, IntrWillReturn, IntrInaccessibleMemOrArgMemOnly, IntrNoCallback, IntrNoFree],
- "", [SDNPMemOperand]>;
+ "", [SDNPMemOperand]>,
+ AMDGPUAsyncCntIntrinsic;
def int_amdgcn_ds_atomic_barrier_arrive_rtn_b64 :
ClangBuiltin<"__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64">,
@@ -4016,7 +4051,7 @@ def int_amdgcn_fdiv_fast : PureIntrinsic<
[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;
// Async instructions increment ASYNCcnt which is modeled as InaccessibleMem.
-class AMDGPUAsyncClusterLoadLDS : Intrinsic <
+class AMDGPUAsyncClusterLoadLDS : AMDGPUAsyncCntIntrinsic, Intrinsic <
[],
[global_ptr_ty, // Base global pointer to load from
local_ptr_ty, // LDS base pointer to store to
@@ -4030,7 +4065,7 @@ class AMDGPUAsyncClusterLoadLDS : Intrinsic <
"", [SDNPMemOperand]
>;
-class AMDGPUAsyncGlobalLoadToLDS : Intrinsic <
+class AMDGPUAsyncGlobalLoadToLDS : AMDGPUAsyncCntIntrinsic, Intrinsic <
[],
[global_ptr_ty, // Base global pointer to load from
local_ptr_ty, // LDS base pointer to store to.
@@ -4043,7 +4078,7 @@ class AMDGPUAsyncGlobalLoadToLDS : Intrinsic <
"", [SDNPMemOperand]
>;
-class AMDGPUAsyncGlobalStoreFromLDS : Intrinsic <
+class AMDGPUAsyncGlobalStoreFromLDS : AMDGPUAsyncCntIntrinsic, Intrinsic <
[],
[global_ptr_ty, // Base global pointer to store to
local_ptr_ty, // LDS base pointer to load from
@@ -4320,7 +4355,7 @@ def int_amdgcn_swmmac_i32_16x16x128_iu8 : AMDGPUSWmmacIntrinsicABIdxClamp<ll
}
class AMDGPUTensorLoadStore:
- Intrinsic<
+ AMDGPUTensorCntIntrinsic, Intrinsic<
[],
[llvm_v4i32_ty, // D# group 0
llvm_v8i32_ty, // D# group 1
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
index c69b0b2871f35..c471f8c3a76fe 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td
@@ -430,103 +430,20 @@ def AMDGPUImageDMaskIntrinsicTable : GenericTable {
let PrimaryKeyEarlyOut = 1;
}
-// DMA transfers between global memory and LDS.
-class LDSDMAIntrinsic<Intrinsic intr> {
- Intrinsic Intr = intr;
-}
-
-def LDSDMAIntrinsics : GenericTable {
- let FilterClass = "LDSDMAIntrinsic";
+class AsyncIntrinsicTable : GenericTable {
let Fields = ["Intr"];
let PrimaryKey = ["Intr"];
- let PrimaryKeyName = "lookupLDSDMAIntrinsic";
-}
-
-def : LDSDMAIntrinsic<int_amdgcn_raw_buffer_load_lds>;
-def : LDSDMAIntrinsic<int_amdgcn_raw_ptr_buffer_load_lds>;
-def : LDSDMAIntrinsic<int_amdgcn_struct_buffer_load_lds>;
-def : LDSDMAIntrinsic<int_amdgcn_struct_ptr_buffer_load_lds>;
-def : LDSDMAIntrinsic<int_amdgcn_load_to_lds>;
-def : LDSDMAIntrinsic<int_amdgcn_global_load_lds>;
-
-// DMA transfers between global memory and LDS tracked with asyncmarks.
-class AsyncLDSDMAIntrinsic<Intrinsic intr> {
- Intrinsic Intr = intr;
}
-def AsyncLDSDMAIntrinsics : GenericTable {
- let FilterClass = "AsyncLDSDMAIntrinsic";
- let Fields = ["Intr"];
-
- let PrimaryKey = ["Intr"];
- let PrimaryKeyName = "lookupAsyncLDSDMAIntrinsic";
-}
-
-def : AsyncLDSDMAIntrinsic<int_amdgcn_raw_buffer_load_async_lds>;
-def : AsyncLDSDMAIntrinsic<int_amdgcn_raw_ptr_buffer_load_async_lds>;
-def : AsyncLDSDMAIntrinsic<int_amdgcn_struct_buffer_load_async_lds>;
-def : AsyncLDSDMAIntrinsic<int_amdgcn_struct_ptr_buffer_load_async_lds>;
-def : AsyncLDSDMAIntrinsic<int_amdgcn_load_async_to_lds>;
-def : AsyncLDSDMAIntrinsic<int_amdgcn_global_load_async_lds>;
-
-// Operations tracked by ASYNCcnt.
-class AsyncIntrinsic<Intrinsic intr> {
- Intrinsic Intr = intr;
-}
-
-def AsyncIntrinsics : GenericTable {
- let FilterClass = "AsyncIntrinsic";
- let Fields = ["Intr"];
-
- let PrimaryKey = ["Intr"];
+def AsyncIntrinsics : AsyncIntrinsicTable {
+ let FilterClass = "AMDGPUAsyncIntrinsic";
+ let CppTypeName = "AsyncIntrinsic";
let PrimaryKeyName = "lookupAsyncIntrinsic";
}
-def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b8>;
-def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b32>;
-def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b64>;
-def : AsyncIntrinsic<int_amdgcn_cluster_load_async_to_lds_b128>;
-def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b8>;
-def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b32>;
-def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b64>;
-def : AsyncIntrinsic<int_amdgcn_global_load_async_to_lds_b128>;
-def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b8>;
-def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b32>;
-def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b64>;
-def : AsyncIntrinsic<int_amdgcn_global_store_async_from_lds_b128>;
-def : AsyncIntrinsic<int_amdgcn_ds_atomic_async_barrier_arrive_b64>;
-def : AsyncIntrinsic<int_amdgcn_s_wait_asynccnt>;
-
-// Operations tracked by TENSORcnt.
-class TensorIntrinsic<Intrinsic intr> {
- Intrinsic Intr = intr;
-}
-
-def TensorIntrinsics : GenericTable {
- let FilterClass = "TensorIntrinsic";
- let Fields = ["Intr"];
-
- let PrimaryKey = ["Intr"];
- let PrimaryKeyName = "lookupTensorIntrinsic";
-}
-
-def : TensorIntrinsic<int_amdgcn_tensor_load_to_lds>;
-def : TensorIntrinsic<int_amdgcn_tensor_store_from_lds>;
-def : TensorIntrinsic<int_amdgcn_s_wait_tensorcnt>;
-
-// Asyncmark production and consumption.
-class AsyncMarkIntrinsic<Intrinsic intr> {
- Intrinsic Intr = intr;
-}
-
-def AsyncMarkIntrinsics : GenericTable {
- let FilterClass = "AsyncMarkIntrinsic";
- let Fields = ["Intr"];
-
- let PrimaryKey = ["Intr"];
- let PrimaryKeyName = "lookupAsyncMarkIntrinsic";
+def AsyncLDSDMAIntrinsics : AsyncIntrinsicTable {
+ let FilterClass = "AMDGPUAsyncLDSDMAIntrinsic";
+ let CppTypeName = "AsyncLDSDMAIntrinsic";
+ let PrimaryKeyName = "lookupAsyncLDSDMAIntrinsic";
}
-
-def : AsyncMarkIntrinsic<int_amdgcn_asyncmark>;
-def : AsyncMarkIntrinsic<int_amdgcn_wait_asyncmark>;
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
index a8e6822c21683..5df232c29eb40 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
@@ -3386,38 +3386,20 @@ struct AlwaysUniform {
};
const AlwaysUniform *lookupAlwaysUniform(unsigned Intr);
-struct LDSDMAIntrinsic {
- unsigned Intr;
-};
-const LDSDMAIntrinsic *lookupLDSDMAIntrinsic(unsigned Intr);
-
-struct AsyncLDSDMAIntrinsic {
- unsigned Intr;
-};
-const AsyncLDSDMAIntrinsic *lookupAsyncLDSDMAIntrinsic(unsigned Intr);
-
struct AsyncIntrinsic {
unsigned Intr;
};
const AsyncIntrinsic *lookupAsyncIntrinsic(unsigned Intr);
-struct TensorIntrinsic {
- unsigned Intr;
-};
-const TensorIntrinsic *lookupTensorIntrinsic(unsigned Intr);
-
-struct AsyncMarkIntrinsic {
+struct AsyncLDSDMAIntrinsic {
unsigned Intr;
};
-const AsyncMarkIntrinsic *lookupAsyncMarkIntrinsic(unsigned Intr);
+const AsyncLDSDMAIntrinsic *lookupAsyncLDSDMAIntrinsic(unsigned Intr);
#define GET_SourcesOfDivergence_IMPL
#define GET_UniformIntrinsics_IMPL
-#define GET_LDSDMAIntrinsics_IMPL
-#define GET_AsyncLDSDMAIntrinsics_IMPL
#define GET_AsyncIntrinsics_IMPL
-#define GET_TensorIntrinsics_IMPL
-#define GET_AsyncMarkIntrinsics_IMPL
+#define GET_AsyncLDSDMAIntrinsics_IMPL
#define GET_Gfx9BufferFormat_IMPL
#define GET_Gfx10BufferFormat_IMPL
#define GET_Gfx11PlusBufferFormat_IMPL
@@ -3438,11 +3420,7 @@ bool isAsyncLDSDMAIntrinsic(unsigned IntrID) {
return lookupAsyncLDSDMAIntrinsic(IntrID);
}
-bool isAsyncIntrinsic(unsigned IntrID) {
- return lookupLDSDMAIntrinsic(IntrID) || lookupAsyncLDSDMAIntrinsic(IntrID) ||
- lookupAsyncIntrinsic(IntrID) || lookupTensorIntrinsic(IntrID) ||
- lookupAsyncMarkIntrinsic(IntrID);
-}
+bool isAsyncIntrinsic(unsigned IntrID) { return lookupAsyncIntrinsic(IntrID); }
const GcnBufferFormatInfo *getGcnBufferFormatInfo(uint8_t BitsPerComp,
uint8_t NumComponents,
More information about the llvm-branch-commits
mailing list