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/3] [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/3] 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, >From a1db281a9b8e643b6814d6c121d70525d6ff3db5 Mon Sep 17 00:00:00 2001 From: Zach Goldthorpe <[email protected]> Date: Wed, 7 Oct 2026 12:49:06 -0500 Subject: [PATCH 3/3] Narrow marker class to LDSDMA intrinsics --- llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 88 ++++++------------- .../AMDGPU/AMDGPUInstructionSelector.cpp | 17 +++- .../Target/AMDGPU/AMDGPUSearchableTables.td | 17 +--- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 21 ++++- .../Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp | 18 ++-- llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h | 3 - 6 files changed, 70 insertions(+), 94 deletions(-) diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 4b7a06de8bcd5..ac766ab192c4e 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -40,25 +40,11 @@ class AMDGPURsrcIntrinsic<int rsrcarg, bit isimage = false> { bit IsImage = isimage; } -// Marker classes for asynchronous operations. -class AMDGPUAsyncIntrinsic { +// Marker class for DMA transfers between global memory and LDS. +class AMDGPULDSDMAIntrinsic { 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,7 +2072,7 @@ def int_amdgcn_struct_tbuffer_store : DefaultAttrsIntrinsic < ImmArg<ArgIndex<5>>, ImmArg<ArgIndex<6>>], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<1>; -class AMDGPURawBufferLoadLDS : Intrinsic < +class AMDGPURawBufferLoadLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [llvm_v4i32_ty, // rsrc(SGPR) LLVMQualPointerType<3>, // LDS base offset @@ -2105,13 +2091,11 @@ 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, - AMDGPULDSDMAIntrinsic; - def int_amdgcn_raw_buffer_load_async_lds : AMDGPURawBufferLoadLDS, - AMDGPUAsyncLDSDMAIntrinsic; + def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS; + def int_amdgcn_raw_buffer_load_async_lds : AMDGPURawBufferLoadLDS; } // End TargetFeatures = "vmem-to-lds-load-insts" -class AMDGPURawPtrBufferLoadLDS : +class AMDGPURawPtrBufferLoadLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [AMDGPUBufferRsrcTy, // rsrc(SGPR) @@ -2135,15 +2119,13 @@ 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">, - AMDGPULDSDMAIntrinsic; + ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_lds">; def int_amdgcn_raw_ptr_buffer_load_async_lds : AMDGPURawPtrBufferLoadLDS, - ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_async_lds">, - AMDGPUAsyncLDSDMAIntrinsic; + ClangBuiltin<"__builtin_amdgcn_raw_ptr_buffer_load_async_lds">; } // End TargetFeatures = "vmem-to-lds-load-insts" -class AMDGPUStructBufferLoadLDS : Intrinsic < +class AMDGPUStructBufferLoadLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [llvm_v4i32_ty, // rsrc(SGPR) LLVMQualPointerType<3>, // LDS base offset @@ -2163,13 +2145,11 @@ 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, - AMDGPULDSDMAIntrinsic; - def int_amdgcn_struct_buffer_load_async_lds : AMDGPUStructBufferLoadLDS, - AMDGPUAsyncLDSDMAIntrinsic; + def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS; + def int_amdgcn_struct_buffer_load_async_lds : AMDGPUStructBufferLoadLDS; } // End TargetFeatures = "vmem-to-lds-load-insts" -class AMDGPUStructPtrBufferLoadLDS : +class AMDGPUStructPtrBufferLoadLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [AMDGPUBufferRsrcTy, // rsrc(SGPR) @@ -2194,11 +2174,9 @@ 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">, - AMDGPULDSDMAIntrinsic; + ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_lds">; def int_amdgcn_struct_ptr_buffer_load_async_lds : AMDGPUStructPtrBufferLoadLDS, - ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_async_lds">, - AMDGPUAsyncLDSDMAIntrinsic; + ClangBuiltin<"__builtin_amdgcn_struct_ptr_buffer_load_async_lds">; } // End TargetFeatures = "vmem-to-lds-load-insts" def int_amdgcn_s_buffer_prefetch_data : DefaultAttrsIntrinsic < @@ -2909,7 +2887,7 @@ def int_amdgcn_perm : /// This allows abstracting over both global pointers (address space 1) and /// the buffer-resource-wrapper pointers (address space 7 and 9). /// TODO: add support for address space 5 and scratch_load_lds. -class AMDGPULoadToLDS : +class AMDGPULoadToLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [llvm_anyptr_ty, // Base pointer to load from. Varies per lane. @@ -2924,10 +2902,10 @@ 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, AMDGPULDSDMAIntrinsic; -def int_amdgcn_load_async_to_lds : AMDGPULoadToLDS, AMDGPUAsyncLDSDMAIntrinsic; +def int_amdgcn_load_to_lds : AMDGPULoadToLDS; +def int_amdgcn_load_async_to_lds : AMDGPULoadToLDS; -class AMDGPUGlobalLoadLDS : Intrinsic< +class AMDGPUGlobalLoadLDS : AMDGPULDSDMAIntrinsic, Intrinsic< [], [LLVMQualPointerType<1>, // Base global pointer to load from LLVMQualPointerType<3>, // LDS base pointer to store to @@ -2943,11 +2921,8 @@ 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">, AMDGPULDSDMAIntrinsic; -def int_amdgcn_global_load_async_lds : AMDGPUGlobalLoadLDS, - ClangBuiltin<"__builtin_amdgcn_global_load_async_lds">, - AMDGPUAsyncLDSDMAIntrinsic; +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">; // This is IntrHasSideEffects because it reads from a volatile hardware register. let TargetFeatures = "pops-exiting-wave-id" in @@ -2957,16 +2932,14 @@ 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]>, - AMDGPUAsyncMarkIntrinsic; + IntrNoCallback, IntrNoFree]>; // 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]>, - AMDGPUAsyncMarkIntrinsic; + IntrWillReturn, IntrNoCallback, IntrNoFree]>; //===----------------------------------------------------------------------===// // GFX10 Intrinsics @@ -3878,19 +3851,16 @@ class AMDGPUWaitAsyncIntrinsic : IntrNoFree]>; def int_amdgcn_s_wait_asynccnt : - ClangBuiltin<"__builtin_amdgcn_s_wait_asynccnt">, AMDGPUWaitAsyncIntrinsic, - AMDGPUAsyncCntIntrinsic; + ClangBuiltin<"__builtin_amdgcn_s_wait_asynccnt">, AMDGPUWaitAsyncIntrinsic; def int_amdgcn_s_wait_tensorcnt : - ClangBuiltin<"__builtin_amdgcn_s_wait_tensorcnt">, AMDGPUWaitAsyncIntrinsic, - AMDGPUTensorCntIntrinsic; + ClangBuiltin<"__builtin_amdgcn_s_wait_tensorcnt">, AMDGPUWaitAsyncIntrinsic; 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]>, - AMDGPUAsyncCntIntrinsic; + "", [SDNPMemOperand]>; def int_amdgcn_ds_atomic_barrier_arrive_rtn_b64 : ClangBuiltin<"__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64">, @@ -4051,7 +4021,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 : AMDGPUAsyncCntIntrinsic, Intrinsic < +class AMDGPUAsyncClusterLoadLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [global_ptr_ty, // Base global pointer to load from local_ptr_ty, // LDS base pointer to store to @@ -4065,7 +4035,7 @@ class AMDGPUAsyncClusterLoadLDS : AMDGPUAsyncCntIntrinsic, Intrinsic < "", [SDNPMemOperand] >; -class AMDGPUAsyncGlobalLoadToLDS : AMDGPUAsyncCntIntrinsic, Intrinsic < +class AMDGPUAsyncGlobalLoadToLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [global_ptr_ty, // Base global pointer to load from local_ptr_ty, // LDS base pointer to store to. @@ -4078,7 +4048,7 @@ class AMDGPUAsyncGlobalLoadToLDS : AMDGPUAsyncCntIntrinsic, Intrinsic < "", [SDNPMemOperand] >; -class AMDGPUAsyncGlobalStoreFromLDS : AMDGPUAsyncCntIntrinsic, Intrinsic < +class AMDGPUAsyncGlobalStoreFromLDS : AMDGPULDSDMAIntrinsic, Intrinsic < [], [global_ptr_ty, // Base global pointer to store to local_ptr_ty, // LDS base pointer to load from @@ -4355,7 +4325,7 @@ def int_amdgcn_swmmac_i32_16x16x128_iu8 : AMDGPUSWmmacIntrinsicABIdxClamp<ll } class AMDGPUTensorLoadStore: - AMDGPUTensorCntIntrinsic, Intrinsic< + AMDGPULDSDMAIntrinsic, Intrinsic< [], [llvm_v4i32_ty, // D# group 0 llvm_v8i32_ty, // D# group 1 diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp index 09544e4b12263..140cf58e8fdd3 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp @@ -3521,6 +3521,19 @@ 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; @@ -3616,7 +3629,7 @@ bool AMDGPUInstructionSelector::selectBufferLoadLds(MachineInstr &MI) const { Aux & (IsGFX12Plus ? AMDGPU::CPol::SWZ : AMDGPU::CPol::SWZ_pregfx12) ? 1 : 0); // swz - MIB.addImm(AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID)); + MIB.addImm(isAsyncLDSDMA(IntrinsicID)); MachineMemOperand *LoadMMO = *MI.memoperands_begin(); // Don't set the offset value here because the pointer points to the base of @@ -3811,7 +3824,7 @@ bool AMDGPUInstructionSelector::selectGlobalLoadLds(MachineInstr &MI) const{ unsigned Aux = MI.getOperand(5).getImm(); MIB.addImm(Aux & ~AMDGPU::CPol::VIRTUAL_BITS); // cpol - MIB.addImm(AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID)); + MIB.addImm(isAsyncLDSDMA(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 c471f8c3a76fe..ffecd6b5c54b1 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td +++ b/llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td @@ -430,20 +430,11 @@ def AMDGPUImageDMaskIntrinsicTable : GenericTable { let PrimaryKeyEarlyOut = 1; } -class AsyncIntrinsicTable : GenericTable { +def LDSDMAIntrinsicTable : GenericTable { + let FilterClass = "AMDGPULDSDMAIntrinsic"; + let CppTypeName = "LDSDMAIntrinsic"; let Fields = ["Intr"]; let PrimaryKey = ["Intr"]; -} - -def AsyncIntrinsics : AsyncIntrinsicTable { - let FilterClass = "AMDGPUAsyncIntrinsic"; - let CppTypeName = "AsyncIntrinsic"; - let PrimaryKeyName = "lookupAsyncIntrinsic"; -} - -def AsyncLDSDMAIntrinsics : AsyncIntrinsicTable { - let FilterClass = "AMDGPUAsyncLDSDMAIntrinsic"; - let CppTypeName = "AsyncLDSDMAIntrinsic"; - let PrimaryKeyName = "lookupAsyncLDSDMAIntrinsic"; + let PrimaryKeyName = "lookupLDSDMAIntrinsic"; } diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index 4d04e98bf408a..f9a4ba4c1bd1d 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -12672,6 +12672,19 @@ 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); @@ -12969,8 +12982,8 @@ SDValue SITargetLowering::LowerINTRINSIC_VOID(SDValue Op, ? 1 : 0, DL, MVT::i8)); // swz - Ops.push_back(DAG.getTargetConstant( - AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID), DL, MVT::i8)); + Ops.push_back( + DAG.getTargetConstant(isAsyncLDSDMA(IntrinsicID), DL, MVT::i8)); Ops.push_back(M0Val.getValue(0)); // Chain Ops.push_back(M0Val.getValue(1)); // Glue @@ -13055,8 +13068,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( - AMDGPU::isAsyncLDSDMAIntrinsic(IntrinsicID), DL, MVT::i8)); + Ops.push_back( + DAG.getTargetConstant(isAsyncLDSDMA(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 5df232c29eb40..4a2c0cdee5fb0 100644 --- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp +++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp @@ -3386,20 +3386,14 @@ struct AlwaysUniform { }; const AlwaysUniform *lookupAlwaysUniform(unsigned Intr); -struct AsyncIntrinsic { +struct LDSDMAIntrinsic { unsigned Intr; }; -const AsyncIntrinsic *lookupAsyncIntrinsic(unsigned Intr); - -struct AsyncLDSDMAIntrinsic { - unsigned Intr; -}; -const AsyncLDSDMAIntrinsic *lookupAsyncLDSDMAIntrinsic(unsigned Intr); +const LDSDMAIntrinsic *lookupLDSDMAIntrinsic(unsigned Intr); #define GET_SourcesOfDivergence_IMPL #define GET_UniformIntrinsics_IMPL -#define GET_AsyncIntrinsics_IMPL -#define GET_AsyncLDSDMAIntrinsics_IMPL +#define GET_LDSDMAIntrinsicTable_IMPL #define GET_Gfx9BufferFormat_IMPL #define GET_Gfx10BufferFormat_IMPL #define GET_Gfx11PlusBufferFormat_IMPL @@ -3416,12 +3410,10 @@ bool isIntrinsicAlwaysUniform(unsigned IntrID) { return lookupAlwaysUniform(IntrID); } -bool isAsyncLDSDMAIntrinsic(unsigned IntrID) { - return lookupAsyncLDSDMAIntrinsic(IntrID); +bool isLDSDMAIntrinsic(unsigned IntrID) { + return lookupLDSDMAIntrinsic(IntrID); } -bool isAsyncIntrinsic(unsigned IntrID) { return lookupAsyncIntrinsic(IntrID); } - const GcnBufferFormatInfo *getGcnBufferFormatInfo(uint8_t BitsPerComp, uint8_t NumComponents, uint8_t NumFormat, diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h index e9ef5542dfd1d..06054954665d8 100644 --- a/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h +++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h @@ -1778,9 +1778,6 @@ 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 LDSDMA operation bool isLDSDMAIntrinsic(unsigned IntrID); _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
