[llvm] AMDGPU: share LDS budget logic and add experimental LDS buffering pass (PR #166388)
Yaxun Liu via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 1 14:46:56 PDT 2026
https://github.com/yxsamliu updated https://github.com/llvm/llvm-project/pull/166388
>From 6cc42f6176d3dba9dcf804dc1a4705e96d653ba9 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 6 Jan 2026 16:41:48 -0500
Subject: [PATCH 1/3] AMDGPU: add LDS buffering pass and shared LDS budget
helper
Partially updating an aggregate can leave an unchanged fragment loaded
from and later stored to the same global address. Potentially aliasing
accesses can keep this pair from being removed, leaving the value live
in VGPRs.
Add an experimental pass that buffers eligible values through LDS to
shorten their VGPR live ranges. Share the LDS budget computation with
alloca promotion so both transforms respect the same occupancy limits.
---
llvm/lib/Target/AMDGPU/AMDGPU.h | 17 ++
llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp | 271 ++++++++++++++++++
llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 11 +-
.../lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp | 121 +-------
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 37 +++
llvm/lib/Target/AMDGPU/CMakeLists.txt | 1 +
.../Target/AMDGPU/Utils/AMDGPULDSUtils.cpp | 235 +++++++++++++++
llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h | 53 ++++
llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt | 1 +
.../test/CodeGen/AMDGPU/lds-buffering-sroa.ll | 83 ++++++
llvm/test/CodeGen/AMDGPU/lds-buffering.ll | 183 ++++++++++++
11 files changed, 899 insertions(+), 114 deletions(-)
create mode 100644 llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
create mode 100644 llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp
create mode 100644 llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering-sroa.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index 9fe4123a27bbd..4fa35f4a92d46 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -292,6 +292,23 @@ struct AMDGPUPromoteAllocaToVectorPass
TargetMachine &TM;
};
+// Buffer selected per-thread global memory through LDS to improve
+// performance in memory-bound kernels. This runs late and is separate
+// from alloca promotion.
+struct AMDGPULDSBufferingPass : OptionalPassInfoMixin<AMDGPULDSBufferingPass> {
+ AMDGPULDSBufferingPass(const AMDGPUTargetMachine &TM, unsigned MaxBytes = 64)
+ : TM(TM), MaxBytes(MaxBytes) {}
+ PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM);
+
+private:
+ const AMDGPUTargetMachine &TM;
+ unsigned MaxBytes;
+};
+
+// Legacy PM wrapper for LDS buffering
+FunctionPass *createAMDGPULDSBufferingLegacyPass();
+void initializeAMDGPULDSBufferingLegacyPass(PassRegistry &);
+
struct AMDGPUAtomicOptimizerPass
: OptionalPassInfoMixin<AMDGPUAtomicOptimizerPass> {
AMDGPUAtomicOptimizerPass(TargetMachine &TM, ScanOptions ScanImpl)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp b/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
new file mode 100644
index 0000000000000..03d3482c247fe
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
@@ -0,0 +1,271 @@
+//===-- AMDGPULDSBuffering.cpp - Per-thread LDS buffering -----------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This pass buffers selected per-thread global memory accesses through LDS
+// (addrspace(3)) to improve performance in memory-bound kernels.
+//
+// SROA can split a partially updated aggregate into independent fragments. An
+// unchanged fragment can become a load whose only use is a store back to the
+// same global location. Intervening potentially-aliasing memory operations can
+// prevent generic load/store forwarding. Buffering the long-lived value in LDS
+// shortens its VGPR live range.
+//
+// The pass runs late in the pipeline, after SROA and AMDGPUPromoteAlloca,
+// using only leftover LDS budget to avoid interfering with other LDS
+// optimizations. It respects the same LDS budget constraints as
+// AMDGPUPromoteAlloca, ensuring that LDS usage remains within occupancy
+// tier limits.
+//
+// Current implementation handles the simplest pattern: a load from global
+// memory whose only use is a store back to the same pointer. This pattern is
+// transformed into a pair of memcpy operations (global->LDS and LDS->global),
+// effectively moving the value through LDS instead of accessing global memory
+// directly.
+//
+// This pass was inspired by finding that some rocrand performance tests
+// show better performance when global memory is buffered through LDS
+// instead of being loaded/stored to registers directly. This optimization
+// is experimental and must be enabled in the default pipeline via the
+// -amdgpu-enable-lds-buffering flag (or explicitly scheduled via
+// -passes='amdgpu-lds-buffering<...>').
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPU.h"
+#include "AMDGPUTargetMachine.h"
+#include "GCNSubtarget.h"
+#include "Utils/AMDGPUBaseInfo.h"
+#include "Utils/AMDGPULDSUtils.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/CodeGen/TargetPassConfig.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
+#include "llvm/IR/PassManager.h"
+#include "llvm/IR/PatternMatch.h"
+#include "llvm/InitializePasses.h"
+#include "llvm/Pass.h"
+#include "llvm/Support/Alignment.h"
+#include "llvm/Support/Debug.h"
+#include <algorithm>
+
+#define DEBUG_TYPE "amdgpu-lds-buffering"
+
+using namespace llvm;
+
+namespace {
+
+class AMDGPULDSBufferingImpl {
+ const AMDGPUTargetMachine &TM;
+ unsigned MaxBytes;
+ Module *Mod = nullptr;
+ const DataLayout *DL = nullptr;
+
+public:
+ AMDGPULDSBufferingImpl(const AMDGPUTargetMachine &TM, unsigned MaxBytes)
+ : TM(TM), MaxBytes(MaxBytes) {}
+
+ bool run(Function &F) {
+ LLVM_DEBUG(dbgs() << "[LDSBuffer] Visit function: " << F.getName() << '\n');
+ if (!AMDGPU::isEntryFunctionCC(F.getCallingConv()))
+ return false;
+
+ Mod = F.getParent();
+ DL = &Mod->getDataLayout();
+
+ AMDGPU::AMDGPULDSBudget Budget = AMDGPU::computeLDSBudget(F, TM);
+ if (!Budget.promotable)
+ return false;
+ uint32_t localUsage = Budget.currentUsage;
+ uint32_t localLimit = Budget.limit;
+
+ const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
+ unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(F).second;
+
+ bool Changed = false;
+ unsigned NumTransformed = 0;
+
+ // Minimal pattern: a load from AS(1) whose only use is a store back to the
+ // exact same pointer later. Replace with global<->LDS memcpy pair to
+ // shorten the live range and free VGPRs.
+ SmallVector<Instruction *> ToErase;
+ for (BasicBlock &BB : F) {
+ for (Instruction &I : llvm::make_early_inc_range(BB)) {
+ auto *LI = dyn_cast<LoadInst>(&I);
+ if (!LI || LI->isVolatile() || LI->isAtomic())
+ continue;
+
+ Type *ValTy = LI->getType();
+ if (!ValTy->isFirstClassType())
+ continue;
+
+ Value *Ptr = LI->getPointerOperand();
+ auto *PtrTy = cast<PointerType>(Ptr->getType());
+ if (PtrTy->getAddressSpace() != AMDGPUAS::GLOBAL_ADDRESS)
+ continue;
+
+ if (!LI->hasOneUse())
+ continue;
+ auto *SI = dyn_cast<StoreInst>(LI->user_back());
+ if (!SI || SI->isVolatile() || SI->isAtomic())
+ continue;
+ if (SI->getValueOperand() != LI)
+ continue;
+
+ Value *SPtr = SI->getPointerOperand();
+ if (SPtr != Ptr)
+ continue;
+
+ TypeSize TS = DL->getTypeStoreSize(ValTy);
+ if (TS.isScalable())
+ continue;
+ uint64_t Size = TS.getFixedValue();
+ if (Size == 0 || Size > MaxBytes)
+ continue;
+ Align MinAlign = Align(16);
+ Align LoadAlign = LI->getAlign();
+ Align StoreAlign = SI->getAlign();
+ Align Alignment = std::min(LoadAlign, StoreAlign);
+ if (Alignment < MinAlign)
+ continue;
+
+ // Create LDS slot near the load and emit memcpy global->LDS.
+ LLVM_DEBUG({
+ dbgs() << "[LDSBuffer] Candidate found: load->store same ptr in "
+ << F.getName() << '\n';
+ dbgs() << " size=" << Size
+ << "B, loadAlign=" << LoadAlign.value()
+ << ", storeAlign=" << StoreAlign.value()
+ << ", chosenAlign=" << Alignment.value()
+ << ", ptr AS=" << PtrTy->getAddressSpace() << "\n";
+ });
+ IRBuilder<> BLoad(LI);
+
+ // Ensure LDS budget allows allocating a per-thread slot.
+ uint32_t NewSize = alignTo(localUsage, Alignment);
+ NewSize += WorkGroupSize * static_cast<uint32_t>(Size);
+ if (NewSize > localLimit)
+ continue;
+ localUsage = NewSize;
+ auto [GV, SlotPtr] =
+ createLDSGlobalAndThreadSlot(F, ValTy, Alignment, "ldsbuf", BLoad);
+ // memcpy p3 <- p1
+ LLVM_DEBUG(dbgs() << "[LDSBuffer] Insert memcpy global->LDS: "
+ << GV->getName() << ", bytes=" << Size
+ << ", align=" << Alignment.value() << '\n');
+ BLoad.CreateMemCpy(SlotPtr, Alignment, Ptr, LoadAlign, TS);
+
+ // Replace the final store with memcpy LDS->global.
+ IRBuilder<> BStore(SI);
+ LLVM_DEBUG(dbgs() << "[LDSBuffer] Insert memcpy LDS->global: "
+ << GV->getName() << ", bytes=" << Size
+ << ", align=" << Alignment.value() << '\n');
+ BStore.CreateMemCpy(SPtr, StoreAlign, SlotPtr, Alignment, TS);
+
+ ToErase.push_back(SI);
+ ToErase.push_back(LI);
+ LLVM_DEBUG(dbgs() << "[LDSBuffer] Erase original load/store pair\n");
+ Changed = true;
+ ++NumTransformed;
+ }
+ }
+
+ for (Instruction *E : ToErase)
+ E->eraseFromParent();
+
+ LLVM_DEBUG(dbgs() << "[LDSBuffer] Transformations applied: "
+ << NumTransformed << "\n");
+
+ return Changed;
+ }
+
+private:
+ // Create an LDS array [WGSize x ElemTy] and return pointer to per-thread
+ // slot.
+ std::pair<GlobalVariable *, Value *>
+ createLDSGlobalAndThreadSlot(Function &F, Type *ElemTy, Align Alignment,
+ StringRef BaseName, IRBuilder<> &Builder) {
+ const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
+ unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(F).second;
+ Type *ArrTy = ArrayType::get(ElemTy, WorkGroupSize);
+ GlobalVariable *GV = new GlobalVariable(
+ *Mod, ArrTy, /*isConstant=*/false, GlobalValue::InternalLinkage,
+ PoisonValue::get(ArrTy), (F.getName() + "." + BaseName).str(), nullptr,
+ GlobalVariable::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS);
+ GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
+ GV->setAlignment(Alignment);
+
+ LLVM_DEBUG({
+ dbgs() << "[LDSBuffer] Create LDS global: name=" << GV->getName()
+ << ", elemTy=" << *ElemTy << ", WGSize=" << WorkGroupSize
+ << ", align=" << Alignment.value() << '\n';
+ });
+
+ Value *LinearTID = AMDGPU::buildLinearThreadId(Builder, *Mod, ST);
+ LLVMContext &Ctx = Mod->getContext();
+ Value *Indices[] = {Constant::getNullValue(Type::getInt32Ty(Ctx)),
+ LinearTID};
+ Value *SlotPtr = Builder.CreateInBoundsGEP(ArrTy, GV, Indices);
+ return {GV, SlotPtr};
+ }
+};
+
+} // end anonymous namespace
+
+PreservedAnalyses AMDGPULDSBufferingPass::run(Function &F,
+ FunctionAnalysisManager &AM) {
+ bool Changed = AMDGPULDSBufferingImpl(TM, MaxBytes).run(F);
+ if (!Changed)
+ return PreservedAnalyses::all();
+
+ PreservedAnalyses PA;
+ PA.preserveSet<CFGAnalyses>();
+ return PA;
+}
+
+//===----------------------------------------------------------------------===//
+// Legacy PM wrapper
+//===----------------------------------------------------------------------===//
+
+namespace {
+
+class AMDGPULDSBufferingLegacy : public FunctionPass {
+public:
+ static char ID;
+ AMDGPULDSBufferingLegacy() : FunctionPass(ID) {}
+
+ StringRef getPassName() const override { return "AMDGPU LDS Buffering"; }
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.setPreservesCFG();
+ FunctionPass::getAnalysisUsage(AU);
+ }
+
+ bool runOnFunction(Function &F) override {
+ if (skipFunction(F))
+ return false;
+ if (TargetPassConfig *TPC = getAnalysisIfAvailable<TargetPassConfig>())
+ return AMDGPULDSBufferingImpl(TPC->getTM<AMDGPUTargetMachine>(),
+ /*MaxBytes=*/64)
+ .run(F);
+ return false;
+ }
+};
+
+} // end anonymous namespace
+
+char AMDGPULDSBufferingLegacy::ID = 0;
+
+INITIALIZE_PASS_BEGIN(AMDGPULDSBufferingLegacy, DEBUG_TYPE,
+ "AMDGPU per-thread LDS buffering", false, false)
+INITIALIZE_PASS_END(AMDGPULDSBufferingLegacy, DEBUG_TYPE,
+ "AMDGPU per-thread LDS buffering", false, false)
+
+FunctionPass *llvm::createAMDGPULDSBufferingLegacyPass() {
+ return new AMDGPULDSBufferingLegacy();
+}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
index 6da139ea0b59c..e174d6b944c14 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
@@ -73,11 +73,14 @@ FUNCTION_PASS("amdgpu-promote-kernel-arguments",
AMDGPUPromoteKernelArgumentsPass())
FUNCTION_PASS("amdgpu-rewrite-undef-for-phi", AMDGPURewriteUndefForPHIPass())
FUNCTION_PASS("amdgpu-simplifylib", AMDGPUSimplifyLibCallsPass())
+FUNCTION_PASS("amdgpu-uniform-intrinsic-combine",
+ AMDGPUUniformIntrinsicCombinePass())
FUNCTION_PASS("amdgpu-unify-divergent-exit-nodes",
AMDGPUUnifyDivergentExitNodesPass())
FUNCTION_PASS("amdgpu-usenative", AMDGPUUseNativeCallsPass())
-FUNCTION_PASS("si-annotate-control-flow", SIAnnotateControlFlowPass(*static_cast<const GCNTargetMachine *>(this)))
-FUNCTION_PASS("amdgpu-uniform-intrinsic-combine", AMDGPUUniformIntrinsicCombinePass())
+FUNCTION_PASS(
+ "si-annotate-control-flow",
+ SIAnnotateControlFlowPass(*static_cast<const GCNTargetMachine *>(this)))
#undef FUNCTION_PASS
#ifndef FUNCTION_ANALYSIS
@@ -95,6 +98,10 @@ FUNCTION_ALIAS_ANALYSIS("amdgpu-aa", AMDGPUAA())
#ifndef FUNCTION_PASS_WITH_PARAMS
#define FUNCTION_PASS_WITH_PARAMS(NAME, CLASS, CREATE_PASS, PARSER, PARAMS)
#endif
+FUNCTION_PASS_WITH_PARAMS(
+ "amdgpu-lds-buffering", "AMDGPULDSBufferingPass",
+ [=](unsigned MaxBytes) { return AMDGPULDSBufferingPass(*this, MaxBytes); },
+ parseAMDGPULDSBufferingMaxBytes, "max-bytes")
FUNCTION_PASS_WITH_PARAMS(
"amdgpu-atomic-optimizer",
"AMDGPUAtomicOptimizerPass",
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp
index 3c1730397ab96..c60541ffa9a36 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp
@@ -28,6 +28,7 @@
#include "AMDGPU.h"
#include "GCNSubtarget.h"
#include "Utils/AMDGPUBaseInfo.h"
+#include "Utils/AMDGPULDSUtils.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/Analysis/CaptureTracking.h"
#include "llvm/Analysis/InstSimplifyFolder.h"
@@ -1460,128 +1461,24 @@ void AMDGPUPromoteAllocaImpl::analyzePromoteToLDS(AllocaAnalysis &AA) const {
}
bool AMDGPUPromoteAllocaImpl::hasSufficientLocalMem(const Function &F) {
-
- FunctionType *FTy = F.getFunctionType();
- const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
-
- // If the function has any arguments in the local address space, then it's
- // possible these arguments require the entire local memory space, so
- // we cannot use local memory in the pass.
- for (Type *ParamTy : FTy->params()) {
- PointerType *PtrTy = dyn_cast<PointerType>(ParamTy);
- if (PtrTy && PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
- LocalMemLimit = 0;
+ AMDGPU::AMDGPULDSBudget Budget = AMDGPU::computeLDSBudget(F, TM);
+ CurrentLocalMemUsage = Budget.currentUsage;
+ LocalMemLimit = Budget.limit;
+ if (!Budget.promotable) {
+ if (Budget.disabledDueToLocalArg) {
LLVM_DEBUG(dbgs() << "Function has local memory argument. Promoting to "
"local memory disabled.\n");
- return false;
- }
- }
-
- LocalMemLimit = ST.getAddressableLocalMemorySize();
- if (LocalMemLimit == 0)
- return false;
-
- SmallVector<const Constant *, 16> Stack;
- SmallPtrSet<const Constant *, 8> VisitedConstants;
- SmallPtrSet<const GlobalVariable *, 8> UsedLDS;
-
- auto visitUsers = [&](const GlobalVariable *GV, const Constant *Val) -> bool {
- for (const User *U : Val->users()) {
- if (const Instruction *Use = dyn_cast<Instruction>(U)) {
- if (Use->getFunction() == &F)
- return true;
- } else {
- const Constant *C = cast<Constant>(U);
- if (VisitedConstants.insert(C).second)
- Stack.push_back(C);
- }
- }
-
- return false;
- };
-
- for (GlobalVariable &GV : Mod.globals()) {
- if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS)
- continue;
-
- if (visitUsers(&GV, &GV)) {
- UsedLDS.insert(&GV);
- Stack.clear();
- continue;
}
-
- // For any ConstantExpr uses, we need to recursively search the users until
- // we see a function.
- while (!Stack.empty()) {
- const Constant *C = Stack.pop_back_val();
- if (visitUsers(&GV, C)) {
- UsedLDS.insert(&GV);
- Stack.clear();
- break;
- }
- }
- }
-
- SmallVector<std::pair<uint64_t, Align>, 16> AllocatedSizes;
- AllocatedSizes.reserve(UsedLDS.size());
-
- for (const GlobalVariable *GV : UsedLDS) {
- Align Alignment =
- DL.getValueOrABITypeAlignment(GV->getAlign(), GV->getValueType());
- uint64_t AllocSize = GV->getGlobalSize(DL);
-
- // HIP uses an extern unsized array in local address space for dynamically
- // allocated shared memory. In that case, we have to disable the promotion.
- if (GV->hasExternalLinkage() && AllocSize == 0) {
- LocalMemLimit = 0;
+ if (Budget.disabledDueToExternDynShared) {
LLVM_DEBUG(dbgs() << "Function has a reference to externally allocated "
"local memory. Promoting to local memory "
"disabled.\n");
- return false;
}
-
- AllocatedSizes.emplace_back(AllocSize, Alignment);
- }
-
- // Sort to try to estimate the worst case alignment padding
- //
- // FIXME: We should really do something to fix the addresses to a more optimal
- // value instead
- llvm::sort(AllocatedSizes, llvm::less_second());
-
- // Check how much local memory is being used by global objects
- CurrentLocalMemUsage = 0;
-
- // FIXME: Try to account for padding here. The real padding and address is
- // currently determined from the inverse order of uses in the function when
- // legalizing, which could also potentially change. We try to estimate the
- // worst case here, but we probably should fix the addresses earlier.
- for (auto Alloc : AllocatedSizes) {
- CurrentLocalMemUsage = alignTo(CurrentLocalMemUsage, Alloc.second);
- CurrentLocalMemUsage += Alloc.first;
- }
-
- unsigned MaxOccupancy =
- ST.getWavesPerEU(ST.getFlatWorkGroupSizes(F), CurrentLocalMemUsage, F)
- .second;
-
- // Round up to the next tier of usage.
- unsigned MaxSizeWithWaveCount =
- ST.getMaxLocalMemSizeWithWaveCount(MaxOccupancy, F);
-
- // Program may already use more LDS than is usable at maximum occupancy.
- if (CurrentLocalMemUsage > MaxSizeWithWaveCount)
return false;
-
- LocalMemLimit = MaxSizeWithWaveCount;
+ }
LLVM_DEBUG(dbgs() << F.getName() << " uses " << CurrentLocalMemUsage
- << " bytes of LDS\n"
- << " Rounding size to " << MaxSizeWithWaveCount
- << " with a maximum occupancy of " << MaxOccupancy << '\n'
- << " and " << (LocalMemLimit - CurrentLocalMemUsage)
- << " available for promotion\n");
-
+ << " bytes of LDS\n");
return true;
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index c8c5548181a64..e3da368d6f7ea 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -600,6 +600,12 @@ static cl::opt<bool> EnableImageIntrinsicOptimizer(
cl::desc("Enable image intrinsic optimizer pass"), cl::init(true),
cl::Hidden);
+// Gate insertion of the AMDGPU LDS Buffering pass into the default pipeline.
+static cl::opt<bool> EnableLDSBuffering(
+ "amdgpu-enable-lds-buffering",
+ cl::desc("Enable AMDGPU LDS Buffering pass in the default pipeline"),
+ cl::init(false), cl::Hidden);
+
static cl::opt<bool>
EnableLoopPrefetch("amdgpu-loop-prefetch",
cl::desc("Enable loop data prefetch on AMDGPU"),
@@ -722,6 +728,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
initializeAMDGPUPreLegalizerCombinerPass(*PR);
initializeAMDGPURegBankCombinerPass(*PR);
initializeAMDGPUPromoteAllocaPass(*PR);
+ initializeAMDGPULDSBufferingLegacyPass(*PR);
initializeAMDGPUCodeGenPreparePass(*PR);
initializeAMDGPULateCodeGenPrepareLegacyPass(*PR);
initializeAMDGPURemoveIncompatibleFunctionsLegacyPass(*PR);
@@ -994,6 +1001,30 @@ parseAMDGPUAtomicOptimizerStrategy(StringRef Params) {
return make_error<StringError>("invalid parameter", inconvertibleErrorCode());
}
+static Expected<unsigned> parseAMDGPULDSBufferingMaxBytes(StringRef Params) {
+ unsigned Result = 64;
+ while (!Params.empty()) {
+ StringRef Param;
+ std::tie(Param, Params) = Params.split(';');
+ if (Param.empty())
+ continue;
+ if (!Param.consume_front("max-bytes=")) {
+ return make_error<StringError>(
+ formatv("invalid AMDGPULDSBuffering pass parameter '{0}' ", Param)
+ .str(),
+ inconvertibleErrorCode());
+ }
+ unsigned Parsed = 0;
+ if (Param.getAsInteger(10, Parsed)) {
+ return make_error<StringError>(
+ formatv("invalid AMDGPULDSBuffering max-bytes '{0}' ", Param).str(),
+ inconvertibleErrorCode());
+ }
+ Result = Parsed;
+ }
+ return Result;
+}
+
Expected<AMDGPUAttributorOptions>
parseAMDGPUAttributorPassOptions(StringRef Params) {
AMDGPUAttributorOptions Result;
@@ -1618,6 +1649,9 @@ void AMDGPUPassConfig::addIRPasses() {
if (TM.getOptLevel() > CodeGenOptLevel::None) {
addPass(createAMDGPUPromoteAlloca());
+ // Run per-thread LDS buffering after promote-alloca to use leftover LDS.
+ if (TM.getTargetTriple().isAMDGCN() && EnableLDSBuffering)
+ addPass(createAMDGPULDSBufferingLegacyPass());
if (isPassEnabled(EnableScalarIRPasses))
addStraightLineScalarOptimizationPasses();
@@ -2403,6 +2437,9 @@ void AMDGPUCodeGenPassBuilder::addIRPasses(PassManagerWrapper &PMW) {
if (TM.getOptLevel() > CodeGenOptLevel::None) {
addFunctionPass(AMDGPUPromoteAllocaPass(TM), PMW);
+ // Run per-thread LDS buffering after promote-alloca to use leftover LDS.
+ if (TM.getTargetTriple().isAMDGCN() && EnableLDSBuffering)
+ addFunctionPass(AMDGPULDSBufferingPass(getTM()), PMW);
if (isPassEnabled(EnableScalarIRPasses))
addStraightLineScalarOptimizationPasses(PMW);
diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt
index b7e679a69a80d..ea89a7cc616c0 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -79,6 +79,7 @@ add_llvm_target(AMDGPUCodeGen
AMDGPULowerKernelArguments.cpp
AMDGPULowerKernelAttributes.cpp
AMDGPULowerModuleLDSPass.cpp
+ AMDGPULDSBuffering.cpp
AMDGPUPrepareAGPRAlloc.cpp
AMDGPULowerExecSync.cpp
AMDGPUSwLowerLDS.cpp
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp b/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp
new file mode 100644
index 0000000000000..8cc80e1c816b0
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp
@@ -0,0 +1,235 @@
+//===-- AMDGPULDSUtils.cpp - AMDGPU LDS utilities ------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Shared helpers for computing LDS usage and limits for an AMDGPU function.
+//
+//===----------------------------------------------------------------------===//
+
+#include "Utils/AMDGPULDSUtils.h"
+
+#include "AMDGPU.h"
+#include "GCNSubtarget.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/DataLayout.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/GlobalVariable.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Support/Alignment.h"
+
+using namespace llvm;
+
+//===----------------------------------------------------------------------===//
+// Work-group / work-item query IR helpers
+//===----------------------------------------------------------------------===//
+
+namespace {
+
+// Read local size Y/Z from the HSA dispatch packet.
+static std::pair<Value *, Value *>
+getLocalSizeYZFromDispatch(IRBuilderBase &Builder, Module &M,
+ const AMDGPUSubtarget &ST) {
+ Function &F = *Builder.GetInsertBlock()->getParent();
+
+ CallInst *DispatchPtr = cast<CallInst>(
+ Builder.CreateIntrinsic(Intrinsic::amdgcn_dispatch_ptr, {}));
+ DispatchPtr->addRetAttr(Attribute::NoAlias);
+ DispatchPtr->addRetAttr(Attribute::NonNull);
+ F.removeFnAttr("amdgpu-no-dispatch-ptr");
+ DispatchPtr->addDereferenceableRetAttr(64);
+
+ Type *I32Ty = Type::getInt32Ty(M.getContext());
+ Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(I32Ty, DispatchPtr, 1);
+ LoadInst *LoadXY = Builder.CreateAlignedLoad(I32Ty, GEPXY, Align(4));
+ Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(I32Ty, DispatchPtr, 2);
+ LoadInst *LoadZU = Builder.CreateAlignedLoad(I32Ty, GEPZU, Align(4));
+
+ MDNode *MD = MDNode::get(M.getContext(), {});
+ LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);
+ LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);
+ ST.makeLIDRangeMetadata(LoadZU);
+
+ Value *Y = Builder.CreateLShr(LoadXY, 16);
+ return {Y, LoadZU};
+}
+
+} // end anonymous namespace
+
+Value *AMDGPU::getWorkitemID(IRBuilderBase &Builder, Module &M,
+ const AMDGPUSubtarget &ST, unsigned N) {
+ Function *F = Builder.GetInsertBlock()->getParent();
+ Intrinsic::ID IntrID = Intrinsic::not_intrinsic;
+ StringRef AttrName;
+
+ switch (N) {
+ case 0:
+ IntrID = Intrinsic::amdgcn_workitem_id_x;
+ AttrName = "amdgpu-no-workitem-id-x";
+ break;
+ case 1:
+ IntrID = Intrinsic::amdgcn_workitem_id_y;
+ AttrName = "amdgpu-no-workitem-id-y";
+ break;
+ case 2:
+ IntrID = Intrinsic::amdgcn_workitem_id_z;
+ AttrName = "amdgpu-no-workitem-id-z";
+ break;
+ default:
+ llvm_unreachable("invalid dimension");
+ }
+
+ Function *WorkitemIdFn = Intrinsic::getOrInsertDeclaration(&M, IntrID);
+ CallInst *CI = cast<CallInst>(Builder.CreateCall(WorkitemIdFn));
+ ST.makeLIDRangeMetadata(CI);
+ F->removeFnAttr(AttrName);
+ return CI;
+}
+
+Value *AMDGPU::buildLinearThreadId(IRBuilderBase &Builder, Module &M,
+ const AMDGPUSubtarget &ST) {
+ Value *TCntY = nullptr;
+ Value *TCntZ = nullptr;
+ std::tie(TCntY, TCntZ) = getLocalSizeYZFromDispatch(Builder, M, ST);
+ Value *TIdX = getWorkitemID(Builder, M, ST, 0);
+ Value *TIdY = getWorkitemID(Builder, M, ST, 1);
+ Value *TIdZ = getWorkitemID(Builder, M, ST, 2);
+
+ Value *Tmp0 = Builder.CreateMul(TCntY, TCntZ, "", true, true);
+ Tmp0 = Builder.CreateMul(Tmp0, TIdX);
+ Value *Tmp1 = Builder.CreateMul(TIdY, TCntZ, "", true, true);
+ Value *TID = Builder.CreateAdd(Tmp0, Tmp1);
+ TID = Builder.CreateAdd(TID, TIdZ);
+ return TID;
+}
+
+//===----------------------------------------------------------------------===//
+// LDS budget computation
+//===----------------------------------------------------------------------===//
+
+AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
+ const TargetMachine &TM) {
+ AMDGPU::AMDGPULDSBudget Result;
+
+ const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
+ const Module *M = F.getParent();
+ const DataLayout &DL = M->getDataLayout();
+
+ // If the function has any arguments in the local address space, then it's
+ // possible these arguments require the entire local memory space, so
+ // we cannot use local memory.
+ FunctionType *FTy = F.getFunctionType();
+ for (Type *ParamTy : FTy->params()) {
+ PointerType *PtrTy = dyn_cast<PointerType>(ParamTy);
+ if (PtrTy && PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
+ Result.limit = 0;
+ Result.promotable = false;
+ Result.disabledDueToLocalArg = true;
+ return Result;
+ }
+ }
+
+ uint32_t LocalMemLimit = ST.getAddressableLocalMemorySize();
+ if (LocalMemLimit == 0) {
+ Result.limit = 0;
+ Result.promotable = false;
+ return Result;
+ }
+
+ SmallVector<const Constant *, 16> Stack;
+ SmallPtrSet<const Constant *, 8> VisitedConstants;
+ SmallPtrSet<const GlobalVariable *, 8> UsedLDS;
+
+ auto visitUsers = [&](const GlobalVariable *GV, const Constant *Val) -> bool {
+ for (const User *U : Val->users()) {
+ if (const Instruction *Use = dyn_cast<Instruction>(U)) {
+ if (Use->getParent()->getParent() == &F)
+ return true;
+ } else {
+ const Constant *C = cast<Constant>(U);
+ if (VisitedConstants.insert(C).second)
+ Stack.push_back(C);
+ }
+ }
+ return false;
+ };
+
+ for (const GlobalVariable &GV : M->globals()) {
+ if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS)
+ continue;
+
+ if (visitUsers(&GV, &GV)) {
+ UsedLDS.insert(&GV);
+ Stack.clear();
+ continue;
+ }
+
+ while (!Stack.empty()) {
+ const Constant *C = Stack.pop_back_val();
+ if (visitUsers(&GV, C)) {
+ UsedLDS.insert(&GV);
+ Stack.clear();
+ break;
+ }
+ }
+ }
+
+ SmallVector<std::pair<uint64_t, Align>, 16> AllocatedSizes;
+ AllocatedSizes.reserve(UsedLDS.size());
+
+ for (const GlobalVariable *GV : UsedLDS) {
+ Align Alignment =
+ DL.getValueOrABITypeAlignment(GV->getAlign(), GV->getValueType());
+ uint64_t AllocSize = DL.getTypeAllocSize(GV->getValueType());
+
+ // HIP uses an extern unsized array in local address space for dynamically
+ // allocated shared memory.
+ if (GV->hasExternalLinkage() && AllocSize == 0) {
+ Result.limit = 0;
+ Result.promotable = false;
+ Result.disabledDueToExternDynShared = true;
+ return Result;
+ }
+
+ AllocatedSizes.emplace_back(AllocSize, Alignment);
+ }
+
+ // Sort to try to estimate the worst case alignment padding.
+ llvm::sort(AllocatedSizes, llvm::less_second());
+
+ uint32_t CurrentLocalMemUsage = 0;
+ for (const std::pair<uint64_t, Align> &Alloc : AllocatedSizes) {
+ CurrentLocalMemUsage = alignTo(CurrentLocalMemUsage, Alloc.second);
+ CurrentLocalMemUsage += Alloc.first;
+ }
+
+ unsigned MaxOccupancy =
+ ST.getWavesPerEU(ST.getFlatWorkGroupSizes(F), CurrentLocalMemUsage, F)
+ .second;
+
+ unsigned MaxSizeWithWaveCount =
+ ST.getMaxLocalMemSizeWithWaveCount(MaxOccupancy, F);
+
+ if (CurrentLocalMemUsage > MaxSizeWithWaveCount) {
+ Result.currentUsage = CurrentLocalMemUsage;
+ Result.limit = MaxSizeWithWaveCount;
+ Result.maxOccupancy = MaxOccupancy;
+ Result.promotable = false;
+ return Result;
+ }
+
+ Result.currentUsage = CurrentLocalMemUsage;
+ Result.limit = MaxSizeWithWaveCount;
+ Result.maxOccupancy = MaxOccupancy;
+ Result.promotable = true;
+ return Result;
+}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h b/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h
new file mode 100644
index 0000000000000..18690a77647c1
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h
@@ -0,0 +1,53 @@
+//===-- AMDGPULDSUtils.h - AMDGPU LDS utilities ----------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Shared helpers for computing LDS usage and limits for an AMDGPU function.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H
+#define LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H
+
+#include <cstdint>
+#include <utility>
+
+namespace llvm {
+
+class AMDGPUSubtarget;
+class Function;
+class IRBuilderBase;
+class Module;
+class TargetMachine;
+class Value;
+
+namespace AMDGPU {
+
+/// Get workitem id for dimension N (0,1,2).
+Value *getWorkitemID(IRBuilderBase &Builder, Module &M,
+ const AMDGPUSubtarget &ST, unsigned N);
+
+/// Compute linear thread id within a workgroup.
+Value *buildLinearThreadId(IRBuilderBase &Builder, Module &M,
+ const AMDGPUSubtarget &ST);
+
+struct AMDGPULDSBudget {
+ uint32_t currentUsage = 0;
+ uint32_t limit = 0;
+ unsigned maxOccupancy = 0;
+ bool promotable = false;
+ bool disabledDueToLocalArg = false;
+ bool disabledDueToExternDynShared = false;
+};
+
+AMDGPULDSBudget computeLDSBudget(const Function &F, const TargetMachine &TM);
+
+} // end namespace AMDGPU
+
+} // end namespace llvm
+
+#endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H
diff --git a/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt b/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
index 7b2200d8bc488..4633763d23f22 100644
--- a/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
@@ -4,6 +4,7 @@ add_llvm_component_library(LLVMAMDGPUUtils
AMDGPUDelayedMCExpr.cpp
AMDGPUPALMetadata.cpp
AMDKernelCodeTUtils.cpp
+ AMDGPULDSUtils.cpp
LINK_COMPONENTS
Analysis
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-sroa.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-sroa.ll
new file mode 100644
index 0000000000000..2fbbb9a2d7088
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-sroa.ll
@@ -0,0 +1,83 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 6
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='default<O3>' -S < %s | FileCheck %s --check-prefix=OPT
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='default<O3>,function(amdgpu-lds-buffering<max-bytes=64>)' -S < %s | FileCheck %s --check-prefix=LDS
+
+target triple = "amdgcn-amd-amdhsa"
+
+%state = type { <4 x i32>, <4 x i32> }
+
+; SROA splits this partially updated aggregate. The unchanged second field
+; becomes a one-use load/store pair that cannot be removed because %out may
+; alias %state.ptr.
+define amdgpu_kernel void @partial_update(ptr addrspace(1) %state.ptr, ptr addrspace(1) %out) {
+; OPT-LABEL: define amdgpu_kernel void @partial_update(
+; OPT-SAME: ptr addrspace(1) nofree captures(none) [[STATE_PTR:%.*]], ptr addrspace(1) nofree writeonly captures(none) initializes((0, 4)) [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+; OPT-NEXT: [[ENTRY:.*:]]
+; OPT-NEXT: [[STATE_SROA_0_0_COPYLOAD:%.*]] = load <4 x i32>, ptr addrspace(1) [[STATE_PTR]], align 16, !amdgpu.noclobber [[META0:![0-9]+]]
+; OPT-NEXT: [[STATE_SROA_4_0_STATE_PTR_SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[STATE_PTR]], i64 16
+; OPT-NEXT: [[STATE_SROA_4_0_COPYLOAD:%.*]] = load <4 x i32>, ptr addrspace(1) [[STATE_SROA_4_0_STATE_PTR_SROA_IDX]], align 16, !amdgpu.noclobber [[META0]]
+; OPT-NEXT: [[UPDATED:%.*]] = xor <4 x i32> [[STATE_SROA_0_0_COPYLOAD]], splat (i32 1)
+; OPT-NEXT: [[VALUE:%.*]] = extractelement <4 x i32> [[UPDATED]], i64 0
+; OPT-NEXT: store i32 [[VALUE]], ptr addrspace(1) [[OUT]], align 4
+; OPT-NEXT: store <4 x i32> [[UPDATED]], ptr addrspace(1) [[STATE_PTR]], align 16
+; OPT-NEXT: store <4 x i32> [[STATE_SROA_4_0_COPYLOAD]], ptr addrspace(1) [[STATE_SROA_4_0_STATE_PTR_SROA_IDX]], align 16
+; OPT-NEXT: ret void
+;
+; LDS-LABEL: define amdgpu_kernel void @partial_update(
+; LDS-SAME: ptr addrspace(1) nofree captures(none) [[STATE_PTR:%.*]], ptr addrspace(1) nofree writeonly captures(none) initializes((0, 4)) [[OUT:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+; LDS-NEXT: [[ENTRY:.*:]]
+; LDS-NEXT: [[STATE_SROA_0_0_COPYLOAD:%.*]] = load <4 x i32>, ptr addrspace(1) [[STATE_PTR]], align 16, !amdgpu.noclobber [[META0:![0-9]+]]
+; LDS-NEXT: [[STATE_SROA_4_0_STATE_PTR_SROA_IDX:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[STATE_PTR]], i64 16
+; LDS-NEXT: [[TMP0:%.*]] = call noalias nonnull dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+; LDS-NEXT: [[TMP1:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP0]], i64 1
+; LDS-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !invariant.load [[META0]]
+; LDS-NEXT: [[TMP3:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[TMP0]], i64 2
+; LDS-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(4) [[TMP3]], align 4, !range [[RNG1:![0-9]+]], !invariant.load [[META0]]
+; LDS-NEXT: [[TMP5:%.*]] = lshr i32 [[TMP2]], 16
+; LDS-NEXT: [[TMP6:%.*]] = call range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.x()
+; LDS-NEXT: [[TMP7:%.*]] = call range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.y()
+; LDS-NEXT: [[TMP8:%.*]] = call range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.z()
+; LDS-NEXT: [[TMP9:%.*]] = mul nuw nsw i32 [[TMP5]], [[TMP4]]
+; LDS-NEXT: [[TMP10:%.*]] = mul i32 [[TMP9]], [[TMP6]]
+; LDS-NEXT: [[TMP11:%.*]] = mul nuw nsw i32 [[TMP7]], [[TMP4]]
+; LDS-NEXT: [[TMP12:%.*]] = add i32 [[TMP10]], [[TMP11]]
+; LDS-NEXT: [[TMP13:%.*]] = add i32 [[TMP12]], [[TMP8]]
+; LDS-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1024 x <4 x i32>], ptr addrspace(3) @partial_update.ldsbuf, i32 0, i32 [[TMP13]]
+; LDS-NEXT: call void @llvm.memcpy.p3.p1.i64(ptr addrspace(3) align 16 [[TMP14]], ptr addrspace(1) align 16 [[STATE_SROA_4_0_STATE_PTR_SROA_IDX]], i64 16, i1 false)
+; LDS-NEXT: [[UPDATED:%.*]] = xor <4 x i32> [[STATE_SROA_0_0_COPYLOAD]], splat (i32 1)
+; LDS-NEXT: [[VALUE:%.*]] = extractelement <4 x i32> [[UPDATED]], i64 0
+; LDS-NEXT: store i32 [[VALUE]], ptr addrspace(1) [[OUT]], align 4
+; LDS-NEXT: store <4 x i32> [[UPDATED]], ptr addrspace(1) [[STATE_PTR]], align 16
+; LDS-NEXT: call void @llvm.memcpy.p1.p3.i64(ptr addrspace(1) align 16 [[STATE_SROA_4_0_STATE_PTR_SROA_IDX]], ptr addrspace(3) align 16 [[TMP14]], i64 16, i1 false)
+; LDS-NEXT: ret void
+;
+entry:
+ %state = alloca %state, align 16, addrspace(5)
+ call void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) align 16 %state,
+ ptr addrspace(1) align 16 %state.ptr,
+ i64 32, i1 false)
+ %updated.ptr = getelementptr inbounds %state, ptr addrspace(5) %state,
+ i32 0, i32 0
+ %old = load <4 x i32>, ptr addrspace(5) %updated.ptr, align 16
+ %updated = xor <4 x i32> %old, splat (i32 1)
+ store <4 x i32> %updated, ptr addrspace(5) %updated.ptr, align 16
+ %value = extractelement <4 x i32> %updated, i64 0
+ store i32 %value, ptr addrspace(1) %out, align 4
+ call void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) align 16 %state.ptr,
+ ptr addrspace(5) align 16 %state,
+ i64 32, i1 false)
+ ret void
+}
+
+declare void @llvm.memcpy.p5.p1.i64(ptr addrspace(5) noalias nocapture writeonly,
+ ptr addrspace(1) noalias nocapture readonly,
+ i64, i1 immarg)
+declare void @llvm.memcpy.p1.p5.i64(ptr addrspace(1) noalias nocapture writeonly,
+ ptr addrspace(5) noalias nocapture readonly,
+ i64, i1 immarg)
+;.
+; OPT: [[META0]] = !{}
+;.
+; LDS: [[META0]] = !{}
+; LDS: [[RNG1]] = !{i32 0, i32 1025}
+;.
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering.ll
new file mode 100644
index 0000000000000..64165c3cb1e40
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering.ll
@@ -0,0 +1,183 @@
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64>' -S %s | FileCheck %s
+
+%state = type { <4 x i32>, <4 x i32> }
+
+;===---------------------------------------------------------------------===//
+; Positive cases
+;===---------------------------------------------------------------------===//
+
+; CHECK: @ldsbuf_test.ldsbuf = internal unnamed_addr addrspace(3) global
+; CHECK: @ldsbuf_complex.ldsbuf = internal unnamed_addr addrspace(3) global
+; CHECK: @ldsbuf_rocrand_preserved_subsequence.ldsbuf = internal unnamed_addr addrspace(3) global
+
+; Basic positive case: a single global load whose only use is a store back to
+; the same pointer. Include an intervening (potentially-aliasing) global store
+; to keep the value live and avoid trivially folding away.
+; CHECK-LABEL: @ldsbuf_test(
+; CHECK: %[[SLOT_BASIC:[^ ]+]] = getelementptr inbounds {{.*}}, ptr addrspace(3) @ldsbuf_test.ldsbuf, i32 0, i32 %
+; CHECK: call void @llvm.memcpy.p3.p1.i64(ptr addrspace(3){{.*}}%[[SLOT_BASIC]], ptr addrspace(1){{.*}}%p, i64 16, i1 false)
+; CHECK: call void @llvm.memcpy.p1.p3.i64(ptr addrspace(1){{.*}}%p, ptr addrspace(3){{.*}}%[[SLOT_BASIC]], i64 16, i1 false)
+define amdgpu_kernel void @ldsbuf_test(ptr addrspace(1) %p, ptr addrspace(1) %q) #0 {
+entry:
+ %ld = load <4 x i32>, ptr addrspace(1) %p, align 16
+ store i32 0, ptr addrspace(1) %q, align 4
+ store <4 x i32> %ld, ptr addrspace(1) %p, align 16
+ ret void
+}
+
+; "Complexity" positive case: keep the loaded value live across control flow
+; and additional global memory operations, then store back to the same pointer.
+; CHECK-LABEL: @ldsbuf_complex(
+; CHECK: %[[SLOT_COMPLEX:[^ ]+]] = getelementptr inbounds {{.*}}, ptr addrspace(3) @ldsbuf_complex.ldsbuf, i32 0, i32 %
+; CHECK: call void @llvm.memcpy.p3.p1.i64(ptr addrspace(3){{.*}}%[[SLOT_COMPLEX]], ptr addrspace(1){{.*}}%p, i64 16, i1 false)
+; CHECK: call void @llvm.memcpy.p1.p3.i64(ptr addrspace(1){{.*}}%p, ptr addrspace(3){{.*}}%[[SLOT_COMPLEX]], i64 16, i1 false)
+define amdgpu_kernel void @ldsbuf_complex(ptr addrspace(1) %p, ptr addrspace(1) %q, i1 %c) #0 {
+entry:
+ %ld = load <4 x i32>, ptr addrspace(1) %p, align 16
+ br i1 %c, label %then, label %else
+
+then:
+ store i32 1, ptr addrspace(1) %q, align 4
+ br label %merge
+
+else:
+ store i32 2, ptr addrspace(1) %q, align 4
+ br label %merge
+
+merge:
+ store <4 x i32> %ld, ptr addrspace(1) %p, align 16
+ ret void
+}
+
+; rocRand-inspired positive case: model a per-thread RNG state where part of
+; the state (a preserved 16B "subsequence" field) is loaded and stored back
+; unchanged, but kept live across a non-trivial loop body.
+; CHECK-LABEL: @ldsbuf_rocrand_preserved_subsequence(
+; CHECK: %[[TAILPTR:[^ ]+]] = getelementptr inbounds %state, ptr addrspace(1) %state_ptr, i32 0, i32 1
+; CHECK: %[[SLOT_TAIL:[^ ]+]] = getelementptr inbounds {{.*}}, ptr addrspace(3) @ldsbuf_rocrand_preserved_subsequence.ldsbuf, i32 0, i32 %
+; CHECK: call void @llvm.memcpy.p3.p1.i64(ptr addrspace(3){{.*}}%[[SLOT_TAIL]], ptr addrspace(1){{.*}}%[[TAILPTR]], i64 16, i1 false)
+; CHECK: call void @llvm.memcpy.p1.p3.i64(ptr addrspace(1){{.*}}%[[TAILPTR]], ptr addrspace(3){{.*}}%[[SLOT_TAIL]], i64 16, i1 false)
+define amdgpu_kernel void @ldsbuf_rocrand_preserved_subsequence(ptr addrspace(1) %state_ptr,
+ ptr addrspace(1) %out) #0 {
+entry:
+ %tailptr = getelementptr inbounds %state, ptr addrspace(1) %state_ptr, i32 0, i32 1
+ %tail = load <4 x i32>, ptr addrspace(1) %tailptr, align 16
+
+ br label %loop
+
+loop:
+ %i = phi i32 [ 0, %entry ], [ %i.next, %loop ]
+ %acc = phi i32 [ 1, %entry ], [ %acc.next, %loop ]
+
+ %shl = shl i32 %acc, 6
+ %xor = xor i32 %shl, %acc
+ %shr = lshr i32 %xor, 13
+ %acc.next = xor i32 %shr, 1234567
+
+ %idx = zext i32 %i to i64
+ %outp = getelementptr inbounds float, ptr addrspace(1) %out, i64 %idx
+ %f = uitofp i32 %acc.next to float
+ store float %f, ptr addrspace(1) %outp, align 4
+
+ %i.next = add nuw i32 %i, 1
+ %done = icmp eq i32 %i.next, 8
+ br i1 %done, label %exit, label %loop
+
+exit:
+ store <4 x i32> %tail, ptr addrspace(1) %tailptr, align 16
+ ret void
+}
+
+;===---------------------------------------------------------------------===//
+; Negative coverage: patterns that must NOT be transformed.
+;===---------------------------------------------------------------------===//
+
+; Negative: atomic operations are excluded (must not transform).
+; CHECK-LABEL: @ldsbuf_atomic(
+; CHECK: load atomic <4 x i32>, ptr addrspace(1) %p unordered, align 16
+; CHECK: store atomic <4 x i32> {{.*}}, ptr addrspace(1) %p unordered, align 16
+define amdgpu_kernel void @ldsbuf_atomic(ptr addrspace(1) %p) #0 {
+entry:
+ %ld = load atomic <4 x i32>, ptr addrspace(1) %p unordered, align 16
+ store atomic <4 x i32> %ld, ptr addrspace(1) %p unordered, align 16
+ ret void
+}
+
+; Negative: volatile operations are excluded (must not transform).
+; CHECK-LABEL: @ldsbuf_volatile(
+; CHECK: load volatile <4 x i32>, ptr addrspace(1) %p, align 16
+; CHECK: store volatile <4 x i32> {{.*}}, ptr addrspace(1) %p, align 16
+define amdgpu_kernel void @ldsbuf_volatile(ptr addrspace(1) %p) #0 {
+entry:
+ %ld = load volatile <4 x i32>, ptr addrspace(1) %p, align 16
+ store volatile <4 x i32> %ld, ptr addrspace(1) %p, align 16
+ ret void
+}
+
+; Negative: alignment requirement (min 16) must be met.
+; CHECK-LABEL: @ldsbuf_misaligned(
+; CHECK: load <4 x i32>, ptr addrspace(1) %p, align 8
+; CHECK: store <4 x i32> {{.*}}, ptr addrspace(1) %p, align 8
+define amdgpu_kernel void @ldsbuf_misaligned(ptr addrspace(1) %p) #0 {
+entry:
+ %ld = load <4 x i32>, ptr addrspace(1) %p, align 8
+ store <4 x i32> %ld, ptr addrspace(1) %p, align 8
+ ret void
+}
+
+; Negative: size limit (max-bytes=64) must be respected.
+; CHECK-LABEL: @ldsbuf_too_large(
+; CHECK: load <20 x i32>, ptr addrspace(1) %p, align 16
+; CHECK: store <20 x i32> {{.*}}, ptr addrspace(1) %p, align 16
+define amdgpu_kernel void @ldsbuf_too_large(ptr addrspace(1) %p) #0 {
+entry:
+ %ld = load <20 x i32>, ptr addrspace(1) %p, align 16
+ store <20 x i32> %ld, ptr addrspace(1) %p, align 16
+ ret void
+}
+
+; Negative: non-global pointers are excluded (only addrspace(1) is supported).
+; CHECK-LABEL: @ldsbuf_non_global_ptr(
+; CHECK: load <4 x i32>, ptr %p, align 16
+; CHECK: store <4 x i32> {{.*}}, ptr %p, align 16
+define amdgpu_kernel void @ldsbuf_non_global_ptr(ptr %p) #0 {
+entry:
+ %ld = load <4 x i32>, ptr %p, align 16
+ store <4 x i32> %ld, ptr %p, align 16
+ ret void
+}
+
+; Negative: the load must have exactly one use (the final store).
+; CHECK-LABEL: @ldsbuf_multiple_uses(
+; CHECK: %ld = load <4 x i32>, ptr addrspace(1) %p, align 16
+; CHECK: extractelement <4 x i32> %ld, i32 0
+; CHECK: store <4 x i32> %ld, ptr addrspace(1) %p, align 16
+define amdgpu_kernel void @ldsbuf_multiple_uses(ptr addrspace(1) %p, ptr addrspace(1) %q) #0 {
+entry:
+ %ld = load <4 x i32>, ptr addrspace(1) %p, align 16
+ %e = extractelement <4 x i32> %ld, i32 0
+ store i32 %e, ptr addrspace(1) %q, align 4
+ store <4 x i32> %ld, ptr addrspace(1) %p, align 16
+ ret void
+}
+
+; Negative: budget rejection. With large pre-existing LDS usage and large
+; workgroup size, per-thread slots should be rejected by the LDS budget check.
+ at ldsbuf_budget_reject.big = internal addrspace(3) global [8192 x i32] zeroinitializer, align 16
+
+; CHECK-LABEL: @ldsbuf_budget_reject(
+; CHECK: load <16 x i32>, ptr addrspace(1) %p, align 16
+; CHECK: store <16 x i32> {{.*}}, ptr addrspace(1) %p, align 16
+define amdgpu_kernel void @ldsbuf_budget_reject(ptr addrspace(1) %p) #1 {
+entry:
+ %x = load i32, ptr addrspace(3) getelementptr inbounds ([8192 x i32], ptr addrspace(3) @ldsbuf_budget_reject.big, i32 0, i32 0), align 16
+ call void @llvm.donothing()
+ %ld = load <16 x i32>, ptr addrspace(1) %p, align 16
+ store <16 x i32> %ld, ptr addrspace(1) %p, align 16
+ ret void
+}
+
+declare void @llvm.donothing()
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "uniform-work-group-size"="true" }
+attributes #1 = { "amdgpu-flat-work-group-size"="1024,1024" "uniform-work-group-size"="true" }
>From 931965f91c627c87ba11e957f954cc6128dda20a Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Fri, 28 Aug 2026 14:41:56 -0400
Subject: [PATCH 2/3] AMDGPU: Fix LDS buffering resource accounting
Use the allocated type size and overflow-safe arithmetic when reserving per-thread LDS slots. Account for alignment padding and reject unsupported targets.\n\nShare thread index construction with promote-alloca and cover the resource and target boundaries.
---
llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp | 58 +++---
.../AMDGPU/{Utils => }/AMDGPULDSUtils.cpp | 140 +++++++++-----
.../AMDGPU/{Utils => }/AMDGPULDSUtils.h | 32 ++--
.../lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp | 181 +++---------------
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 4 +-
llvm/lib/Target/AMDGPU/CMakeLists.txt | 1 +
llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt | 1 -
.../CodeGen/AMDGPU/lds-buffering-budget.ll | 50 +++++
.../AMDGPU/lds-buffering-invalid-params.ll | 9 +
.../CodeGen/AMDGPU/lds-buffering-layout.ll | 22 +++
.../CodeGen/AMDGPU/lds-buffering-targets.ll | 23 +++
llvm/test/CodeGen/AMDGPU/lds-buffering.ll | 3 +-
12 files changed, 266 insertions(+), 258 deletions(-)
rename llvm/lib/Target/AMDGPU/{Utils => }/AMDGPULDSUtils.cpp (59%)
rename llvm/lib/Target/AMDGPU/{Utils => }/AMDGPULDSUtils.h (55%)
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering-budget.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering-layout.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering-targets.ll
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp b/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
index 03d3482c247fe..11de8d98a8240 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
@@ -24,8 +24,7 @@
// Current implementation handles the simplest pattern: a load from global
// memory whose only use is a store back to the same pointer. This pattern is
// transformed into a pair of memcpy operations (global->LDS and LDS->global),
-// effectively moving the value through LDS instead of accessing global memory
-// directly.
+// parking the value in LDS to shorten its VGPR live range.
//
// This pass was inspired by finding that some rocrand performance tests
// show better performance when global memory is buffered through LDS
@@ -37,22 +36,21 @@
//===----------------------------------------------------------------------===//
#include "AMDGPU.h"
+#include "AMDGPULDSUtils.h"
#include "AMDGPUTargetMachine.h"
#include "GCNSubtarget.h"
#include "Utils/AMDGPUBaseInfo.h"
-#include "Utils/AMDGPULDSUtils.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/CodeGen/TargetPassConfig.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instructions.h"
-#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/PassManager.h"
-#include "llvm/IR/PatternMatch.h"
#include "llvm/InitializePasses.h"
#include "llvm/Pass.h"
#include "llvm/Support/Alignment.h"
#include "llvm/Support/Debug.h"
#include <algorithm>
+#include <utility>
#define DEBUG_TYPE "amdgpu-lds-buffering"
@@ -72,18 +70,16 @@ class AMDGPULDSBufferingImpl {
bool run(Function &F) {
LLVM_DEBUG(dbgs() << "[LDSBuffer] Visit function: " << F.getName() << '\n');
- if (!AMDGPU::isEntryFunctionCC(F.getCallingConv()))
+ if (!TM.getTargetTriple().isAMDGCN() ||
+ !AMDGPU::isEntryFunctionCC(F.getCallingConv()))
return false;
Mod = F.getParent();
DL = &Mod->getDataLayout();
AMDGPU::AMDGPULDSBudget Budget = AMDGPU::computeLDSBudget(F, TM);
- if (!Budget.promotable)
+ if (!Budget.Promotable)
return false;
- uint32_t localUsage = Budget.currentUsage;
- uint32_t localLimit = Budget.limit;
-
const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(F).second;
@@ -121,12 +117,14 @@ class AMDGPULDSBufferingImpl {
if (SPtr != Ptr)
continue;
- TypeSize TS = DL->getTypeStoreSize(ValTy);
- if (TS.isScalable())
+ TypeSize StoreSize = DL->getTypeStoreSize(ValTy);
+ TypeSize AllocSize = DL->getTypeAllocSize(ValTy);
+ if (StoreSize.isScalable() || AllocSize.isScalable())
continue;
- uint64_t Size = TS.getFixedValue();
- if (Size == 0 || Size > MaxBytes)
+ uint64_t CopySize = StoreSize.getFixedValue();
+ if (CopySize == 0 || CopySize > MaxBytes)
continue;
+ uint64_t SlotSize = AllocSize.getFixedValue();
Align MinAlign = Align(16);
Align LoadAlign = LI->getAlign();
Align StoreAlign = SI->getAlign();
@@ -138,7 +136,7 @@ class AMDGPULDSBufferingImpl {
LLVM_DEBUG({
dbgs() << "[LDSBuffer] Candidate found: load->store same ptr in "
<< F.getName() << '\n';
- dbgs() << " size=" << Size
+ dbgs() << " size=" << CopySize
<< "B, loadAlign=" << LoadAlign.value()
<< ", storeAlign=" << StoreAlign.value()
<< ", chosenAlign=" << Alignment.value()
@@ -147,25 +145,24 @@ class AMDGPULDSBufferingImpl {
IRBuilder<> BLoad(LI);
// Ensure LDS budget allows allocating a per-thread slot.
- uint32_t NewSize = alignTo(localUsage, Alignment);
- NewSize += WorkGroupSize * static_cast<uint32_t>(Size);
- if (NewSize > localLimit)
+ if (WorkGroupSize == 0 || SlotSize > Budget.Limit / WorkGroupSize)
+ continue;
+ if (!Budget.tryReserve(SlotSize * WorkGroupSize, Alignment))
continue;
- localUsage = NewSize;
- auto [GV, SlotPtr] =
- createLDSGlobalAndThreadSlot(F, ValTy, Alignment, "ldsbuf", BLoad);
+ auto [GV, SlotPtr] = createLDSGlobalAndThreadSlot(
+ F, ValTy, WorkGroupSize, Alignment, BLoad);
// memcpy p3 <- p1
LLVM_DEBUG(dbgs() << "[LDSBuffer] Insert memcpy global->LDS: "
- << GV->getName() << ", bytes=" << Size
+ << GV->getName() << ", bytes=" << CopySize
<< ", align=" << Alignment.value() << '\n');
- BLoad.CreateMemCpy(SlotPtr, Alignment, Ptr, LoadAlign, TS);
+ BLoad.CreateMemCpy(SlotPtr, Alignment, Ptr, LoadAlign, StoreSize);
// Replace the final store with memcpy LDS->global.
IRBuilder<> BStore(SI);
LLVM_DEBUG(dbgs() << "[LDSBuffer] Insert memcpy LDS->global: "
- << GV->getName() << ", bytes=" << Size
+ << GV->getName() << ", bytes=" << CopySize
<< ", align=" << Alignment.value() << '\n');
- BStore.CreateMemCpy(SPtr, StoreAlign, SlotPtr, Alignment, TS);
+ BStore.CreateMemCpy(SPtr, StoreAlign, SlotPtr, Alignment, StoreSize);
ToErase.push_back(SI);
ToErase.push_back(LI);
@@ -188,14 +185,13 @@ class AMDGPULDSBufferingImpl {
// Create an LDS array [WGSize x ElemTy] and return pointer to per-thread
// slot.
std::pair<GlobalVariable *, Value *>
- createLDSGlobalAndThreadSlot(Function &F, Type *ElemTy, Align Alignment,
- StringRef BaseName, IRBuilder<> &Builder) {
- const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
- unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(F).second;
+ createLDSGlobalAndThreadSlot(Function &F, Type *ElemTy,
+ unsigned WorkGroupSize, Align Alignment,
+ IRBuilder<> &Builder) {
Type *ArrTy = ArrayType::get(ElemTy, WorkGroupSize);
GlobalVariable *GV = new GlobalVariable(
*Mod, ArrTy, /*isConstant=*/false, GlobalValue::InternalLinkage,
- PoisonValue::get(ArrTy), (F.getName() + "." + BaseName).str(), nullptr,
+ PoisonValue::get(ArrTy), (F.getName() + ".ldsbuf").str(), nullptr,
GlobalVariable::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS);
GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
GV->setAlignment(Alignment);
@@ -206,7 +202,7 @@ class AMDGPULDSBufferingImpl {
<< ", align=" << Alignment.value() << '\n';
});
- Value *LinearTID = AMDGPU::buildLinearThreadId(Builder, *Mod, ST);
+ Value *LinearTID = AMDGPU::buildLinearThreadId(Builder, TM);
LLVMContext &Ctx = Mod->getContext();
Value *Indices[] = {Constant::getNullValue(Type::getInt32Ty(Ctx)),
LinearTID};
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp b/llvm/lib/Target/AMDGPU/AMDGPULDSUtils.cpp
similarity index 59%
rename from llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp
rename to llvm/lib/Target/AMDGPU/AMDGPULDSUtils.cpp
index 8cc80e1c816b0..5299a81bf5f2a 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULDSUtils.cpp
@@ -1,4 +1,4 @@
-//===-- AMDGPULDSUtils.cpp - AMDGPU LDS utilities ------------------------===//
+//===-- AMDGPULDSUtils.cpp - AMDGPU LDS utilities -------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -10,10 +10,9 @@
//
//===----------------------------------------------------------------------===//
-#include "Utils/AMDGPULDSUtils.h"
+#include "AMDGPULDSUtils.h"
-#include "AMDGPU.h"
-#include "GCNSubtarget.h"
+#include "AMDGPUSubtarget.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/SmallVector.h"
@@ -24,8 +23,12 @@
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicsAMDGPU.h"
+#include "llvm/IR/IntrinsicsR600.h"
#include "llvm/IR/Module.h"
+#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/Alignment.h"
+#include "llvm/Target/TargetMachine.h"
+#include <utility>
using namespace llvm;
@@ -35,14 +38,25 @@ using namespace llvm;
namespace {
-// Read local size Y/Z from the HSA dispatch packet.
-static std::pair<Value *, Value *>
-getLocalSizeYZFromDispatch(IRBuilderBase &Builder, Module &M,
- const AMDGPUSubtarget &ST) {
+static std::pair<Value *, Value *> getLocalSizeYZ(IRBuilderBase &Builder,
+ const TargetMachine &TM,
+ const AMDGPUSubtarget &ST) {
Function &F = *Builder.GetInsertBlock()->getParent();
+ Module &M = *F.getParent();
- CallInst *DispatchPtr = cast<CallInst>(
- Builder.CreateIntrinsic(Intrinsic::amdgcn_dispatch_ptr, {}));
+ if (TM.getTargetTriple().getOS() != Triple::AMDHSA) {
+ CallInst *LocalSizeY = Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::r600_read_local_size_y, {});
+ CallInst *LocalSizeZ = Builder.CreateIntrinsicWithoutFolding(
+ Intrinsic::r600_read_local_size_z, {});
+
+ ST.makeLIDRangeMetadata(LocalSizeY);
+ ST.makeLIDRangeMetadata(LocalSizeZ);
+ return {LocalSizeY, LocalSizeZ};
+ }
+
+ CallInst *DispatchPtr =
+ Builder.CreateIntrinsicWithoutFolding(Intrinsic::amdgcn_dispatch_ptr, {});
DispatchPtr->addRetAttr(Attribute::NoAlias);
DispatchPtr->addRetAttr(Attribute::NonNull);
F.removeFnAttr("amdgpu-no-dispatch-ptr");
@@ -63,46 +77,55 @@ getLocalSizeYZFromDispatch(IRBuilderBase &Builder, Module &M,
return {Y, LoadZU};
}
-} // end anonymous namespace
-
-Value *AMDGPU::getWorkitemID(IRBuilderBase &Builder, Module &M,
- const AMDGPUSubtarget &ST, unsigned N) {
+static Value *getWorkitemID(IRBuilderBase &Builder, const TargetMachine &TM,
+ const AMDGPUSubtarget &ST, unsigned N) {
Function *F = Builder.GetInsertBlock()->getParent();
Intrinsic::ID IntrID = Intrinsic::not_intrinsic;
StringRef AttrName;
switch (N) {
case 0:
- IntrID = Intrinsic::amdgcn_workitem_id_x;
+ IntrID = TM.getTargetTriple().isAMDGCN()
+ ? static_cast<Intrinsic::ID>(Intrinsic::amdgcn_workitem_id_x)
+ : static_cast<Intrinsic::ID>(Intrinsic::r600_read_tidig_x);
AttrName = "amdgpu-no-workitem-id-x";
break;
case 1:
- IntrID = Intrinsic::amdgcn_workitem_id_y;
+ IntrID = TM.getTargetTriple().isAMDGCN()
+ ? static_cast<Intrinsic::ID>(Intrinsic::amdgcn_workitem_id_y)
+ : static_cast<Intrinsic::ID>(Intrinsic::r600_read_tidig_y);
AttrName = "amdgpu-no-workitem-id-y";
break;
case 2:
- IntrID = Intrinsic::amdgcn_workitem_id_z;
+ IntrID = TM.getTargetTriple().isAMDGCN()
+ ? static_cast<Intrinsic::ID>(Intrinsic::amdgcn_workitem_id_z)
+ : static_cast<Intrinsic::ID>(Intrinsic::r600_read_tidig_z);
AttrName = "amdgpu-no-workitem-id-z";
break;
default:
llvm_unreachable("invalid dimension");
}
- Function *WorkitemIdFn = Intrinsic::getOrInsertDeclaration(&M, IntrID);
- CallInst *CI = cast<CallInst>(Builder.CreateCall(WorkitemIdFn));
+ Function *WorkitemIdFn =
+ Intrinsic::getOrInsertDeclaration(F->getParent(), IntrID);
+ CallInst *CI = Builder.CreateCall(WorkitemIdFn);
ST.makeLIDRangeMetadata(CI);
F->removeFnAttr(AttrName);
return CI;
}
-Value *AMDGPU::buildLinearThreadId(IRBuilderBase &Builder, Module &M,
- const AMDGPUSubtarget &ST) {
+} // end anonymous namespace
+
+Value *AMDGPU::buildLinearThreadId(IRBuilderBase &Builder,
+ const TargetMachine &TM) {
+ Function &F = *Builder.GetInsertBlock()->getParent();
+ const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
Value *TCntY = nullptr;
Value *TCntZ = nullptr;
- std::tie(TCntY, TCntZ) = getLocalSizeYZFromDispatch(Builder, M, ST);
- Value *TIdX = getWorkitemID(Builder, M, ST, 0);
- Value *TIdY = getWorkitemID(Builder, M, ST, 1);
- Value *TIdZ = getWorkitemID(Builder, M, ST, 2);
+ std::tie(TCntY, TCntZ) = getLocalSizeYZ(Builder, TM, ST);
+ Value *TIdX = getWorkitemID(Builder, TM, ST, 0);
+ Value *TIdY = getWorkitemID(Builder, TM, ST, 1);
+ Value *TIdZ = getWorkitemID(Builder, TM, ST, 2);
Value *Tmp0 = Builder.CreateMul(TCntY, TCntZ, "", true, true);
Tmp0 = Builder.CreateMul(Tmp0, TIdX);
@@ -116,6 +139,25 @@ Value *AMDGPU::buildLinearThreadId(IRBuilderBase &Builder, Module &M,
// LDS budget computation
//===----------------------------------------------------------------------===//
+bool AMDGPU::AMDGPULDSBudget::tryReserve(uint64_t AllocSize, Align Alignment) {
+ if (!Promotable || CurrentUsage > Limit)
+ return false;
+
+ // The backend may allocate LDS globals in a different order than the IR
+ // pass visits them. Reserve the maximum possible leading padding so the
+ // budget remains valid for any order.
+ uint64_t Padding = CurrentUsage == 0 ? 0 : Alignment.value() - 1;
+ if (Padding > Limit - CurrentUsage)
+ return false;
+
+ uint64_t PaddedUsage = CurrentUsage + Padding;
+ if (AllocSize > Limit - PaddedUsage)
+ return false;
+
+ CurrentUsage = PaddedUsage + AllocSize;
+ return true;
+}
+
AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
const TargetMachine &TM) {
AMDGPU::AMDGPULDSBudget Result;
@@ -131,25 +173,20 @@ AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
for (Type *ParamTy : FTy->params()) {
PointerType *PtrTy = dyn_cast<PointerType>(ParamTy);
if (PtrTy && PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
- Result.limit = 0;
- Result.promotable = false;
- Result.disabledDueToLocalArg = true;
+ Result.DisabledDueToLocalArg = true;
return Result;
}
}
uint32_t LocalMemLimit = ST.getAddressableLocalMemorySize();
- if (LocalMemLimit == 0) {
- Result.limit = 0;
- Result.promotable = false;
+ if (LocalMemLimit == 0)
return Result;
- }
SmallVector<const Constant *, 16> Stack;
SmallPtrSet<const Constant *, 8> VisitedConstants;
SmallPtrSet<const GlobalVariable *, 8> UsedLDS;
- auto visitUsers = [&](const GlobalVariable *GV, const Constant *Val) -> bool {
+ auto VisitUsers = [&](const Constant *Val) -> bool {
for (const User *U : Val->users()) {
if (const Instruction *Use = dyn_cast<Instruction>(U)) {
if (Use->getParent()->getParent() == &F)
@@ -167,7 +204,7 @@ AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS)
continue;
- if (visitUsers(&GV, &GV)) {
+ if (VisitUsers(&GV)) {
UsedLDS.insert(&GV);
Stack.clear();
continue;
@@ -175,7 +212,7 @@ AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
while (!Stack.empty()) {
const Constant *C = Stack.pop_back_val();
- if (visitUsers(&GV, C)) {
+ if (VisitUsers(C)) {
UsedLDS.insert(&GV);
Stack.clear();
break;
@@ -194,9 +231,7 @@ AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
// HIP uses an extern unsized array in local address space for dynamically
// allocated shared memory.
if (GV->hasExternalLinkage() && AllocSize == 0) {
- Result.limit = 0;
- Result.promotable = false;
- Result.disabledDueToExternDynShared = true;
+ Result.DisabledDueToExternDynShared = true;
return Result;
}
@@ -206,30 +241,33 @@ AMDGPU::AMDGPULDSBudget AMDGPU::computeLDSBudget(const Function &F,
// Sort to try to estimate the worst case alignment padding.
llvm::sort(AllocatedSizes, llvm::less_second());
- uint32_t CurrentLocalMemUsage = 0;
+ Result.Limit = LocalMemLimit;
+ Result.Promotable = true;
for (const std::pair<uint64_t, Align> &Alloc : AllocatedSizes) {
- CurrentLocalMemUsage = alignTo(CurrentLocalMemUsage, Alloc.second);
- CurrentLocalMemUsage += Alloc.first;
+ uint64_t NewUsage = alignTo(Result.CurrentUsage, Alloc.second);
+ if (NewUsage > Result.Limit || Alloc.first > Result.Limit - NewUsage) {
+ Result.Promotable = false;
+ return Result;
+ }
+ Result.CurrentUsage = NewUsage + Alloc.first;
}
unsigned MaxOccupancy =
- ST.getWavesPerEU(ST.getFlatWorkGroupSizes(F), CurrentLocalMemUsage, F)
+ ST.getWavesPerEU(ST.getFlatWorkGroupSizes(F),
+ static_cast<uint32_t>(Result.CurrentUsage), F)
.second;
unsigned MaxSizeWithWaveCount =
ST.getMaxLocalMemSizeWithWaveCount(MaxOccupancy, F);
- if (CurrentLocalMemUsage > MaxSizeWithWaveCount) {
- Result.currentUsage = CurrentLocalMemUsage;
- Result.limit = MaxSizeWithWaveCount;
- Result.maxOccupancy = MaxOccupancy;
- Result.promotable = false;
+ if (Result.CurrentUsage > MaxSizeWithWaveCount) {
+ Result.Limit = MaxSizeWithWaveCount;
+ Result.MaxOccupancy = MaxOccupancy;
+ Result.Promotable = false;
return Result;
}
- Result.currentUsage = CurrentLocalMemUsage;
- Result.limit = MaxSizeWithWaveCount;
- Result.maxOccupancy = MaxOccupancy;
- Result.promotable = true;
+ Result.Limit = MaxSizeWithWaveCount;
+ Result.MaxOccupancy = MaxOccupancy;
return Result;
}
diff --git a/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h b/llvm/lib/Target/AMDGPU/AMDGPULDSUtils.h
similarity index 55%
rename from llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h
rename to llvm/lib/Target/AMDGPU/AMDGPULDSUtils.h
index 18690a77647c1..71ae853e78c3a 100644
--- a/llvm/lib/Target/AMDGPU/Utils/AMDGPULDSUtils.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPULDSUtils.h
@@ -10,38 +10,34 @@
//
//===----------------------------------------------------------------------===//
-#ifndef LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H
-#define LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H
+#ifndef LLVM_LIB_TARGET_AMDGPU_AMDGPULDSUTILS_H
+#define LLVM_LIB_TARGET_AMDGPU_AMDGPULDSUTILS_H
+#include "llvm/Support/Alignment.h"
#include <cstdint>
-#include <utility>
namespace llvm {
-class AMDGPUSubtarget;
class Function;
class IRBuilderBase;
-class Module;
class TargetMachine;
class Value;
namespace AMDGPU {
-/// Get workitem id for dimension N (0,1,2).
-Value *getWorkitemID(IRBuilderBase &Builder, Module &M,
- const AMDGPUSubtarget &ST, unsigned N);
-
/// Compute linear thread id within a workgroup.
-Value *buildLinearThreadId(IRBuilderBase &Builder, Module &M,
- const AMDGPUSubtarget &ST);
+Value *buildLinearThreadId(IRBuilderBase &Builder, const TargetMachine &TM);
struct AMDGPULDSBudget {
- uint32_t currentUsage = 0;
- uint32_t limit = 0;
- unsigned maxOccupancy = 0;
- bool promotable = false;
- bool disabledDueToLocalArg = false;
- bool disabledDueToExternDynShared = false;
+ uint64_t CurrentUsage = 0;
+ uint64_t Limit = 0;
+ unsigned MaxOccupancy = 0;
+ bool Promotable = false;
+ bool DisabledDueToLocalArg = false;
+ bool DisabledDueToExternDynShared = false;
+
+ /// Reserve an allocation while allowing for any possible leading padding.
+ bool tryReserve(uint64_t AllocSize, Align Alignment);
};
AMDGPULDSBudget computeLDSBudget(const Function &F, const TargetMachine &TM);
@@ -50,4 +46,4 @@ AMDGPULDSBudget computeLDSBudget(const Function &F, const TargetMachine &TM);
} // end namespace llvm
-#endif // LLVM_LIB_TARGET_AMDGPU_UTILS_AMDGPULDSUTILS_H
+#endif // LLVM_LIB_TARGET_AMDGPU_AMDGPULDSUTILS_H
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp b/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp
index c60541ffa9a36..4c6b48fee027e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPromoteAlloca.cpp
@@ -26,9 +26,9 @@
//===----------------------------------------------------------------------===//
#include "AMDGPU.h"
+#include "AMDGPULDSUtils.h"
#include "GCNSubtarget.h"
#include "Utils/AMDGPUBaseInfo.h"
-#include "Utils/AMDGPULDSUtils.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/Analysis/CaptureTracking.h"
#include "llvm/Analysis/InstSimplifyFolder.h"
@@ -38,8 +38,6 @@
#include "llvm/CodeGen/TargetPassConfig.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/IntrinsicInst.h"
-#include "llvm/IR/IntrinsicsAMDGPU.h"
-#include "llvm/IR/IntrinsicsR600.h"
#include "llvm/IR/PatternMatch.h"
#include "llvm/InitializePasses.h"
#include "llvm/Pass.h"
@@ -134,17 +132,13 @@ class AMDGPUPromoteAllocaImpl {
const DataLayout &DL;
// FIXME: This should be per-kernel.
- uint32_t LocalMemLimit = 0;
- uint32_t CurrentLocalMemUsage = 0;
+ uint64_t LocalMemLimit = 0;
+ uint64_t CurrentLocalMemUsage = 0;
unsigned MaxVGPRs;
unsigned VGPRBudgetRatio;
unsigned MaxVectorRegs;
bool IsAMDGCN = false;
- bool IsAMDHSA = false;
-
- std::pair<Value *, Value *> getLocalSizeYZ(IRBuilder<> &Builder);
- Value *getWorkitemID(IRBuilder<> &Builder, unsigned N);
bool collectAllocaUses(AllocaAnalysis &AA) const;
@@ -177,7 +171,6 @@ class AMDGPUPromoteAllocaImpl {
: TM(TM), LI(LI), Mod(M), DL(M.getDataLayout()) {
const Triple &TT = M.getTargetTriple();
IsAMDGCN = TT.isAMDGCN();
- IsAMDHSA = TT.getOS() == Triple::AMDHSA;
}
bool run(Function &F, bool PromoteToLDS);
@@ -368,7 +361,9 @@ bool AMDGPUPromoteAllocaImpl::run(Function &F, bool PromoteToLDS) {
return false;
bool SufficientLDS = PromoteToLDS && hasSufficientLocalMem(F);
- MaxVGPRs = IsAMDGCN ? getMaxVGPRs(CurrentLocalMemUsage, TM, F) : 128;
+ MaxVGPRs =
+ IsAMDGCN ? getMaxVGPRs(static_cast<unsigned>(CurrentLocalMemUsage), TM, F)
+ : 128;
setFunctionLimits(F);
unsigned VectorizationBudget =
@@ -1191,124 +1186,6 @@ void AMDGPUPromoteAllocaImpl::promoteAllocaToVector(AllocaAnalysis &AA) {
AA.Alloca->eraseFromParent();
}
-std::pair<Value *, Value *>
-AMDGPUPromoteAllocaImpl::getLocalSizeYZ(IRBuilder<> &Builder) {
- Function &F = *Builder.GetInsertBlock()->getParent();
- const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
-
- if (!IsAMDHSA) {
- CallInst *LocalSizeY = Builder.CreateIntrinsicWithoutFolding(
- Intrinsic::r600_read_local_size_y, {});
- CallInst *LocalSizeZ = Builder.CreateIntrinsicWithoutFolding(
- Intrinsic::r600_read_local_size_z, {});
-
- ST.makeLIDRangeMetadata(LocalSizeY);
- ST.makeLIDRangeMetadata(LocalSizeZ);
-
- return std::pair(LocalSizeY, LocalSizeZ);
- }
-
- // We must read the size out of the dispatch pointer.
- assert(IsAMDGCN);
-
- // We are indexing into this struct, and want to extract the workgroup_size_*
- // fields.
- //
- // typedef struct hsa_kernel_dispatch_packet_s {
- // uint16_t header;
- // uint16_t setup;
- // uint16_t workgroup_size_x ;
- // uint16_t workgroup_size_y;
- // uint16_t workgroup_size_z;
- // uint16_t reserved0;
- // uint32_t grid_size_x ;
- // uint32_t grid_size_y ;
- // uint32_t grid_size_z;
- //
- // uint32_t private_segment_size;
- // uint32_t group_segment_size;
- // uint64_t kernel_object;
- //
- // #ifdef HSA_LARGE_MODEL
- // void *kernarg_address;
- // #elif defined HSA_LITTLE_ENDIAN
- // void *kernarg_address;
- // uint32_t reserved1;
- // #else
- // uint32_t reserved1;
- // void *kernarg_address;
- // #endif
- // uint64_t reserved2;
- // hsa_signal_t completion_signal; // uint64_t wrapper
- // } hsa_kernel_dispatch_packet_t
- //
- CallInst *DispatchPtr =
- Builder.CreateIntrinsicWithoutFolding(Intrinsic::amdgcn_dispatch_ptr, {});
- DispatchPtr->addRetAttr(Attribute::NoAlias);
- DispatchPtr->addRetAttr(Attribute::NonNull);
- F.removeFnAttr("amdgpu-no-dispatch-ptr");
-
- // Size of the dispatch packet struct.
- DispatchPtr->addDereferenceableRetAttr(64);
-
- Type *I32Ty = Type::getInt32Ty(Mod.getContext());
-
- // We could do a single 64-bit load here, but it's likely that the basic
- // 32-bit and extract sequence is already present, and it is probably easier
- // to CSE this. The loads should be mergeable later anyway.
- Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(I32Ty, DispatchPtr, 1);
- LoadInst *LoadXY = Builder.CreateAlignedLoad(I32Ty, GEPXY, Align(4));
-
- Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(I32Ty, DispatchPtr, 2);
- LoadInst *LoadZU = Builder.CreateAlignedLoad(I32Ty, GEPZU, Align(4));
-
- MDNode *MD = MDNode::get(Mod.getContext(), {});
- LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);
- LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);
- ST.makeLIDRangeMetadata(LoadZU);
-
- // Extract y component. Upper half of LoadZU should be zero already.
- Value *Y = Builder.CreateLShr(LoadXY, 16);
-
- return std::pair(Y, LoadZU);
-}
-
-Value *AMDGPUPromoteAllocaImpl::getWorkitemID(IRBuilder<> &Builder,
- unsigned N) {
- Function *F = Builder.GetInsertBlock()->getParent();
- const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, *F);
- Intrinsic::ID IntrID = Intrinsic::not_intrinsic;
- StringRef AttrName;
-
- switch (N) {
- case 0:
- IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_x
- : (Intrinsic::ID)Intrinsic::r600_read_tidig_x;
- AttrName = "amdgpu-no-workitem-id-x";
- break;
- case 1:
- IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_y
- : (Intrinsic::ID)Intrinsic::r600_read_tidig_y;
- AttrName = "amdgpu-no-workitem-id-y";
- break;
-
- case 2:
- IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_z
- : (Intrinsic::ID)Intrinsic::r600_read_tidig_z;
- AttrName = "amdgpu-no-workitem-id-z";
- break;
- default:
- llvm_unreachable("invalid dimension");
- }
-
- Function *WorkitemIdFn = Intrinsic::getOrInsertDeclaration(&Mod, IntrID);
- CallInst *CI = Builder.CreateCall(WorkitemIdFn);
- ST.makeLIDRangeMetadata(CI);
- F->removeFnAttr(AttrName);
-
- return CI;
-}
-
static bool isCallPromotable(CallInst *CI) {
IntrinsicInst *II = dyn_cast<IntrinsicInst>(CI);
if (!II)
@@ -1462,14 +1339,14 @@ void AMDGPUPromoteAllocaImpl::analyzePromoteToLDS(AllocaAnalysis &AA) const {
bool AMDGPUPromoteAllocaImpl::hasSufficientLocalMem(const Function &F) {
AMDGPU::AMDGPULDSBudget Budget = AMDGPU::computeLDSBudget(F, TM);
- CurrentLocalMemUsage = Budget.currentUsage;
- LocalMemLimit = Budget.limit;
- if (!Budget.promotable) {
- if (Budget.disabledDueToLocalArg) {
+ CurrentLocalMemUsage = Budget.CurrentUsage;
+ LocalMemLimit = Budget.Limit;
+ if (!Budget.Promotable) {
+ if (Budget.DisabledDueToLocalArg) {
LLVM_DEBUG(dbgs() << "Function has local memory argument. Promoting to "
"local memory disabled.\n");
}
- if (Budget.disabledDueToExternDynShared) {
+ if (Budget.DisabledDueToExternDynShared) {
LLVM_DEBUG(dbgs() << "Function has a reference to externally allocated "
"local memory. Promoting to local memory "
"disabled.\n");
@@ -1478,7 +1355,12 @@ bool AMDGPUPromoteAllocaImpl::hasSufficientLocalMem(const Function &F) {
}
LLVM_DEBUG(dbgs() << F.getName() << " uses " << CurrentLocalMemUsage
- << " bytes of LDS\n");
+ << " bytes of LDS\n"
+ << " Rounding size to " << LocalMemLimit
+ << " with a maximum occupancy of " << Budget.MaxOccupancy
+ << '\n'
+ << " and " << (LocalMemLimit - CurrentLocalMemUsage)
+ << " available for promotion\n");
return true;
}
@@ -1506,18 +1388,20 @@ bool AMDGPUPromoteAllocaImpl::tryPromoteAllocaToLDS(
// FIXME: It is also possible that if we're allowed to use all of the memory
// could end up using more than the maximum due to alignment padding.
- uint32_t NewSize = alignTo(CurrentLocalMemUsage, Alignment);
std::optional<TypeSize> ElemSize = AA.Alloca->getAllocationSize(DL);
if (!ElemSize || ElemSize->isScalable())
return false;
- TypeSize AllocSize = WorkGroupSize * *ElemSize;
- NewSize += AllocSize.getFixedValue();
-
- if (NewSize > LocalMemLimit) {
- LLVM_DEBUG(dbgs() << " " << AllocSize
- << " bytes of local memory not available to promote\n");
+ uint64_t ElemBytes = ElemSize->getFixedValue();
+ if (Alignment.value() > LocalMemLimit)
+ return false;
+ uint64_t NewSize = alignTo(CurrentLocalMemUsage, Alignment);
+ if (NewSize > LocalMemLimit || WorkGroupSize == 0 ||
+ ElemBytes > (LocalMemLimit - NewSize) / WorkGroupSize) {
+ LLVM_DEBUG(dbgs() << " local memory not available to promote\n");
return false;
}
+ uint64_t AllocSize = ElemBytes * WorkGroupSize;
+ NewSize += AllocSize;
CurrentLocalMemUsage = NewSize;
@@ -1533,18 +1417,7 @@ bool AMDGPUPromoteAllocaImpl::tryPromoteAllocaToLDS(
GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
GV->setAlignment(AA.Alloca->getAlign());
- Value *TCntY, *TCntZ;
-
- std::tie(TCntY, TCntZ) = getLocalSizeYZ(Builder);
- Value *TIdX = getWorkitemID(Builder, 0);
- Value *TIdY = getWorkitemID(Builder, 1);
- Value *TIdZ = getWorkitemID(Builder, 2);
-
- Value *Tmp0 = Builder.CreateMul(TCntY, TCntZ, "", true, true);
- Tmp0 = Builder.CreateMul(Tmp0, TIdX);
- Value *Tmp1 = Builder.CreateMul(TIdY, TCntZ, "", true, true);
- Value *TID = Builder.CreateAdd(Tmp0, Tmp1);
- TID = Builder.CreateAdd(TID, TIdZ);
+ Value *TID = AMDGPU::buildLinearThreadId(Builder, TM);
LLVMContext &Context = Mod.getContext();
Value *Indices[] = {Constant::getNullValue(Type::getInt32Ty(Context)), TID};
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index e3da368d6f7ea..3705c57e5b6ff 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -1010,14 +1010,14 @@ static Expected<unsigned> parseAMDGPULDSBufferingMaxBytes(StringRef Params) {
continue;
if (!Param.consume_front("max-bytes=")) {
return make_error<StringError>(
- formatv("invalid AMDGPULDSBuffering pass parameter '{0}' ", Param)
+ formatv("invalid AMDGPU LDS buffering pass parameter '{0}'", Param)
.str(),
inconvertibleErrorCode());
}
unsigned Parsed = 0;
if (Param.getAsInteger(10, Parsed)) {
return make_error<StringError>(
- formatv("invalid AMDGPULDSBuffering max-bytes '{0}' ", Param).str(),
+ formatv("invalid AMDGPU LDS buffering max-bytes '{0}'", Param).str(),
inconvertibleErrorCode());
}
Result = Parsed;
diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt
index ea89a7cc616c0..34d7101085255 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -80,6 +80,7 @@ add_llvm_target(AMDGPUCodeGen
AMDGPULowerKernelAttributes.cpp
AMDGPULowerModuleLDSPass.cpp
AMDGPULDSBuffering.cpp
+ AMDGPULDSUtils.cpp
AMDGPUPrepareAGPRAlloc.cpp
AMDGPULowerExecSync.cpp
AMDGPUSwLowerLDS.cpp
diff --git a/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt b/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
index 4633763d23f22..7b2200d8bc488 100644
--- a/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/Utils/CMakeLists.txt
@@ -4,7 +4,6 @@ add_llvm_component_library(LLVMAMDGPUUtils
AMDGPUDelayedMCExpr.cpp
AMDGPUPALMetadata.cpp
AMDKernelCodeTUtils.cpp
- AMDGPULDSUtils.cpp
LINK_COMPONENTS
Analysis
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-budget.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-budget.ll
new file mode 100644
index 0000000000000..44b530d84e453
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-budget.ll
@@ -0,0 +1,50 @@
+; NOTE: Do not autogenerate. The checks cover LDS allocation rejection only.
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64>' -S < %s | FileCheck %s --check-prefix=PADDING
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=4194304>' -S < %s | FileCheck %s --check-prefix=OVERFLOW
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64>' -S < %s | FileCheck %s --check-prefix=EXISTING-OVERFLOW
+
+target triple = "amdgcn-amd-amdhsa"
+
+ at padding.used = internal addrspace(3) global [102400 x i8] poison, align 16
+ at existing_overflow.a = internal addrspace(3) global [2147483700 x i8] poison, align 1
+ at existing_overflow.b = internal addrspace(3) global [2147483700 x i8] poison, align 1
+
+; The store size of <9 x i32> is 36 bytes, but its allocation size is 64
+; bytes. Account for the allocation size of every work-item slot.
+; PADDING-NOT: @padding.ldsbuf
+; PADDING-LABEL: define amdgpu_kernel void @padding(
+; PADDING: %value = load <9 x i32>, ptr addrspace(1) %ptr, align 16
+; PADDING: store <9 x i32> %value, ptr addrspace(1) %ptr, align 16
+define amdgpu_kernel void @padding(ptr addrspace(1) %ptr) #0 {
+ %used = load volatile i8, ptr addrspace(3) @padding.used, align 1
+ %value = load <9 x i32>, ptr addrspace(1) %ptr, align 16
+ store <9 x i32> %value, ptr addrspace(1) %ptr, align 16
+ ret void
+}
+
+; The per-work-group allocation does not fit in the target's LDS limit. The
+; size calculation must not wrap when max-bytes permits this candidate.
+; OVERFLOW-NOT: @overflow.ldsbuf
+; OVERFLOW-LABEL: define amdgpu_kernel void @overflow(
+; OVERFLOW: %value = load <1048576 x i32>, ptr addrspace(1) %ptr, align 16
+; OVERFLOW: store <1048576 x i32> %value, ptr addrspace(1) %ptr, align 16
+define amdgpu_kernel void @overflow(ptr addrspace(1) %ptr) #0 {
+ %value = load <1048576 x i32>, ptr addrspace(1) %ptr, align 16
+ store <1048576 x i32> %value, ptr addrspace(1) %ptr, align 16
+ ret void
+}
+
+; Accumulating existing LDS allocations must not wrap at 32 bits.
+; EXISTING-OVERFLOW-NOT: @existing_overflow.ldsbuf
+; EXISTING-OVERFLOW-LABEL: define amdgpu_kernel void @existing_overflow(
+; EXISTING-OVERFLOW: %value = load <4 x i32>, ptr addrspace(1) %ptr, align 16
+; EXISTING-OVERFLOW: store <4 x i32> %value, ptr addrspace(1) %ptr, align 16
+define amdgpu_kernel void @existing_overflow(ptr addrspace(1) %ptr) #0 {
+ %used.a = load volatile i8, ptr addrspace(3) @existing_overflow.a, align 1
+ %used.b = load volatile i8, ptr addrspace(3) @existing_overflow.b, align 1
+ %value = load <4 x i32>, ptr addrspace(1) %ptr, align 16
+ store <4 x i32> %value, ptr addrspace(1) %ptr, align 16
+ ret void
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1024,1024" }
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll
new file mode 100644
index 0000000000000..b3b559025ed88
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll
@@ -0,0 +1,9 @@
+; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<unknown=1>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=UNKNOWN
+; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<max-bytes=invalid>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=INVALID
+
+; UNKNOWN: amdgpu-lds-buffering: invalid AMDGPU LDS buffering pass parameter 'unknown=1'
+; INVALID: amdgpu-lds-buffering: invalid AMDGPU LDS buffering max-bytes 'invalid'
+
+define amdgpu_kernel void @kernel() {
+ ret void
+}
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-layout.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-layout.ll
new file mode 100644
index 0000000000000..6fad62b0537df
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-layout.ll
@@ -0,0 +1,22 @@
+; NOTE: Do not autogenerate. This checks the final LDS layout, not instructions.
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-enable-lds-buffering < %s | FileCheck %s
+
+target triple = "amdgcn-amd-amdhsa"
+
+ at used = internal addrspace(3) global [130500 x i8] poison, align 4
+
+; The backend lays out LDS globals created after module LDS lowering in inverse
+; use order. Include worst-case leading padding when reserving each slot so
+; this does not exceed the gfx950 LDS limit during code generation.
+; CHECK: .amdhsa_group_segment_fixed_size 131536
+define amdgpu_kernel void @layout_order(ptr addrspace(1) %high,
+ ptr addrspace(1) %low) #0 {
+ %used = load volatile i8, ptr addrspace(3) @used, align 1
+ %high.value = load i8, ptr addrspace(1) %high, align 131072
+ store i8 %high.value, ptr addrspace(1) %high, align 131072
+ %low.value = load i8, ptr addrspace(1) %low, align 16
+ store i8 %low.value, ptr addrspace(1) %low, align 16
+ ret void
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1024,1024" }
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-targets.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-targets.ll
new file mode 100644
index 0000000000000..0e54050867959
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-targets.ll
@@ -0,0 +1,23 @@
+; NOTE: Do not autogenerate. The checks cover target-specific work-item inputs.
+; RUN: opt -mtriple=amdgcn-amd-amdpal -mcpu=gfx900 -passes='amdgpu-lds-buffering<max-bytes=64>' -S < %s | FileCheck %s --check-prefix=PAL
+; RUN: opt -mtriple=r600-amd-unknown -mcpu=redwood -passes='amdgpu-lds-buffering<max-bytes=64>' -S < %s | FileCheck %s --check-prefix=R600
+
+; PAL-NOT: @llvm.amdgcn.dispatch.ptr
+; PAL-LABEL: define amdgpu_kernel void @kernel(
+; PAL-NOT: @llvm.amdgcn.dispatch.ptr
+; PAL: call{{.*}}i32 @llvm.r600.read.local.size.y()
+; PAL: call{{.*}}i32 @llvm.r600.read.local.size.z()
+; PAL: call{{.*}}i32 @llvm.amdgcn.workitem.id.x()
+; R600-NOT: @kernel.ldsbuf
+; R600-LABEL: define amdgpu_kernel void @kernel(
+; R600: %value = load <4 x i32>, ptr addrspace(1) %ptr, align 16
+; R600: store <4 x i32> %value, ptr addrspace(1) %ptr, align 16
+define amdgpu_kernel void @kernel(ptr addrspace(1) %ptr,
+ ptr addrspace(1) %other) #0 {
+ %value = load <4 x i32>, ptr addrspace(1) %ptr, align 16
+ store i32 0, ptr addrspace(1) %other, align 4
+ store <4 x i32> %value, ptr addrspace(1) %ptr, align 16
+ ret void
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" }
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering.ll
index 64165c3cb1e40..5b3fba2253cac 100644
--- a/llvm/test/CodeGen/AMDGPU/lds-buffering.ll
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering.ll
@@ -1,4 +1,5 @@
-; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64>' -S %s | FileCheck %s
+; NOTE: Do not autogenerate. The checks focus on which patterns are transformed.
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64>' -S < %s | FileCheck %s
%state = type { <4 x i32>, <4 x i32> }
>From 94c2cfbee8c7e73559a337fe57c0a287bd7fb42d Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 1 Sep 2026 17:46:00 -0400
Subject: [PATCH 3/3] AMDGPU: add experimental LDS-buffering and pre-loop
waitcnt diagnostics
These options support an investigation into whether the AMDGPU LDS
buffering transform's performance gain generalizes. They are all hidden
and default-off; none changes default codegen.
- LDS buffering A/B controls in `AMDGPULDSBuffering`: a `shadow-lds` and
an `irrelevant-lds` mode and an `-amdgpu-lds-buffering-min-align`
option, so a standalone reproducer can separate live-range shortening
from generic LDS traffic and from the pre-loop machine-code shape.
- Two pre-loop waitcnt experiments in `SIInsertWaitcnts`:
`-amdgpu-enable-pre-loop-vmem-wait` retires older VMEM loads not used
by a loop at its preheader, and `-amdgpu-preheader-flush-loop-carried-load`
flushes an outside-loop VMEM load whose value is used in the loop even
when the loop has its own VMEM traffic. Both are performance-only
(flushing is always correct).
Findings: the PRNG-shaped gain needs the register reassignment, an actual
LDS store, and the partial wait together; no single factor or pair
reproduces it. The stream-shaped preheader-wait win does not reproduce on
the current toolchain and the flush heuristic regresses. These stay
experimental diagnostics and are not proposed for upstream.
Adds `waitcnt-pre-loop-vmem.mir` and `lds-buffering-options.ll`.
---
llvm/lib/Target/AMDGPU/AMDGPU.h | 20 ++-
llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp | 63 +++++--
llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def | 7 +-
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 102 +++++++++--
llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp | 106 +++++++++++-
.../AMDGPU/lds-buffering-invalid-params.ll | 6 +
.../CodeGen/AMDGPU/lds-buffering-options.ll | 64 +++++++
.../CodeGen/AMDGPU/waitcnt-pre-loop-vmem.mir | 160 ++++++++++++++++++
8 files changed, 489 insertions(+), 39 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/lds-buffering-options.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/waitcnt-pre-loop-vmem.mir
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index 4fa35f4a92d46..1195fadac2feb 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -295,18 +295,30 @@ struct AMDGPUPromoteAllocaToVectorPass
// Buffer selected per-thread global memory through LDS to improve
// performance in memory-bound kernels. This runs late and is separate
// from alloca promotion.
+enum class AMDGPULDSBufferingMode { Buffer, ShadowLDS, IrrelevantLDS };
+
+struct AMDGPULDSBufferingOptions {
+ unsigned MaxBytes = 64;
+ unsigned MinAlignment = 16;
+ int OnlyCandidate = -1;
+ AMDGPULDSBufferingMode Mode = AMDGPULDSBufferingMode::Buffer;
+};
+
struct AMDGPULDSBufferingPass : OptionalPassInfoMixin<AMDGPULDSBufferingPass> {
- AMDGPULDSBufferingPass(const AMDGPUTargetMachine &TM, unsigned MaxBytes = 64)
- : TM(TM), MaxBytes(MaxBytes) {}
+ AMDGPULDSBufferingPass(
+ const AMDGPUTargetMachine &TM,
+ AMDGPULDSBufferingOptions Options = AMDGPULDSBufferingOptions())
+ : TM(TM), Options(Options) {}
PreservedAnalyses run(Function &F, FunctionAnalysisManager &AM);
private:
const AMDGPUTargetMachine &TM;
- unsigned MaxBytes;
+ AMDGPULDSBufferingOptions Options;
};
// Legacy PM wrapper for LDS buffering
-FunctionPass *createAMDGPULDSBufferingLegacyPass();
+FunctionPass *createAMDGPULDSBufferingLegacyPass(
+ AMDGPULDSBufferingOptions Options = AMDGPULDSBufferingOptions());
void initializeAMDGPULDSBufferingLegacyPass(PassRegistry &);
struct AMDGPUAtomicOptimizerPass
diff --git a/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp b/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
index 11de8d98a8240..d885ef12fb6f6 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPULDSBuffering.cpp
@@ -60,13 +60,14 @@ namespace {
class AMDGPULDSBufferingImpl {
const AMDGPUTargetMachine &TM;
- unsigned MaxBytes;
+ AMDGPULDSBufferingOptions Options;
Module *Mod = nullptr;
const DataLayout *DL = nullptr;
public:
- AMDGPULDSBufferingImpl(const AMDGPUTargetMachine &TM, unsigned MaxBytes)
- : TM(TM), MaxBytes(MaxBytes) {}
+ AMDGPULDSBufferingImpl(const AMDGPUTargetMachine &TM,
+ AMDGPULDSBufferingOptions Options)
+ : TM(TM), Options(Options) {}
bool run(Function &F) {
LLVM_DEBUG(dbgs() << "[LDSBuffer] Visit function: " << F.getName() << '\n');
@@ -84,6 +85,7 @@ class AMDGPULDSBufferingImpl {
unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(F).second;
bool Changed = false;
+ unsigned CandidateIndex = 0;
unsigned NumTransformed = 0;
// Minimal pattern: a load from AS(1) whose only use is a store back to the
@@ -122,21 +124,27 @@ class AMDGPULDSBufferingImpl {
if (StoreSize.isScalable() || AllocSize.isScalable())
continue;
uint64_t CopySize = StoreSize.getFixedValue();
- if (CopySize == 0 || CopySize > MaxBytes)
+ if (CopySize == 0 || CopySize > Options.MaxBytes)
continue;
uint64_t SlotSize = AllocSize.getFixedValue();
- Align MinAlign = Align(16);
+ Align MinAlign = Align(Options.MinAlignment);
Align LoadAlign = LI->getAlign();
Align StoreAlign = SI->getAlign();
Align Alignment = std::min(LoadAlign, StoreAlign);
if (Alignment < MinAlign)
continue;
+ unsigned ThisCandidate = CandidateIndex++;
+ if (Options.OnlyCandidate >= 0 &&
+ ThisCandidate != static_cast<unsigned>(Options.OnlyCandidate))
+ continue;
+
// Create LDS slot near the load and emit memcpy global->LDS.
LLVM_DEBUG({
dbgs() << "[LDSBuffer] Candidate found: load->store same ptr in "
<< F.getName() << '\n';
- dbgs() << " size=" << CopySize
+ dbgs() << " index=" << ThisCandidate
+ << ", size=" << CopySize
<< "B, loadAlign=" << LoadAlign.value()
<< ", storeAlign=" << StoreAlign.value()
<< ", chosenAlign=" << Alignment.value()
@@ -151,6 +159,34 @@ class AMDGPULDSBufferingImpl {
continue;
auto [GV, SlotPtr] = createLDSGlobalAndThreadSlot(
F, ValTy, WorkGroupSize, Alignment, BLoad);
+
+ if (Options.Mode != AMDGPULDSBufferingMode::Buffer) {
+ Value *ControlValue =
+ Options.Mode == AMDGPULDSBufferingMode::ShadowLDS
+ ? static_cast<Value *>(LI)
+ : Constant::getNullValue(ValTy);
+ IRBuilder<> BAfterLoad(LI->getNextNode());
+ StoreInst *ControlStore = BAfterLoad.CreateStore(
+ ControlValue, SlotPtr, /*isVolatile=*/true);
+ ControlStore->setAlignment(Alignment);
+
+ IRBuilder<> BAfterStore(SI->getNextNode());
+ LoadInst *ControlLoad = BAfterStore.CreateLoad(
+ ValTy, SlotPtr, /*isVolatile=*/true, "ldsbuf.control");
+ ControlLoad->setAlignment(Alignment);
+
+ LLVM_DEBUG(dbgs()
+ << "[LDSBuffer] Insert "
+ << (Options.Mode == AMDGPULDSBufferingMode::ShadowLDS
+ ? "shadow"
+ : "irrelevant")
+ << " LDS control: " << GV->getName() << ", bytes="
+ << CopySize << ", align=" << Alignment.value() << '\n');
+ Changed = true;
+ ++NumTransformed;
+ continue;
+ }
+
// memcpy p3 <- p1
LLVM_DEBUG(dbgs() << "[LDSBuffer] Insert memcpy global->LDS: "
<< GV->getName() << ", bytes=" << CopySize
@@ -215,7 +251,7 @@ class AMDGPULDSBufferingImpl {
PreservedAnalyses AMDGPULDSBufferingPass::run(Function &F,
FunctionAnalysisManager &AM) {
- bool Changed = AMDGPULDSBufferingImpl(TM, MaxBytes).run(F);
+ bool Changed = AMDGPULDSBufferingImpl(TM, Options).run(F);
if (!Changed)
return PreservedAnalyses::all();
@@ -231,9 +267,12 @@ PreservedAnalyses AMDGPULDSBufferingPass::run(Function &F,
namespace {
class AMDGPULDSBufferingLegacy : public FunctionPass {
+ AMDGPULDSBufferingOptions Options;
+
public:
static char ID;
- AMDGPULDSBufferingLegacy() : FunctionPass(ID) {}
+ AMDGPULDSBufferingLegacy(AMDGPULDSBufferingOptions Options)
+ : FunctionPass(ID), Options(Options) {}
StringRef getPassName() const override { return "AMDGPU LDS Buffering"; }
@@ -246,8 +285,7 @@ class AMDGPULDSBufferingLegacy : public FunctionPass {
if (skipFunction(F))
return false;
if (TargetPassConfig *TPC = getAnalysisIfAvailable<TargetPassConfig>())
- return AMDGPULDSBufferingImpl(TPC->getTM<AMDGPUTargetMachine>(),
- /*MaxBytes=*/64)
+ return AMDGPULDSBufferingImpl(TPC->getTM<AMDGPUTargetMachine>(), Options)
.run(F);
return false;
}
@@ -262,6 +300,7 @@ INITIALIZE_PASS_BEGIN(AMDGPULDSBufferingLegacy, DEBUG_TYPE,
INITIALIZE_PASS_END(AMDGPULDSBufferingLegacy, DEBUG_TYPE,
"AMDGPU per-thread LDS buffering", false, false)
-FunctionPass *llvm::createAMDGPULDSBufferingLegacyPass() {
- return new AMDGPULDSBufferingLegacy();
+FunctionPass *
+llvm::createAMDGPULDSBufferingLegacyPass(AMDGPULDSBufferingOptions Options) {
+ return new AMDGPULDSBufferingLegacy(Options);
}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
index e174d6b944c14..315f3ae9740c6 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
+++ b/llvm/lib/Target/AMDGPU/AMDGPUPassRegistry.def
@@ -100,8 +100,11 @@ FUNCTION_ALIAS_ANALYSIS("amdgpu-aa", AMDGPUAA())
#endif
FUNCTION_PASS_WITH_PARAMS(
"amdgpu-lds-buffering", "AMDGPULDSBufferingPass",
- [=](unsigned MaxBytes) { return AMDGPULDSBufferingPass(*this, MaxBytes); },
- parseAMDGPULDSBufferingMaxBytes, "max-bytes")
+ [=](AMDGPULDSBufferingOptions Options) {
+ return AMDGPULDSBufferingPass(*this, Options);
+ },
+ parseAMDGPULDSBufferingOptions,
+ "max-bytes;min-align;only-candidate;mode")
FUNCTION_PASS_WITH_PARAMS(
"amdgpu-atomic-optimizer",
"AMDGPUAtomicOptimizerPass",
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 3705c57e5b6ff..9964ba9464752 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -114,6 +114,7 @@
#include "llvm/Passes/PassBuilder.h"
#include "llvm/Support/Compiler.h"
#include "llvm/Support/FormatVariadic.h"
+#include "llvm/Support/MathExtras.h"
#include "llvm/TargetParser/AMDGPUTargetParser.h"
#include "llvm/Transforms/HipStdPar/HipStdPar.h"
#include "llvm/Transforms/IPO.h"
@@ -606,6 +607,29 @@ static cl::opt<bool> EnableLDSBuffering(
cl::desc("Enable AMDGPU LDS Buffering pass in the default pipeline"),
cl::init(false), cl::Hidden);
+static cl::opt<unsigned> LDSBufferingMinAlignment(
+ "amdgpu-lds-buffering-min-align",
+ cl::desc("Minimum load/store alignment for AMDGPU LDS buffering"),
+ cl::init(16), cl::Hidden);
+
+static cl::opt<int> LDSBufferingOnlyCandidate(
+ "amdgpu-lds-buffering-only-candidate",
+ cl::desc("Transform only the selected zero-based LDS buffering candidate "
+ "(-1 transforms all candidates)"),
+ cl::init(-1), cl::Hidden);
+
+static cl::opt<AMDGPULDSBufferingMode> LDSBufferingMode(
+ "amdgpu-lds-buffering-mode",
+ cl::desc("Select AMDGPU LDS buffering experiment mode"),
+ cl::values(
+ clEnumValN(AMDGPULDSBufferingMode::Buffer, "buffer",
+ "Buffer the candidate value through LDS"),
+ clEnumValN(AMDGPULDSBufferingMode::ShadowLDS, "shadow-lds",
+ "Copy the candidate through LDS without replacing it"),
+ clEnumValN(AMDGPULDSBufferingMode::IrrelevantLDS, "irrelevant-lds",
+ "Keep the candidate live and add unrelated LDS traffic")),
+ cl::init(AMDGPULDSBufferingMode::Buffer), cl::Hidden);
+
static cl::opt<bool>
EnableLoopPrefetch("amdgpu-loop-prefetch",
cl::desc("Enable loop data prefetch on AMDGPU"),
@@ -1001,30 +1025,81 @@ parseAMDGPUAtomicOptimizerStrategy(StringRef Params) {
return make_error<StringError>("invalid parameter", inconvertibleErrorCode());
}
-static Expected<unsigned> parseAMDGPULDSBufferingMaxBytes(StringRef Params) {
- unsigned Result = 64;
+static Expected<AMDGPULDSBufferingOptions>
+parseAMDGPULDSBufferingOptions(StringRef Params) {
+ AMDGPULDSBufferingOptions Result;
while (!Params.empty()) {
StringRef Param;
std::tie(Param, Params) = Params.split(';');
if (Param.empty())
continue;
- if (!Param.consume_front("max-bytes=")) {
+ if (Param.consume_front("max-bytes=")) {
+ unsigned Parsed = 0;
+ if (Param.getAsInteger(10, Parsed)) {
+ return make_error<StringError>(
+ formatv("invalid AMDGPU LDS buffering max-bytes '{0}'", Param)
+ .str(),
+ inconvertibleErrorCode());
+ }
+ Result.MaxBytes = Parsed;
+ continue;
+ }
+ if (Param.consume_front("min-align=")) {
+ unsigned Parsed = 0;
+ if (Param.getAsInteger(10, Parsed) || !isPowerOf2_64(Parsed)) {
+ return make_error<StringError>(
+ formatv("invalid AMDGPU LDS buffering min-align '{0}'", Param)
+ .str(),
+ inconvertibleErrorCode());
+ }
+ Result.MinAlignment = Parsed;
+ continue;
+ }
+ if (Param.consume_front("only-candidate=")) {
+ int Parsed = 0;
+ if (Param.getAsInteger(10, Parsed) || Parsed < -1) {
+ return make_error<StringError>(
+ formatv("invalid AMDGPU LDS buffering only-candidate '{0}'", Param)
+ .str(),
+ inconvertibleErrorCode());
+ }
+ Result.OnlyCandidate = Parsed;
+ continue;
+ }
+ if (Param.consume_front("mode=")) {
+ if (Param == "buffer")
+ Result.Mode = AMDGPULDSBufferingMode::Buffer;
+ else if (Param == "shadow-lds")
+ Result.Mode = AMDGPULDSBufferingMode::ShadowLDS;
+ else if (Param == "irrelevant-lds")
+ Result.Mode = AMDGPULDSBufferingMode::IrrelevantLDS;
+ else
+ return make_error<StringError>(
+ formatv("invalid AMDGPU LDS buffering mode '{0}'", Param).str(),
+ inconvertibleErrorCode());
+ continue;
+ }
+ {
return make_error<StringError>(
formatv("invalid AMDGPU LDS buffering pass parameter '{0}'", Param)
.str(),
inconvertibleErrorCode());
}
- unsigned Parsed = 0;
- if (Param.getAsInteger(10, Parsed)) {
- return make_error<StringError>(
- formatv("invalid AMDGPU LDS buffering max-bytes '{0}'", Param).str(),
- inconvertibleErrorCode());
- }
- Result = Parsed;
}
return Result;
}
+static AMDGPULDSBufferingOptions getPipelineLDSBufferingOptions() {
+ if (!isPowerOf2_64(LDSBufferingMinAlignment))
+ report_fatal_error("AMDGPU LDS buffering minimum alignment must be a "
+ "nonzero power of two");
+ if (LDSBufferingOnlyCandidate < -1)
+ report_fatal_error(
+ "AMDGPU LDS buffering candidate index must be -1 or nonnegative");
+ return {/*MaxBytes=*/64, LDSBufferingMinAlignment, LDSBufferingOnlyCandidate,
+ LDSBufferingMode};
+}
+
Expected<AMDGPUAttributorOptions>
parseAMDGPUAttributorPassOptions(StringRef Params) {
AMDGPUAttributorOptions Result;
@@ -1651,7 +1726,8 @@ void AMDGPUPassConfig::addIRPasses() {
addPass(createAMDGPUPromoteAlloca());
// Run per-thread LDS buffering after promote-alloca to use leftover LDS.
if (TM.getTargetTriple().isAMDGCN() && EnableLDSBuffering)
- addPass(createAMDGPULDSBufferingLegacyPass());
+ addPass(
+ createAMDGPULDSBufferingLegacyPass(getPipelineLDSBufferingOptions()));
if (isPassEnabled(EnableScalarIRPasses))
addStraightLineScalarOptimizationPasses();
@@ -2439,7 +2515,9 @@ void AMDGPUCodeGenPassBuilder::addIRPasses(PassManagerWrapper &PMW) {
addFunctionPass(AMDGPUPromoteAllocaPass(TM), PMW);
// Run per-thread LDS buffering after promote-alloca to use leftover LDS.
if (TM.getTargetTriple().isAMDGCN() && EnableLDSBuffering)
- addFunctionPass(AMDGPULDSBufferingPass(getTM()), PMW);
+ addFunctionPass(
+ AMDGPULDSBufferingPass(getTM(), getPipelineLDSBufferingOptions()),
+ PMW);
if (isPassEnabled(EnableScalarIRPasses))
addStraightLineScalarOptimizationPasses(PMW);
diff --git a/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp b/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
index deace610d0f6b..8750590e52d66 100644
--- a/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInsertWaitcnts.cpp
@@ -59,6 +59,17 @@ static cl::opt<bool> ForceEmitZeroLoadFlag(
cl::desc("Force all waitcnt load counters to wait until 0"),
cl::init(false), cl::Hidden);
+static cl::opt<bool> EnablePreLoopVMEMWait(
+ "amdgpu-enable-pre-loop-vmem-wait",
+ cl::desc("Experimentally retire older VMEM loads before entering a loop"),
+ cl::init(false), cl::Hidden);
+
+static cl::opt<bool> EnablePreheaderFlushLoopCarriedLoad(
+ "amdgpu-preheader-flush-loop-carried-load",
+ cl::desc("Experimentally flush an outside-loop VMEM load in the preheader "
+ "even when the loop has its own VMEM traffic"),
+ cl::init(false), cl::Hidden);
+
static cl::opt<bool> ExpertSchedulingModeFlag(
"amdgpu-expert-scheduling-mode",
cl::desc("Enable expert scheduling mode 2 for all functions (GFX12+ only)"),
@@ -394,6 +405,8 @@ class SIInsertWaitcnts {
const WaitcntBrackets &Brackets);
PreheaderFlushFlags isPreheaderToFlush(MachineBasicBlock &MBB,
const WaitcntBrackets &ScoreBrackets);
+ unsigned getPreLoopVmemWait(MachineBasicBlock &MBB,
+ const WaitcntBrackets &ScoreBrackets) const;
bool isVMEMOrFlatVMEM(const MachineInstr &MI) const;
bool isDSRead(const MachineInstr &MI) const;
bool mayStoreIncrementingDSCNT(const MachineInstr &MI) const;
@@ -432,7 +445,8 @@ class SIInsertWaitcnts {
bool generateWaitcntInstBefore(MachineInstr &MI,
WaitcntBrackets &ScoreBrackets,
MachineInstr *OldWaitcntInstr,
- PreheaderFlushFlags FlushFlags);
+ PreheaderFlushFlags FlushFlags,
+ unsigned PreLoopVmemWait);
bool generateWaitcnt(AMDGPU::Waitcnt Wait,
MachineBasicBlock::instr_iterator It,
MachineBasicBlock &Block, WaitcntBrackets &ScoreBrackets,
@@ -2292,9 +2306,11 @@ bool WaitcntGeneratorGFX12Plus::createNewWaitcnt(
/// If FlushFlags.FlushVmCnt is true, we want to flush the vmcnt counter here.
/// If FlushFlags.FlushDsCnt is true, we want to flush the dscnt counter here
/// (GFX12+ only, where DS_CNT is a separate counter).
-bool SIInsertWaitcnts::generateWaitcntInstBefore(
- MachineInstr &MI, WaitcntBrackets &ScoreBrackets,
- MachineInstr *OldWaitcntInstr, PreheaderFlushFlags FlushFlags) {
+bool SIInsertWaitcnts::generateWaitcntInstBefore(MachineInstr &MI,
+ WaitcntBrackets &ScoreBrackets,
+ MachineInstr *OldWaitcntInstr,
+ PreheaderFlushFlags FlushFlags,
+ unsigned PreLoopVmemWait) {
LLVM_DEBUG(dbgs() << "\n*** GenerateWaitcntInstBefore: "; MI.print(dbgs()););
assert(!isNonWaitcntMetaInst(MI));
@@ -2593,6 +2609,9 @@ bool SIInsertWaitcnts::generateWaitcntInstBefore(
Wait.set(T, 0);
}
+ if (PreLoopVmemWait != ~0u)
+ Wait.add(AMDGPU::LOAD_CNT, PreLoopVmemWait);
+
if (FlushFlags.FlushDsCnt && ScoreBrackets.hasPendingEvent(AMDGPU::DS_CNT))
Wait.set(AMDGPU::DS_CNT, 0);
@@ -3061,6 +3080,7 @@ bool SIInsertWaitcnts::insertWaitcntInBlock(MachineFunction &MF,
// Walk over the instructions.
MachineInstr *OldWaitcntInstr = nullptr;
+ unsigned PreLoopVmemWait = getPreLoopVmemWait(Block, ScoreBrackets);
// NOTE: We may append instrs after Inst while iterating.
ScoreBrackets.verify();
@@ -3085,8 +3105,9 @@ bool SIInsertWaitcnts::insertWaitcntInBlock(MachineFunction &MF,
// Generate an s_waitcnt instruction to be placed before Inst, if needed.
Modified |= generateWaitcntInstBefore(Inst, ScoreBrackets, OldWaitcntInstr,
- FlushFlags);
+ FlushFlags, PreLoopVmemWait);
OldWaitcntInstr = nullptr;
+ PreLoopVmemWait = ~0u;
if (Inst.getOpcode() == AMDGPU::ASYNCMARK) {
// Asyncmarks record the current wait state and so should not allow
@@ -3142,6 +3163,9 @@ bool SIInsertWaitcnts::insertWaitcntInBlock(MachineFunction &MF,
Wait.set(AMDGPU::DS_CNT, 0);
}
+ if (PreLoopVmemWait != ~0u)
+ Wait.add(AMDGPU::LOAD_CNT, PreLoopVmemWait);
+
// Combine or remove any redundant waitcnts at the end of the block.
Modified |= generateWaitcnt(Wait, Block.instr_end(), Block, ScoreBrackets,
OldWaitcntInstr);
@@ -3215,6 +3239,55 @@ SIInsertWaitcnts::isPreheaderToFlush(MachineBasicBlock &MBB,
return PreheaderFlushFlags();
}
+unsigned SIInsertWaitcnts::getPreLoopVmemWait(
+ MachineBasicBlock &MBB, const WaitcntBrackets &ScoreBrackets) const {
+ if (!EnablePreLoopVMEMWait || ScoreBrackets.empty(AMDGPU::LOAD_CNT) ||
+ ScoreBrackets.counterOutOfOrder(AMDGPU::LOAD_CNT))
+ return ~0u;
+
+ unsigned Outstanding = ScoreBrackets.getOutstanding(AMDGPU::LOAD_CNT);
+ if (Outstanding < 2)
+ return ~0u;
+
+ MachineBasicBlock *Succ = MBB.getSingleSuccessor();
+ if (!Succ)
+ return ~0u;
+
+ MachineLoop *Loop = MLI.getLoopFor(Succ);
+ if (!Loop || Loop->getLoopPreheader() != &MBB)
+ return ~0u;
+
+ // Leave every pending load used by the loop outstanding. Retiring one more
+ // request drains only older loads that do not contribute to the loop.
+ unsigned MaxRequiredWait = 0;
+ bool UsesPendingLoad = false;
+ for (MachineBasicBlock *LoopMBB : Loop->blocks()) {
+ for (const MachineInstr &MI : *LoopMBB) {
+ for (const MachineOperand &Op : MI.all_uses()) {
+ if (Op.isDebug() || !TRI.isVectorRegister(MRI, Op.getReg()))
+ continue;
+ if (Op.isImplicit() && MI.mayLoadOrStore())
+ continue;
+
+ AMDGPU::Waitcnt Wait;
+ ScoreBrackets.determineWaitForPhysReg(AMDGPU::LOAD_CNT,
+ Op.getReg().asMCReg(), Wait, MI);
+ unsigned RequiredWait = Wait.get(AMDGPU::LOAD_CNT);
+ if (RequiredWait == ~0u)
+ continue;
+
+ UsesPendingLoad = true;
+ MaxRequiredWait = std::max(MaxRequiredWait, RequiredWait);
+ }
+ }
+ }
+
+ if (!UsesPendingLoad || MaxRequiredWait + 1 >= Outstanding)
+ return ~0u;
+
+ return MaxRequiredWait + 1;
+}
+
bool SIInsertWaitcnts::isVMEMOrFlatVMEM(const MachineInstr &MI) const {
if (SIInstrInfo::isFLAT(MI))
return TII.mayAccessVMEMThroughFlat(MI);
@@ -3329,8 +3402,11 @@ SIInsertWaitcnts::getPreheaderFlushFlags(MachineLoop *ML,
if (VgprDefDS.contains(RU))
TrackSimpleDSOpt = false;
- // Early exit if all optimizations are invalidated
- if (VMemInvalidated && !TrackSimpleDSOpt && !TrackDSFlushPoint)
+ // Early exit if all optimizations are invalidated. Keep scanning when
+ // the loop-carried-load flush experiment is on so that
+ // UsesVgprVMEMLoadedOutside is still computed.
+ if (VMemInvalidated && !TrackSimpleDSOpt && !TrackDSFlushPoint &&
+ !EnablePreheaderFlushLoopCarriedLoad)
return Flags;
// Check for flush points (DS read used in same iteration)
@@ -3363,8 +3439,11 @@ SIInsertWaitcnts::getPreheaderFlushFlags(MachineLoop *ML,
VgprDefVMEM.insert(RU);
}
}
- // Early exit if all optimizations are invalidated
- if (VMemInvalidated && !TrackSimpleDSOpt && !TrackDSFlushPoint)
+ // Early exit if all optimizations are invalidated. Keep scanning when
+ // the loop-carried-load flush experiment is on so that
+ // UsesVgprVMEMLoadedOutside is still computed.
+ if (VMemInvalidated && !TrackSimpleDSOpt && !TrackDSFlushPoint &&
+ !EnablePreheaderFlushLoopCarriedLoad)
return Flags;
}
@@ -3400,6 +3479,15 @@ SIInsertWaitcnts::getPreheaderFlushFlags(MachineLoop *ML,
(HasVMemLoad && ST.hasVmemWriteVgprInOrder())))
Flags.FlushVmCnt = true;
+ // Experiment: also flush when the loop uses a value loaded by a VMEM op
+ // outside the loop, even if the loop has its own VMEM traffic (which the
+ // decision above excludes). Retiring the outside load before the loop can
+ // reduce dynamic scoreboard stalls for long loops. Flushing is always
+ // correct, so this only affects performance; it is off by default because it
+ // pessimizes short loops.
+ if (EnablePreheaderFlushLoopCarriedLoad && UsesVgprVMEMLoadedOutside)
+ Flags.FlushVmCnt = true;
+
// DS flush decision:
// Simple DS Opt: flush if loop uses DS read values from outside
// and either has no DS reads in the loop, or DS reads whose results
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll
index b3b559025ed88..41d518f66c717 100644
--- a/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-invalid-params.ll
@@ -1,8 +1,14 @@
; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<unknown=1>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=UNKNOWN
; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<max-bytes=invalid>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=INVALID
+; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<min-align=3>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=BAD-ALIGN
+; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<only-candidate=-2>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=BAD-CANDIDATE
+; RUN: not opt -mtriple=amdgcn-amd-amdhsa -passes='amdgpu-lds-buffering<mode=unknown>' -disable-output < %s 2>&1 | FileCheck %s --check-prefix=BAD-MODE
; UNKNOWN: amdgpu-lds-buffering: invalid AMDGPU LDS buffering pass parameter 'unknown=1'
; INVALID: amdgpu-lds-buffering: invalid AMDGPU LDS buffering max-bytes 'invalid'
+; BAD-ALIGN: amdgpu-lds-buffering: invalid AMDGPU LDS buffering min-align '3'
+; BAD-CANDIDATE: amdgpu-lds-buffering: invalid AMDGPU LDS buffering only-candidate '-2'
+; BAD-MODE: amdgpu-lds-buffering: invalid AMDGPU LDS buffering mode 'unknown'
define amdgpu_kernel void @kernel() {
ret void
diff --git a/llvm/test/CodeGen/AMDGPU/lds-buffering-options.ll b/llvm/test/CodeGen/AMDGPU/lds-buffering-options.ll
new file mode 100644
index 0000000000000..fca5329cf99ca
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/lds-buffering-options.ll
@@ -0,0 +1,64 @@
+; NOTE: Do not autogenerate. This tests experimental candidate controls.
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64>' -S < %s | FileCheck %s --check-prefix=DEFAULT
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64;min-align=4>' -S < %s | FileCheck %s --check-prefix=RELAXED
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64;min-align=4;only-candidate=1>' -S < %s | FileCheck %s --check-prefix=SECOND
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64;min-align=4;only-candidate=1;mode=shadow-lds>' -S < %s | FileCheck %s --check-prefix=SHADOW
+; RUN: opt -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -passes='amdgpu-lds-buffering<max-bytes=64;min-align=4;only-candidate=1;mode=irrelevant-lds>' -S < %s | FileCheck %s --check-prefix=CONTROL
+
+; DEFAULT-LABEL: define amdgpu_kernel void @natural_alignment(
+; DEFAULT: %first = load i32, ptr addrspace(1) %p, align 4
+; DEFAULT: store i32 %first, ptr addrspace(1) %p, align 4
+; DEFAULT: %second = load i32, ptr addrspace(1) %q, align 4
+; DEFAULT: store i32 %second, ptr addrspace(1) %q, align 4
+
+; RELAXED: @natural_alignment.ldsbuf = internal unnamed_addr addrspace(3) global
+; RELAXED: @natural_alignment.ldsbuf.1 = internal unnamed_addr addrspace(3) global
+; RELAXED-LABEL: define amdgpu_kernel void @natural_alignment(
+; RELAXED-NOT: load i32, ptr addrspace(1) %p
+; RELAXED: call void @llvm.memcpy.p3.p1.i64({{.*}}%p, i64 4, i1 false)
+; RELAXED: call void @llvm.memcpy.p1.p3.i64({{.*}}%p, {{.*}}i64 4, i1 false)
+; RELAXED: call void @llvm.memcpy.p3.p1.i64({{.*}}%q, i64 4, i1 false)
+; RELAXED: call void @llvm.memcpy.p1.p3.i64({{.*}}%q, {{.*}}i64 4, i1 false)
+
+; SECOND: @natural_alignment.ldsbuf = internal unnamed_addr addrspace(3) global
+; SECOND-NOT: @natural_alignment.ldsbuf.1 =
+; SECOND-LABEL: define amdgpu_kernel void @natural_alignment(
+; SECOND: %first = load i32, ptr addrspace(1) %p, align 4
+; SECOND: store i32 %first, ptr addrspace(1) %p, align 4
+; SECOND-NOT: load i32, ptr addrspace(1) %q
+; SECOND: call void @llvm.memcpy.p3.p1.i64({{.*}}%q, i64 4, i1 false)
+; SECOND: call void @llvm.memcpy.p1.p3.i64({{.*}}%q, {{.*}}i64 4, i1 false)
+
+; SHADOW: @natural_alignment.ldsbuf = internal unnamed_addr addrspace(3) global
+; SHADOW-LABEL: define amdgpu_kernel void @natural_alignment(
+; SHADOW: %second = load i32, ptr addrspace(1) %q, align 4
+; SHADOW: store volatile i32 %second, ptr addrspace(3) {{.*}}, align 4
+; SHADOW: store i32 %second, ptr addrspace(1) %q, align 4
+; SHADOW: %ldsbuf.control = load volatile i32, ptr addrspace(3) {{.*}}, align 4
+; SHADOW-NOT: call void @llvm.memcpy
+
+; CONTROL: @natural_alignment.ldsbuf = internal unnamed_addr addrspace(3) global
+; CONTROL-NOT: @natural_alignment.ldsbuf.1 =
+; CONTROL-LABEL: define amdgpu_kernel void @natural_alignment(
+; CONTROL: %first = load i32, ptr addrspace(1) %p, align 4
+; CONTROL: store i32 %first, ptr addrspace(1) %p, align 4
+; CONTROL: %second = load i32, ptr addrspace(1) %q, align 4
+; CONTROL: store volatile i32 0, ptr addrspace(3) {{.*}}, align 4
+; CONTROL: store i32 %second, ptr addrspace(1) %q, align 4
+; CONTROL: %ldsbuf.control = load volatile i32, ptr addrspace(3) {{.*}}, align 4
+; CONTROL-NOT: call void @llvm.memcpy
+
+define amdgpu_kernel void @natural_alignment(ptr addrspace(1) %p,
+ ptr addrspace(1) %q,
+ ptr addrspace(1) %out) #0 {
+entry:
+ %first = load i32, ptr addrspace(1) %p, align 4
+ store i32 1, ptr addrspace(1) %out, align 4
+ store i32 %first, ptr addrspace(1) %p, align 4
+ %second = load i32, ptr addrspace(1) %q, align 4
+ store i32 2, ptr addrspace(1) %out, align 4
+ store i32 %second, ptr addrspace(1) %q, align 4
+ ret void
+}
+
+attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "uniform-work-group-size"="true" }
diff --git a/llvm/test/CodeGen/AMDGPU/waitcnt-pre-loop-vmem.mir b/llvm/test/CodeGen/AMDGPU/waitcnt-pre-loop-vmem.mir
new file mode 100644
index 0000000000000..62b332af9b501
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/waitcnt-pre-loop-vmem.mir
@@ -0,0 +1,160 @@
+# RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -run-pass=si-insert-waitcnts -verify-machineinstrs %s -o - | FileCheck %s --check-prefix=DISABLED
+# RUN: llc -mtriple=amdgpu9.50-amd-amdhsa -run-pass=si-insert-waitcnts -verify-machineinstrs -amdgpu-enable-pre-loop-vmem-wait %s -o - | FileCheck %s --check-prefix=ENABLED
+# RUN: llc -mtriple=amdgpu12.00-amd-amdhsa -run-pass=si-insert-waitcnts -verify-machineinstrs -amdgpu-enable-pre-loop-vmem-wait %s -o - | FileCheck %s --check-prefix=GFX12
+
+# The older load is not used in the loop, while the younger load is. With the
+# optimization enabled, retire only the older request at the start of the
+# preheader and preserve the younger request for latency hiding. Its normal
+# dependency wait remains at the first use in the loop.
+
+# DISABLED-LABEL: name: retire_older_load
+# DISABLED-LABEL: bb.1:
+# DISABLED-NOT: S_WAITCNT
+# DISABLED: S_NOP 0
+# DISABLED-NEXT: S_BRANCH %bb.2
+# DISABLED-LABEL: bb.2:
+# DISABLED: S_WAITCNT .Vmcnt_0{{$}}
+
+# ENABLED-LABEL: name: retire_older_load
+# ENABLED-LABEL: bb.1:
+# ENABLED: S_WAITCNT .Vmcnt_1{{$}}
+# ENABLED-NEXT: S_NOP 0
+# ENABLED-NEXT: S_BRANCH %bb.2
+# ENABLED-LABEL: bb.2:
+# ENABLED: S_WAITCNT .Vmcnt_0{{$}}
+
+# GFX12-LABEL: name: retire_older_load
+# GFX12-LABEL: bb.1:
+# GFX12: S_WAIT_LOADCNT 1
+# GFX12-NEXT: S_NOP 0
+# GFX12-NEXT: S_BRANCH %bb.2
+# GFX12-LABEL: bb.2:
+# GFX12: S_WAIT_LOADCNT 0
+
+---
+name: retire_older_load
+machineFunctionInfo:
+ isEntryFunction: true
+body: |
+ bb.0:
+ successors: %bb.1
+
+ $vgpr0 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr0, $sgpr0_sgpr1_sgpr2_sgpr3, 0, 0, 0, 0, implicit $exec
+ $vgpr1 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr1, $sgpr4_sgpr5_sgpr6_sgpr7, 0, 0, 0, 0, implicit $exec
+ S_BRANCH %bb.1
+
+ bb.1:
+ successors: %bb.2
+
+ S_NOP 0
+ S_BRANCH %bb.2
+
+ bb.2:
+ successors: %bb.2, %bb.3
+
+ $vgpr2 = V_ADD_U32_e32 $vgpr1, $vgpr3, implicit $exec
+ $vgpr4 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr4, $sgpr4_sgpr5_sgpr6_sgpr7, 0, 0, 0, 0, implicit $exec
+ S_CMP_LG_U32 killed $sgpr8, $sgpr9, implicit-def $scc
+ S_CBRANCH_SCC1 %bb.2, implicit killed $scc
+ S_BRANCH %bb.3
+
+ bb.3:
+ $vgpr5 = V_ADD_U32_e32 $vgpr0, $vgpr6, implicit $exec
+ S_ENDPGM 0
+
+...
+---
+
+# An empty fallthrough preheader has no first real instruction. Insert the
+# partial wait at the end of the block instead of dropping it.
+
+# ENABLED-LABEL: name: empty_preheader
+# ENABLED-LABEL: bb.1:
+# ENABLED: S_WAITCNT .Vmcnt_1{{$}}
+# ENABLED-LABEL: bb.2:
+
+# GFX12-LABEL: name: empty_preheader
+# GFX12-LABEL: bb.1:
+# GFX12: S_WAIT_LOADCNT 1
+# GFX12-LABEL: bb.2:
+
+name: empty_preheader
+machineFunctionInfo:
+ isEntryFunction: true
+body: |
+ bb.0:
+ successors: %bb.1
+
+ $vgpr0 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr0, $sgpr0_sgpr1_sgpr2_sgpr3, 0, 0, 0, 0, implicit $exec
+ $vgpr1 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr1, $sgpr4_sgpr5_sgpr6_sgpr7, 0, 0, 0, 0, implicit $exec
+ S_BRANCH %bb.1
+
+ bb.1:
+ successors: %bb.2
+
+ bb.2:
+ successors: %bb.2, %bb.3
+
+ $vgpr2 = V_ADD_U32_e32 $vgpr1, $vgpr3, implicit $exec
+ $vgpr4 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr4, $sgpr4_sgpr5_sgpr6_sgpr7, 0, 0, 0, 0, implicit $exec
+ S_CMP_LG_U32 killed $sgpr8, $sgpr9, implicit-def $scc
+ S_CBRANCH_SCC1 %bb.2, implicit killed $scc
+ S_BRANCH %bb.3
+
+ bb.3:
+ $vgpr5 = V_ADD_U32_e32 $vgpr0, $vgpr6, implicit $exec
+ S_ENDPGM 0
+
+...
+---
+
+# If the oldest outstanding load is used in the loop, no earlier request can
+# be retired without also waiting for a loop input.
+
+# ENABLED-LABEL: name: keep_oldest_load_pending
+# ENABLED-LABEL: bb.1:
+# ENABLED-NOT: S_WAITCNT
+# ENABLED: S_NOP 0
+# ENABLED-NEXT: S_BRANCH %bb.2
+# ENABLED-LABEL: bb.2:
+# ENABLED: S_WAITCNT .Vmcnt_1{{$}}
+
+# GFX12-LABEL: name: keep_oldest_load_pending
+# GFX12-LABEL: bb.1:
+# GFX12-NOT: S_WAIT_LOADCNT
+# GFX12: S_NOP 0
+# GFX12-NEXT: S_BRANCH %bb.2
+# GFX12-LABEL: bb.2:
+# GFX12: S_WAIT_LOADCNT 1
+
+name: keep_oldest_load_pending
+machineFunctionInfo:
+ isEntryFunction: true
+body: |
+ bb.0:
+ successors: %bb.1
+
+ $vgpr0 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr0, $sgpr0_sgpr1_sgpr2_sgpr3, 0, 0, 0, 0, implicit $exec
+ $vgpr1 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr1, $sgpr4_sgpr5_sgpr6_sgpr7, 0, 0, 0, 0, implicit $exec
+ S_BRANCH %bb.1
+
+ bb.1:
+ successors: %bb.2
+
+ S_NOP 0
+ S_BRANCH %bb.2
+
+ bb.2:
+ successors: %bb.2, %bb.3
+
+ $vgpr2 = V_ADD_U32_e32 $vgpr0, $vgpr3, implicit $exec
+ $vgpr4 = BUFFER_LOAD_FORMAT_X_IDXEN killed $vgpr4, $sgpr4_sgpr5_sgpr6_sgpr7, 0, 0, 0, 0, implicit $exec
+ S_CMP_LG_U32 killed $sgpr8, $sgpr9, implicit-def $scc
+ S_CBRANCH_SCC1 %bb.2, implicit killed $scc
+ S_BRANCH %bb.3
+
+ bb.3:
+ $vgpr5 = V_ADD_U32_e32 $vgpr1, $vgpr6, implicit $exec
+ S_ENDPGM 0
+
+...
More information about the llvm-commits
mailing list