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 <[email protected]> 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 <[email protected]> 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, _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
