[clang] [llvm] [mlir] [AMDGPU]Pin VGPR (PR #215177)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 10 20:47:24 PDT 2026
https://github.com/demonsan updated https://github.com/llvm/llvm-project/pull/215177
>From dc041b17d047d69a0b500caaf91a8c312f61af4e Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Thu, 2 Jul 2026 09:19:09 +0000
Subject: [PATCH 01/19] [AMDGPU] Add register-pinning
intrinsics/builtins/attribute for MFMA operands
Adds a way to pin a value to a specific VGPR/AGPR (and to control the register
FILE of MFMA operands) from C++/HIP, without inline asm.
Surface:
- llvm.amdgcn.pin.vgpr / llvm.amdgcn.pin.agpr intrinsics (overloaded on the
value type; register number is an immarg).
- __builtin_amdgcn_pin_{vgpr,agpr}[_v4f32,_v16f32] builtins.
- __attribute__((amdgpu_pin_vgpr(N))) / amdgpu_pin_agpr(N) declaration
attributes on local variables (N is a constant expression, incl. template
parameters), so every store to the variable is pinned automatically.
Backend (new pass SIPreColorPins, pre-RA):
- Hard-pins single-def, single-BB values to the requested physical tuple
(physreg substitution; tie/MFMA-accumulator-edge component; subreg copy-out;
cross-BB live-in recompute); falls back to a soft allocation hint otherwise.
- Drives the occupancy target from the pinned register range so a wide pinned
accumulator fits without __launch_bounds__.
- Constrains the register file and converts accumulator MFMAs to the vgprcd
form (new getMFMASrcCVDstVGPROp mapping) so the accumulator stays in VGPRs
while inputs are in AGPRs (the mixed v[D], a[A], a[B] form).
- amdgpu_pin_agpr marks the function as maybe-agpr so the AGPR MFMA form is
available.
---
clang/include/clang/Basic/Attr.td | 14 +
clang/include/clang/Basic/BuiltinsAMDGPU.td | 6 +
clang/include/clang/Sema/SemaAMDGPU.h | 4 +
clang/lib/CodeGen/CGDecl.cpp | 10 +
clang/lib/CodeGen/CGExpr.cpp | 74 ++-
clang/lib/CodeGen/CodeGenFunction.h | 11 +
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 16 +
clang/lib/Sema/SemaAMDGPU.cpp | 38 ++
clang/lib/Sema/SemaDeclAttr.cpp | 6 +
.../lib/Sema/SemaTemplateInstantiateDecl.cpp | 16 +
...a-attribute-supported-attributes-list.test | 2 +
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 15 +
llvm/lib/Target/AMDGPU/AMDGPU.h | 4 +
llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp | 12 +
.../lib/Target/AMDGPU/AMDGPUTargetMachine.cpp | 9 +
llvm/lib/Target/AMDGPU/CMakeLists.txt | 1 +
llvm/lib/Target/AMDGPU/SIInstrInfo.h | 5 +
llvm/lib/Target/AMDGPU/SIInstrInfo.td | 10 +
llvm/lib/Target/AMDGPU/SIInstructions.td | 29 ++
.../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 4 +
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 427 ++++++++++++++++++
llvm/test/CodeGen/AMDGPU/llc-pipeline.ll | 5 +
22 files changed, 717 insertions(+), 1 deletion(-)
create mode 100644 llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index d874ccb1b8653..26251a6d746d0 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -2525,6 +2525,20 @@ def AMDGPUNumVGPR : InheritableAttr {
let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
}
+def AMDGPUPinVGPR : InheritableAttr {
+ let Spellings = [Clang<"amdgpu_pin_vgpr", 0>];
+ let Args = [ExprArgument<"Reg">];
+ let Documentation = [Undocumented];
+ let Subjects = SubjectList<[Var]>;
+}
+
+def AMDGPUPinAGPR : InheritableAttr {
+ let Spellings = [Clang<"amdgpu_pin_agpr", 0>];
+ let Args = [ExprArgument<"Reg">];
+ let Documentation = [Undocumented];
+ let Subjects = SubjectList<[Var]>;
+}
+
def AMDGPUMaxNumWorkGroups : InheritableAttr {
let Spellings = [Clang<"amdgpu_max_num_work_groups", 0>];
let Args = [ExprArgument<"MaxNumWorkGroupsX">, ExprArgument<"MaxNumWorkGroupsY", 1>, ExprArgument<"MaxNumWorkGroupsZ", 1>];
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 77b402a8a63b2..2853f9084db2e 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -213,6 +213,12 @@ def __builtin_amdgcn_ds_permute : AMDGPUBuiltin<"int(int, int)", [Const]>;
def __builtin_amdgcn_ds_bpermute : AMDGPUBuiltin<"int(int, int)", [Const]>;
def __builtin_amdgcn_readfirstlane : AMDGPUBuiltin<"int(int)", [Const]>;
def __builtin_amdgcn_readlane : AMDGPUBuiltin<"int(int, int)", [Const]>;
+def __builtin_amdgcn_pin_vgpr : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
+def __builtin_amdgcn_pin_agpr : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
+def __builtin_amdgcn_pin_vgpr_v4f32 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, float>, _Constant int)", [Const]>;
+def __builtin_amdgcn_pin_agpr_v4f32 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, float>, _Constant int)", [Const]>;
+def __builtin_amdgcn_pin_vgpr_v16f32 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, float>, _Constant int)", [Const]>;
+def __builtin_amdgcn_pin_agpr_v16f32 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, float>, _Constant int)", [Const]>;
def __builtin_amdgcn_wave_shuffle : AMDGPUBuiltin<"int(int, int)", [Const]> {
let Documentation = [DocWaveShuffle];
let ArgNames = ["src", "idx"];
diff --git a/clang/include/clang/Sema/SemaAMDGPU.h b/clang/include/clang/Sema/SemaAMDGPU.h
index 0cf4f8adcf108..8909f32c7ae75 100644
--- a/clang/include/clang/Sema/SemaAMDGPU.h
+++ b/clang/include/clang/Sema/SemaAMDGPU.h
@@ -77,6 +77,10 @@ class SemaAMDGPU : public SemaBase {
void handleAMDGPUWavesPerEUAttr(Decl *D, const ParsedAttr &AL);
void handleAMDGPUNumSGPRAttr(Decl *D, const ParsedAttr &AL);
void handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL);
+ void handleAMDGPUPinVGPRAttr(Decl *D, const ParsedAttr &AL);
+ void handleAMDGPUPinAGPRAttr(Decl *D, const ParsedAttr &AL);
+ void addAMDGPUPinVGPRAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Reg);
+ void addAMDGPUPinAGPRAttr(Decl *D, const AttributeCommonInfo &CI, Expr *Reg);
void handleAMDGPUMaxNumWorkGroupsAttr(Decl *D, const ParsedAttr &AL);
void handleAMDGPUFlatWorkGroupSizeAttr(Decl *D, const ParsedAttr &AL);
diff --git a/clang/lib/CodeGen/CGDecl.cpp b/clang/lib/CodeGen/CGDecl.cpp
index 29bc47130c4cd..ee4982e8de0e0 100644
--- a/clang/lib/CodeGen/CGDecl.cpp
+++ b/clang/lib/CodeGen/CGDecl.cpp
@@ -1760,6 +1760,16 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
emission.getAllocatedAddress());
}
+ // Record amdgpu_pin_{vgpr,agpr} locals so stores to them get pinned.
+ if (D.hasAttr<AMDGPUPinVGPRAttr>() || D.hasAttr<AMDGPUPinAGPRAttr>()) {
+ bool IsAGPR = D.hasAttr<AMDGPUPinAGPRAttr>();
+ const Expr *RegE = IsAGPR ? D.getAttr<AMDGPUPinAGPRAttr>()->getReg()
+ : D.getAttr<AMDGPUPinVGPRAttr>()->getReg();
+ unsigned Reg = RegE->EvaluateKnownConstInt(getContext()).getZExtValue();
+ AMDGPUPinnedLocals[emission.getAllocatedAddress().getBasePointer()] = {
+ IsAGPR, Reg};
+ }
+
return emission;
}
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 9201e40bc13a1..68c8a10177cf8 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -45,6 +45,7 @@
#include "llvm/IR/Constants.h"
#include "llvm/IR/DataLayout.h"
#include "llvm/IR/Intrinsics.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/IntrinsicsWebAssembly.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/MDBuilder.h"
@@ -3048,7 +3049,78 @@ void CodeGenFunction::EmitStoreThroughLValue(RValue Src, LValue Dst,
}
assert(Src.isScalar() && "Can't emit an agg store with this method");
- EmitStoreOfScalar(Src.getScalarVal(), Dst, isInit);
+ llvm::Value *SV = Src.getScalarVal();
+ if (Dst.isSimple() && !AMDGPUPinnedLocals.empty())
+ SV = emitAMDGPUPinnedValue(SV, Dst.getPointer(*this));
+ EmitStoreOfScalar(SV, Dst, isInit);
+}
+
+llvm::Value *CodeGenFunction::emitAMDGPUPinnedValue(llvm::Value *V,
+ llvm::Value *Addr) {
+ auto It = AMDGPUPinnedLocals.find(Addr);
+ if (It == AMDGPUPinnedLocals.end())
+ return V;
+ bool IsAGPR = It->second.first;
+ unsigned Reg = It->second.second;
+ llvm::Type *Ty = V->getType();
+ unsigned Bits = CGM.getDataLayout().getTypeSizeInBits(Ty);
+ if (Bits == 0 || (Bits % 32) != 0)
+ return V; // only whole-dword values are pinnable
+ unsigned Lanes = Bits / 32;
+
+ llvm::Intrinsic::ID IID =
+ IsAGPR ? llvm::Intrinsic::amdgcn_pin_agpr : llvm::Intrinsic::amdgcn_pin_vgpr;
+ auto *F32 = llvm::Type::getFloatTy(getLLVMContext());
+
+ // Pin a value whose type is exactly W dwords (W in {1,4,8,16}).
+ auto pinExact = [&](llvm::Value *Chunk, unsigned RegNo) -> llvm::Value * {
+ llvm::Function *Fn = CGM.getIntrinsic(IID, {Chunk->getType()});
+ return Builder.CreateCall(
+ Fn, {Chunk, llvm::ConstantInt::get(Int32Ty, RegNo)});
+ };
+
+ // Single supported-width value: bitcast to <Lanes x float> and pin directly.
+ auto pinWidth = [&](llvm::Value *In, unsigned RegNo,
+ unsigned W) -> llvm::Value * {
+ llvm::Type *VecTy =
+ W == 1 ? (llvm::Type *)F32 : llvm::FixedVectorType::get(F32, W);
+ llvm::Value *C = Builder.CreateBitCast(In, VecTy);
+ C = pinExact(C, RegNo);
+ return C;
+ };
+
+ // Value fits a single pin (<=16 dwords: 1/4/8/16)?
+ auto roundWidth = [](unsigned L) -> unsigned {
+ if (L == 1) return 1;
+ if (L <= 4) return 4;
+ if (L <= 8) return 8;
+ return 16;
+ };
+
+ if (Lanes <= 16 && (Lanes == 1 || Lanes == 4 || Lanes == 8 || Lanes == 16)) {
+ llvm::Value *Pinned = pinWidth(V, Reg, Lanes);
+ return Builder.CreateBitCast(Pinned, Ty);
+ }
+
+ // Wide value: chunk into 16-dword pieces (register-resident via
+ // llvm.vector.{extract,insert}), pinning each to consecutive registers.
+ auto *VecF = llvm::FixedVectorType::get(F32, Lanes);
+ llvm::Value *Vec = Builder.CreateBitCast(V, VecF);
+ auto *V16 = llvm::FixedVectorType::get(F32, 16);
+ unsigned Off = 0;
+ while (Off < Lanes) {
+ unsigned W = Lanes - Off >= 16 ? 16 : roundWidth(Lanes - Off);
+ if (Off + W > Lanes)
+ break; // leave a tiny non-power tail unpinned
+ llvm::Value *Idx = llvm::ConstantInt::get(Int64Ty, Off);
+ llvm::Type *SubTy =
+ W == 16 ? (llvm::Type *)V16 : llvm::FixedVectorType::get(F32, W);
+ llvm::Value *Sub = Builder.CreateExtractVector(SubTy, Vec, Idx);
+ Sub = pinExact(Sub, Reg + Off);
+ Vec = Builder.CreateInsertVector(VecF, Vec, Sub, Idx);
+ Off += W;
+ }
+ return Builder.CreateBitCast(Vec, Ty);
}
void CodeGenFunction::EmitStoreThroughBitfieldLValue(RValue Src, LValue Dst,
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index 783f97dc354eb..b469840fc7f26 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -1554,6 +1554,17 @@ class CodeGenFunction : public CodeGenTypeCache {
/// decls.
DeclMapTy LocalDeclMap;
+ /// Local variables carrying an amdgpu_pin_{vgpr,agpr} attribute, keyed by the
+ /// variable's storage pointer. Maps to (isAGPR, startRegNo). Every value
+ /// stored to such a variable is wrapped with llvm.amdgcn.pin.* so it is
+ /// register-pinned. \sa emitAMDGPUPinnedValue
+ llvm::DenseMap<llvm::Value *, std::pair<bool, unsigned>> AMDGPUPinnedLocals;
+
+ /// If \p Addr is a pinned local's storage, return \p V wrapped with the
+ /// appropriate llvm.amdgcn.pin.* intrinsic(s) (chunked for wide values);
+ /// otherwise return \p V unchanged.
+ llvm::Value *emitAMDGPUPinnedValue(llvm::Value *V, llvm::Value *Addr);
+
// Keep track of the cleanups for callee-destructed parameters pushed to the
// cleanup stack so that they can be deactivated later.
llvm::DenseMap<const ParmVarDecl *, EHScopeStack::stable_iterator>
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index e4be3869700cd..2b9be8b75afb2 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -695,6 +695,22 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_readfirstlane:
return emitBuiltinWithOneOverloadedType<1>(*this, E,
Intrinsic::amdgcn_readfirstlane);
+ case AMDGPU::BI__builtin_amdgcn_pin_vgpr:
+ case AMDGPU::BI__builtin_amdgcn_pin_vgpr_v4f32:
+ case AMDGPU::BI__builtin_amdgcn_pin_vgpr_v16f32:
+ case AMDGPU::BI__builtin_amdgcn_pin_agpr:
+ case AMDGPU::BI__builtin_amdgcn_pin_agpr_v4f32:
+ case AMDGPU::BI__builtin_amdgcn_pin_agpr_v16f32: {
+ bool IsVGPR = BuiltinID == AMDGPU::BI__builtin_amdgcn_pin_vgpr ||
+ BuiltinID == AMDGPU::BI__builtin_amdgcn_pin_vgpr_v4f32 ||
+ BuiltinID == AMDGPU::BI__builtin_amdgcn_pin_vgpr_v16f32;
+ Intrinsic::ID IID =
+ IsVGPR ? Intrinsic::amdgcn_pin_vgpr : Intrinsic::amdgcn_pin_agpr;
+ llvm::Value *Val = EmitScalarExpr(E->getArg(0));
+ llvm::Value *Reg = EmitScalarExpr(E->getArg(1));
+ llvm::Function *F = CGM.getIntrinsic(IID, {Val->getType()});
+ return Builder.CreateCall(F, {Val, Reg});
+ }
case AMDGPU::BI__builtin_amdgcn_div_fixup:
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
case AMDGPU::BI__builtin_amdgcn_div_fixuph:
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 2208865e3237b..78768f4608f87 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -726,6 +726,44 @@ void SemaAMDGPU::handleAMDGPUNumVGPRAttr(Decl *D, const ParsedAttr &AL) {
AMDGPUNumVGPRAttr(getASTContext(), AL, NumVGPR));
}
+// Validate a pin register operand. Value-dependent expressions (e.g. template
+// parameters) are accepted as-is and re-checked at instantiation; otherwise the
+// expression must be a non-negative integer constant.
+static Expr *checkPinRegArg(Sema &S, const AttributeCommonInfo &CI, Expr *E) {
+ if (E->isValueDependent())
+ return E;
+ llvm::APSInt Val;
+ ExprResult R = S.VerifyIntegerConstantExpression(E, &Val);
+ if (R.isInvalid())
+ return nullptr;
+ if (Val.isNegative()) {
+ S.Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
+ << CI << /*non-negative*/ 1;
+ return nullptr;
+ }
+ return R.get();
+}
+
+void SemaAMDGPU::addAMDGPUPinVGPRAttr(Decl *D, const AttributeCommonInfo &CI,
+ Expr *RegExpr) {
+ if (Expr *E = checkPinRegArg(SemaRef, CI, RegExpr))
+ D->addAttr(::new (getASTContext()) AMDGPUPinVGPRAttr(getASTContext(), CI, E));
+}
+
+void SemaAMDGPU::addAMDGPUPinAGPRAttr(Decl *D, const AttributeCommonInfo &CI,
+ Expr *RegExpr) {
+ if (Expr *E = checkPinRegArg(SemaRef, CI, RegExpr))
+ D->addAttr(::new (getASTContext()) AMDGPUPinAGPRAttr(getASTContext(), CI, E));
+}
+
+void SemaAMDGPU::handleAMDGPUPinVGPRAttr(Decl *D, const ParsedAttr &AL) {
+ addAMDGPUPinVGPRAttr(D, AL, AL.getArgAsExpr(0));
+}
+
+void SemaAMDGPU::handleAMDGPUPinAGPRAttr(Decl *D, const ParsedAttr &AL) {
+ addAMDGPUPinAGPRAttr(D, AL, AL.getArgAsExpr(0));
+}
+
static bool
checkAMDGPUMaxNumWorkGroupsArguments(Sema &S, Expr *XExpr, Expr *YExpr,
Expr *ZExpr,
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 0645f99492433..b7267034961db 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -7721,6 +7721,12 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_AMDGPUNumVGPR:
S.AMDGPU().handleAMDGPUNumVGPRAttr(D, AL);
break;
+ case ParsedAttr::AT_AMDGPUPinVGPR:
+ S.AMDGPU().handleAMDGPUPinVGPRAttr(D, AL);
+ break;
+ case ParsedAttr::AT_AMDGPUPinAGPR:
+ S.AMDGPU().handleAMDGPUPinAGPRAttr(D, AL);
+ break;
case ParsedAttr::AT_AMDGPUMaxNumWorkGroups:
S.AMDGPU().handleAMDGPUMaxNumWorkGroupsAttr(D, AL);
break;
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index f1f97ca125f46..99b2aeb4d7492 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -974,6 +974,22 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
*CUDAClusterDims, New);
}
+ if (const auto *Pin = dyn_cast<AMDGPUPinVGPRAttr>(TmplAttr)) {
+ EnterExpressionEvaluationContext Unevaluated(
+ *this, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+ ExprResult R = SubstExpr(Pin->getReg(), TemplateArgs);
+ if (!R.isInvalid())
+ AMDGPU().addAMDGPUPinVGPRAttr(New, *Pin, R.get());
+ }
+
+ if (const auto *Pin = dyn_cast<AMDGPUPinAGPRAttr>(TmplAttr)) {
+ EnterExpressionEvaluationContext Unevaluated(
+ *this, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+ ExprResult R = SubstExpr(Pin->getReg(), TemplateArgs);
+ if (!R.isInvalid())
+ AMDGPU().addAMDGPUPinAGPRAttr(New, *Pin, R.get());
+ }
+
if (const auto *ParamAttr = dyn_cast<HLSLParamModifierAttr>(TmplAttr)) {
instantiateDependentHLSLParamModifierAttr(*this, TemplateArgs, ParamAttr,
Tmpl, New);
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 8bca68e2119e7..cb1bdf26a6cfd 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -7,6 +7,8 @@
// CHECK-NEXT: AMDGPUMaxNumWorkGroups (SubjectMatchRule_function)
// CHECK-NEXT: AMDGPUNumSGPR (SubjectMatchRule_function)
// CHECK-NEXT: AMDGPUNumVGPR (SubjectMatchRule_function)
+// CHECK-NEXT: AMDGPUPinAGPR (SubjectMatchRule_variable)
+// CHECK-NEXT: AMDGPUPinVGPR (SubjectMatchRule_variable)
// CHECK-NEXT: AMDGPUWavesPerEU (SubjectMatchRule_function)
// CHECK-NEXT: AVRSignal (SubjectMatchRule_function)
// CHECK-NEXT: AbiTag (SubjectMatchRule_record_not_is_union, SubjectMatchRule_variable, SubjectMatchRule_function, SubjectMatchRule_namespace)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 4ca50ff95eeca..c57baa8611e56 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2558,6 +2558,21 @@ def int_amdgcn_readfirstlane :
Intrinsic<[llvm_any_ty], [LLVMMatchType<0>],
[IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison]>;
+// Register-pinning hint. Requests that the value operand be kept in the physical
+// VGPR (int_amdgcn_pin_vgpr) or AGPR (int_amdgcn_pin_agpr) tuple starting at the
+// number given by the second (immediate) operand. Overloaded on the value type:
+// a 32/64/128-bit value pins to 1/2/4 consecutive registers starting at that
+// number. This is a soft register-allocation hint: the allocator prefers those
+// registers when feasible and falls back under pressure. Value passed unchanged.
+def int_amdgcn_pin_vgpr :
+ Intrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty],
+ [IntrNoMem, IntrWillReturn, IntrNoCallback, IntrNoFree,
+ ImmArg<ArgIndex<1>>]>;
+def int_amdgcn_pin_agpr :
+ Intrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty],
+ [IntrNoMem, IntrWillReturn, IntrNoCallback, IntrNoFree,
+ ImmArg<ArgIndex<1>>]>;
+
// The lane argument must be uniform across the currently active threads of the
// current wave. Otherwise, the result is undefined.
def int_amdgcn_readlane :
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.h b/llvm/lib/Target/AMDGPU/AMDGPU.h
index c72fa69aa1419..8fadab0ad8764 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.h
@@ -55,6 +55,7 @@ FunctionPass *createSIMemoryLegalizerPass();
FunctionPass *createSIInsertWaitcntsPass();
FunctionPass *createSIPreAllocateWWMRegsLegacyPass();
FunctionPass *createSIFormMemoryClausesLegacyPass();
+FunctionPass *createSIPreColorPinsPass();
FunctionPass *createSIPostRABundlerPass();
FunctionPass *createAMDGPUImageIntrinsicOptimizerPass(const TargetMachine *);
@@ -247,6 +248,9 @@ extern char &SIOptimizeExecMaskingLegacyID;
void initializeSIPreAllocateWWMRegsLegacyPass(PassRegistry &);
extern char &SIPreAllocateWWMRegsLegacyID;
+void initializeSIPreColorPinsPass(PassRegistry &);
+extern char &SIPreColorPinsID;
+
void initializeAMDGPUImageIntrinsicOptimizerPass(PassRegistry &);
extern char &AMDGPUImageIntrinsicOptimizerID;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp b/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
index 630ffad96e451..b31607c039e8d 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
@@ -1360,6 +1360,18 @@ struct AAAMDGPUMinAGPRAlloc
return true;
}
+ // llvm.amdgcn.pin.agpr is an explicit request to keep a value in an AGPR.
+ // Force a nonzero AGPR requirement (the top of the pinned tuple) so the
+ // function is not inferred as needing no AGPRs, which would lower the MFMA
+ // accumulator to the VGPR form and defeat the pin.
+ case Intrinsic::amdgcn_pin_agpr: {
+ unsigned RegNo =
+ cast<ConstantInt>(CB.getArgOperand(1))->getZExtValue();
+ unsigned NumRegs = divideCeil(
+ CB.getArgOperand(0)->getType()->getPrimitiveSizeInBits(), 32);
+ Maximum.takeAssumedMaximum(std::min(RegNo + NumRegs, 256u));
+ return true;
+ }
// Trap-like intrinsics such as llvm.trap and llvm.debugtrap do not have
// the nocallback attribute, so the AMDGPU attributor can conservatively
// drop all implicitly-known inputs and AGPR allocation information. Make
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
index 30ea659cbf322..666a104aefd4e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -726,6 +726,7 @@ extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializeAMDGPUTarget() {
initializeSIMemoryLegalizerLegacyPass(*PR);
initializeSIOptimizeExecMaskingLegacyPass(*PR);
initializeSIPreAllocateWWMRegsLegacyPass(*PR);
+ initializeSIPreColorPinsPass(*PR);
initializeSIFormMemoryClausesLegacyPass(*PR);
initializeSIPostRABundlerLegacyPass(*PR);
initializeGCNCreateVOPDLegacyPass(*PR);
@@ -1816,6 +1817,10 @@ bool GCNPassConfig::addGlobalInstructionSelect() {
}
void GCNPassConfig::addFastRegAlloc() {
+ // Hard-pin llvm.amdgcn.pin.* values while still in SSA form, before
+ // PHIElimination / TwoAddressInstruction.
+ addPass(createSIPreColorPinsPass());
+
// FIXME: We have to disable the verifier here because of PHIElimination +
// TwoAddressInstructions disabling it.
@@ -1835,6 +1840,10 @@ void GCNPassConfig::addPreRegAlloc() {
}
void GCNPassConfig::addOptimizedRegAlloc() {
+ // Hard-pin llvm.amdgcn.pin.* values while still in SSA form, before
+ // PHIElimination / TwoAddressInstruction / LiveIntervals.
+ addPass(createSIPreColorPinsPass());
+
if (EnableDCEInRA)
insertPass(&DetectDeadLanesID, &DeadMachineInstructionElimID);
diff --git a/llvm/lib/Target/AMDGPU/CMakeLists.txt b/llvm/lib/Target/AMDGPU/CMakeLists.txt
index b7e679a69a80d..8bd4fb10f5212 100644
--- a/llvm/lib/Target/AMDGPU/CMakeLists.txt
+++ b/llvm/lib/Target/AMDGPU/CMakeLists.txt
@@ -186,6 +186,7 @@ add_llvm_target(AMDGPUCodeGen
SIPeepholeSDWA.cpp
SIPostRABundler.cpp
SIPreAllocateWWMRegs.cpp
+ SIPreColorPins.cpp
SIPreEmitPeephole.cpp
SIProgramInfo.cpp
SIRegisterInfo.cpp
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 3a38b517415ae..9d2f58fcc0371 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -1903,6 +1903,11 @@ namespace AMDGPU {
LLVM_READONLY
int32_t getAGPRFormOp(uint32_t Opcode);
+ /// \returns the VGPR (vgprcd) form of an MFMA that uses AGPRs for srcC/vdst,
+ /// or -1. Lets an accumulator be pinned into VGPRs with AGPR inputs.
+ LLVM_READONLY
+ int getMFMASrcCVDstVGPROp(uint16_t Opcode);
+
/// \returns v_cmpx version of a v_cmp instruction.
LLVM_READONLY
int32_t getVCMPXOpFromVCMP(uint32_t Opcode);
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index 5baf52dd12407..df87b3a26352b 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -3530,6 +3530,16 @@ def getAGPRFormOp : InstrMapping {
let ValueCols = [["AGPR"]];
}
+// Map from an mfma using AGPRs for srcC/vdst to the VGPR (vgprcd) form. Used to
+// pin an accumulator into VGPRs while its inputs stay in AGPRs.
+def getMFMASrcCVDstVGPROp : InstrMapping {
+ let FilterClass = "MFMATable";
+ let RowFields = ["AGPRFormOp"];
+ let ColFields = ["AGPRFormKind"];
+ let KeyCol = ["AGPR"];
+ let ValueCols = [["VGPR"]];
+}
+
// Maps an v_cmp instruction to its v_cmpx equivalent.
def getVCMPXOpFromVCMP : InstrMapping {
let FilterClass = "VCMPVCMPXTable";
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 36cacfd47f62c..7d9a40a7e1e20 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -422,6 +422,35 @@ foreach Op = Operations in {
Op.VT, Op.RetReg, Op.Reg>;
}
+// Register-pinning hints. Lowered by EmitInstrWithCustomInserter into a COPY
+// plus a register-allocation hint requesting the numbered VGPR/AGPR tuple.
+// One pseudo per register width; the inserter derives the tuple physreg from the
+// destination register class, so it works for any width/alignment.
+// Expanded by the SIPreColorPins pass (pre-RA, still in SSA form) into either a
+// hard physical-register assignment or a soft COPY + allocation hint.
+class PinPseudo<RegisterClass DstRC, RegisterClass SrcRC> :
+ VPseudoInstSI <(outs DstRC:$vdst), (ins SrcRC:$src, i32imm:$regno), []> {
+ let hasSideEffects = 0;
+ let mayLoad = 0;
+ let mayStore = 0;
+}
+
+foreach w = [32,64,96,128,160,192,224,256,288,320,352,384,512,1024] in {
+ defvar VRC = !cast<RegisterClass>(!if(!eq(w,32), "VGPR_32", "VReg_"#w));
+ defvar ARC = !cast<RegisterClass>(!if(!eq(w,32), "AGPR_32", "AReg_"#w));
+ def PIN_VGPR_B#w : PinPseudo<VRC, VRC>;
+ def PIN_AGPR_B#w : PinPseudo<ARC, VRC>;
+}
+
+// Map each supported value type to the width-appropriate pseudo via its size.
+foreach vt = [i32, f32, v2i32, v4i32, v4f32, v8f16,
+ v8i32, v8f32, v16i32, v16f32] in {
+ def : GCNPat<(vt (int_amdgcn_pin_vgpr vt:$s, (i32 timm:$r))),
+ (!cast<Instruction>("PIN_VGPR_B"#vt.Size) $s, $r)>;
+ def : GCNPat<(vt (int_amdgcn_pin_agpr vt:$s, (i32 timm:$r))),
+ (!cast<Instruction>("PIN_AGPR_B"#vt.Size) $s, $r)>;
+}
+
let usesCustomInserter = 1, Defs = [VCC] in {
def V_ADD_U64_PSEUDO : VPseudoInstSI <
(outs VReg_64_AlignTarget:$vdst), (ins VSrc_b64:$src0, VSrc_b64:$src1),
diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
index 0207c728ea9b8..595f91bbb7ba8 100644
--- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h
@@ -1192,6 +1192,10 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunctionInfo,
return WavesPerEU;
}
+ /// Override the waves-per-EU bounds (used e.g. to let register-pinning drive
+ /// the occupancy target / VGPR budget).
+ void setWavesPerEU(unsigned Min, unsigned Max) { WavesPerEU = {Min, Max}; }
+
/// \returns Default/requested minimum number of waves per execution unit.
unsigned getMinWavesPerEU() const {
return WavesPerEU.first;
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
new file mode 100644
index 0000000000000..ed65b0781483f
--- /dev/null
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -0,0 +1,427 @@
+//===-- SIPreColorPins.cpp - Hard register pinning ------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+/// \file
+/// Lowers the PIN_{VGPR,AGPR}_B* pseudos produced from
+/// llvm.amdgcn.pin.{vgpr,agpr} into a hard register assignment ("pre-coloring").
+///
+/// The value being pinned is rewritten so that its def and all its uses
+/// reference the requested physical VGPR/AGPR tuple directly. Because the value
+/// is then a physical register in the MIR, the register allocator treats it as
+/// fixed interference and can never place it elsewhere or let another value
+/// clobber it -- unlike the soft allocation hint, this cannot be overridden by
+/// competing coalescer copy-hints (e.g. an MFMA accumulator chain).
+///
+/// Tied operands (e.g. the in-place MFMA accumulator, whose vdst is tied to
+/// src2) require care: both ends of a tie must share the same register. The
+/// pass therefore rewrites the whole *tie-connected component* of virtual
+/// registers, so a pin placed on the accumulator input also pins the tied
+/// output. Subregister references are rewritten to the corresponding physical
+/// subregister.
+///
+/// When hard pinning is not safe (a def in the component is a PHI, REG_SEQUENCE
+/// or IMPLICIT_DEF, the physical (sub)register is not a legal member of some
+/// rewritten operand's register class, or the tuple conflicts with an already
+/// hard-pinned value) the pass falls back to the soft behaviour: a COPY plus a
+/// register-allocation hint. This guarantees the pass never regresses
+/// correctness.
+///
+/// Runs pre-RA while the function is still in SSA form (before PHIElimination /
+/// TwoAddressInstruction), so each value has a single reaching def.
+//
+//===----------------------------------------------------------------------===//
+
+#include "AMDGPU.h"
+#include "GCNSubtarget.h"
+#include "SIMachineFunctionInfo.h"
+#include "SIRegisterInfo.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/DenseSet.h"
+#include "llvm/CodeGen/LivePhysRegs.h"
+#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/MachineRegisterInfo.h"
+#include "llvm/Support/CommandLine.h"
+
+using namespace llvm;
+
+#define DEBUG_TYPE "si-pre-color-pins"
+
+static cl::opt<bool> EnableHardPin(
+ "amdgpu-hard-pin-regs", cl::init(true), cl::Hidden,
+ cl::desc("Use hard register pre-coloring for llvm.amdgcn.pin.* (else soft "
+ "allocation hints only)"));
+
+namespace {
+
+class SIPreColorPins : public MachineFunctionPass {
+public:
+ static char ID;
+
+ SIPreColorPins() : MachineFunctionPass(ID) {}
+
+ bool runOnMachineFunction(MachineFunction &MF) override;
+
+ StringRef getPassName() const override {
+ return "SI pre-color pinned registers";
+ }
+
+ void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.setPreservesCFG();
+ MachineFunctionPass::getAnalysisUsage(AU);
+ }
+};
+
+} // end anonymous namespace
+
+char SIPreColorPins::ID = 0;
+
+char &llvm::SIPreColorPinsID = SIPreColorPins::ID;
+
+INITIALIZE_PASS(SIPreColorPins, DEBUG_TYPE, "SI pre-color pinned registers",
+ false, false)
+
+FunctionPass *llvm::createSIPreColorPinsPass() { return new SIPreColorPins(); }
+
+static bool isPinPseudo(const SIInstrInfo *TII, const MachineInstr &MI) {
+ StringRef N = TII->getName(MI.getOpcode());
+ return N.starts_with("PIN_VGPR_B") || N.starts_with("PIN_AGPR_B");
+}
+
+// Physical register tuple a pin targets, or 0 if it is not a legal member of the
+// destination register class (e.g. a misaligned start on a target that requires
+// aligned tuples).
+static MCRegister getPinPhysReg(const SIRegisterInfo *TRI,
+ const TargetRegisterClass *RC, unsigned RegNo) {
+ unsigned First =
+ (TRI->isAGPRClass(RC) ? AMDGPU::AGPR0 : AMDGPU::VGPR0) + RegNo;
+ MCRegister PR = TRI->getRegSizeInBits(*RC) == 32
+ ? MCRegister(First)
+ : TRI->getMatchingSuperReg(First, AMDGPU::sub0, RC);
+ if (PR && RC->contains(PR))
+ return PR;
+ return MCRegister();
+}
+
+bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
+ const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
+ const SIInstrInfo *TII = ST.getInstrInfo();
+ const SIRegisterInfo *TRI = ST.getRegisterInfo();
+ MachineRegisterInfo &MRI = MF.getRegInfo();
+
+ SmallVector<MachineInstr *, 8> Pins;
+ for (MachineBasicBlock &MBB : MF)
+ for (MachineInstr &MI : MBB)
+ if (isPinPseudo(TII, MI))
+ Pins.push_back(&MI);
+
+ if (Pins.empty())
+ return false;
+
+ // Regunits already claimed by a hard pin. A later pin whose tuple overlaps any
+ // claimed unit falls back to soft, so two distinct simultaneously-live values
+ // can never be forced into the same physical register. (Legitimate reuse of a
+ // register by a single value -- e.g. an accumulation chain -- is absorbed by
+ // the tie-connected component of the first pin, after which the value is
+ // already physical and later pins on it are no-ops.)
+ DenseSet<MCRegUnit> Claimed;
+ bool NeedRecomputeLiveIns = false;
+ unsigned ReqVGPRs = 0, ReqAGPRs = 0; // highest register a pin needs, +1
+
+ for (MachineInstr *Pin : Pins) {
+ Register Dst = Pin->getOperand(0).getReg();
+ Register Src = Pin->getOperand(1).getReg();
+ unsigned RegNo = Pin->getOperand(2).getImm();
+ const TargetRegisterClass *RC = MRI.getRegClass(Dst);
+ MCRegister PR = getPinPhysReg(TRI, RC, RegNo);
+
+ // Record how many registers this pin needs so the pin itself can drive the
+ // occupancy target (the register budget must cover the pinned range).
+ unsigned NumRegs = TRI->getRegSizeInBits(*RC) / 32;
+ bool WantAGPR = TRI->isAGPRClass(RC);
+ if (WantAGPR)
+ ReqAGPRs = std::max(ReqAGPRs, RegNo + NumRegs);
+ else
+ ReqVGPRs = std::max(ReqVGPRs, RegNo + NumRegs);
+
+ // Constrain the register *file* of the pinned value and every vreg reachable
+ // through copies / REG_SEQUENCE / the MFMA accumulator edge to VGPR (for
+ // pin_vgpr) or AGPR (for pin_agpr). This keeps a VGPR-pinned accumulator in
+ // VGPRs even when its MFMA inputs are pinned to AGPRs (the MFMA then uses the
+ // mixed v[D], a[A], a[B] form). Unlike a physreg pin this is just a class
+ // narrowing, so it works for loop-carried PHI values too. constrainRegClass
+ // is a no-op when the target file is incompatible (e.g. a VGPR load feeding
+ // an AGPR-pinned input keeps its VGPR def and gets a copy).
+ {
+ DenseSet<Register> Seen;
+ SmallVector<Register, 8> WL;
+ SmallPtrSet<MachineInstr *, 8> AccMFMAs; // MFMAs whose vdst is pinned
+ auto AddC = [&](Register R) {
+ if (R.isVirtual() && Seen.insert(R).second)
+ WL.push_back(R);
+ };
+ AddC(Src);
+ AddC(Dst);
+ for (unsigned I = 0; I < WL.size(); ++I) {
+ for (MachineOperand &MO : MRI.reg_operands(WL[I])) {
+ MachineInstr *MI = MO.getParent();
+ if (MI->isCopy() || MI->isRegSequence()) {
+ for (MachineOperand &O : MI->operands())
+ if (O.isReg())
+ AddC(O.getReg());
+ }
+ if (MO.isTied())
+ AddC(MI->getOperand(MI->findTiedOperandIdx(MO.getOperandNo()))
+ .getReg());
+ if (TII->isMAI(*MI)) {
+ int S2 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
+ AMDGPU::OpName::src2);
+ if (S2 >= 0) {
+ if (MI->getOperand(0).isReg())
+ AddC(MI->getOperand(0).getReg());
+ if (MI->getOperand(S2).isReg())
+ AddC(MI->getOperand(S2).getReg());
+ }
+ // vdst of this MFMA is in the pinned component.
+ if (MO.isDef())
+ AccMFMAs.insert(MI);
+ }
+ }
+ }
+ // Pinning an accumulator to VGPR while its MFMA inputs are in AGPR needs
+ // the vgprcd MFMA form (VGPR dst/srcC, AGPR-or-VGPR srcA/B). ISel picks
+ // the all-AGPR form because the function needs AGPRs; convert the reached
+ // accumulator MFMAs to the vgprcd form, then re-derive the component's
+ // register classes from the rewritten (VGPR-producing) defs. src0/src1
+ // (the AGPR-pinned inputs) stay put -- vgprcd's AVSrc accepts them.
+ bool Converted = false;
+ if (!WantAGPR) {
+ for (MachineInstr *MI : AccMFMAs) {
+ int VOp = AMDGPU::getMFMASrcCVDstVGPROp(MI->getOpcode());
+ if (VOp != -1) {
+ MI->setDesc(TII->get(VOp));
+ Converted = true;
+ }
+ }
+ }
+ for (Register R : WL) {
+ // constrainRegClass cannot cross register files (AGPR<->VGPR are
+ // disjoint); after an opcode conversion the class is re-derived instead.
+ if (Converted)
+ MRI.recomputeRegClass(R);
+ unsigned Sz = TRI->getRegSizeInBits(*MRI.getRegClass(R));
+ const TargetRegisterClass *Want =
+ WantAGPR ? TRI->getAGPRClassForBitWidth(Sz)
+ : TRI->getVGPRClassForBitWidth(Sz);
+ if (Want)
+ MRI.constrainRegClass(R, Want);
+ }
+ }
+
+ bool Hard = EnableHardPin && PR && Src.isVirtual() && Dst.isVirtual();
+
+ // Grow the set of virtual registers that must share PR by following tie
+ // edges (both ends of a tied operand pair must be the same register).
+ SmallVector<Register, 8> Comp;
+ if (Hard) {
+ DenseSet<Register> Seen;
+ auto Add = [&](Register R) {
+ if (R.isVirtual() && Seen.insert(R).second)
+ Comp.push_back(R);
+ };
+ Add(Src);
+ Add(Dst);
+ for (unsigned I = 0; I < Comp.size(); ++I) {
+ Register R = Comp[I];
+ for (MachineOperand &MO : MRI.reg_operands(R)) {
+ MachineInstr *MI = MO.getParent();
+ // Follow tie edges (both ends of a tie must share the register).
+ if (MO.isTied())
+ Add(MI->getOperand(MI->findTiedOperandIdx(MO.getOperandNo()))
+ .getReg());
+ // Follow the MFMA accumulator edge (src2 <-> vdst). The VGPR (vgprcd)
+ // MFMA form is 3-address, so an accumulation chain is connected by
+ // src2->vdst def-use rather than ties; pin the whole chain as a unit.
+ if (TII->isMAI(*MI)) {
+ int Src2 =
+ AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::src2);
+ if (Src2 >= 0) {
+ const MachineOperand &V2 = MI->getOperand(Src2);
+ const MachineOperand &VD = MI->getOperand(0);
+ if ((unsigned)Src2 == MO.getOperandNo() && MI->getNumDefs() > 0 &&
+ VD.isReg())
+ Add(VD.getReg()); // src2 -> vdst
+ else if (MO.isDef() && V2.isReg())
+ Add(V2.getReg()); // vdst -> src2
+ }
+ }
+ }
+ }
+ }
+
+ // Collect and validate every operand referencing a component register.
+ SmallVector<MachineOperand *, 16> ToRewrite;
+ if (Hard) {
+ for (Register R : Comp) {
+ for (MachineInstr &DefMI : MRI.def_instructions(R)) {
+ if (DefMI.isPHI() || DefMI.isRegSequence() || DefMI.isImplicitDef()) {
+ Hard = false;
+ break;
+ }
+ }
+ if (!Hard)
+ break;
+ for (MachineOperand &MO : MRI.reg_operands(R)) {
+ if (MO.getParent() == Pin)
+ continue; // the pin itself is erased
+ MCRegister Tgt =
+ MO.getSubReg() ? TRI->getSubReg(PR, MO.getSubReg()) : PR;
+ if (!Tgt) {
+ Hard = false;
+ break;
+ }
+ const TargetRegisterClass *OpRC = MO.getParent()->getRegClassConstraint(
+ MO.getOperandNo(), TII, TRI);
+ if (OpRC && !OpRC->contains(Tgt)) {
+ Hard = false;
+ break;
+ }
+ ToRewrite.push_back(&MO);
+ }
+ if (!Hard)
+ break;
+ }
+ }
+
+ // Track whether any operand crosses basic blocks; if so we must recompute
+ // physreg live-ins after rewriting (done once at the end).
+ if (Hard) {
+ MachineBasicBlock *MBB = nullptr;
+ for (MachineOperand *MO : ToRewrite) {
+ MachineBasicBlock *B = MO->getParent()->getParent();
+ if (!MBB)
+ MBB = B;
+ else if (B != MBB) {
+ NeedRecomputeLiveIns = true;
+ break;
+ }
+ }
+ }
+
+ // Conflict with an existing hard pin on overlapping regunits?
+ if (Hard) {
+ for (MCRegUnit U : TRI->regunits(PR)) {
+ if (Claimed.contains(U)) {
+ Hard = false;
+ break;
+ }
+ }
+ }
+
+ // Partition operands. Non-tied *subregister uses* (e.g. the per-lane reads a
+ // wide accumulator feeds into stores) are not rewritten to physical
+ // subregisters -- that yields fragile physical-subreg live ranges. Instead
+ // they read a virtual copy-out of the whole tuple.
+ SmallVector<MachineOperand *, 16> DirectOps, SubUses;
+ if (Hard) {
+ for (MachineOperand *MO : ToRewrite) {
+ if (MO->isUse() && MO->getSubReg() && !MO->isTied())
+ SubUses.push_back(MO);
+ else
+ DirectOps.push_back(MO);
+ }
+ }
+
+ // If there are subregister uses, insert one "%out = COPY PR" that dominates
+ // them. Only handle the single-block case; otherwise fall back to soft.
+ MachineBasicBlock *CopyMBB = nullptr;
+ MachineBasicBlock::iterator CopyPt;
+ if (Hard && !SubUses.empty()) {
+ CopyMBB = SubUses.front()->getParent()->getParent();
+ for (MachineOperand *MO : SubUses)
+ if (MO->getParent()->getParent() != CopyMBB) {
+ Hard = false;
+ break;
+ }
+ if (Hard) {
+ // Earliest sub-use in program order becomes the insertion point.
+ DenseSet<MachineInstr *> SubMIs;
+ for (MachineOperand *MO : SubUses)
+ SubMIs.insert(MO->getParent());
+ CopyPt = CopyMBB->end();
+ for (MachineInstr &MI : *CopyMBB)
+ if (SubMIs.contains(&MI)) {
+ CopyPt = MI.getIterator();
+ break;
+ }
+ }
+ }
+
+ if (Hard) {
+ for (MachineOperand *MO : DirectOps) {
+ MCRegister Tgt =
+ MO->getSubReg() ? TRI->getSubReg(PR, MO->getSubReg()) : PR;
+ MO->setReg(Tgt);
+ MO->setSubReg(0);
+ MO->setIsRenamable(false);
+ }
+ if (!SubUses.empty()) {
+ Register Out = MRI.createVirtualRegister(RC);
+ BuildMI(*CopyMBB, CopyPt, CopyPt->getDebugLoc(),
+ TII->get(TargetOpcode::COPY), Out)
+ .addReg(PR);
+ for (MachineOperand *MO : SubUses)
+ MO->setReg(Out); // keep the subregister index
+ }
+ for (MCRegUnit U : TRI->regunits(PR))
+ Claimed.insert(U);
+ Pin->eraseFromParent();
+ continue;
+ }
+
+ // Soft fallback: COPY + register-allocation hint (a no-op hint if the
+ // physical tuple was illegal).
+ BuildMI(*Pin->getParent(), Pin, Pin->getDebugLoc(),
+ TII->get(TargetOpcode::COPY), Dst)
+ .addReg(Src);
+ if (PR) {
+ MRI.setSimpleHint(Dst, PR);
+ if (Src.isVirtual())
+ MRI.setSimpleHint(Src, PR);
+ }
+ Pin->eraseFromParent();
+ }
+
+ // Cross-BB hard pins introduce physical registers that are live across basic
+ // block boundaries; recompute physreg live-in lists so the verifier and the
+ // allocator see correct liveness.
+ if (NeedRecomputeLiveIns) {
+ SmallVector<MachineBasicBlock *, 16> MBBs;
+ for (MachineBasicBlock &MBB : MF)
+ MBBs.push_back(&MBB);
+ fullyRecomputeLiveIns(MBBs);
+ }
+
+ // Let the pins drive occupancy: the register budget must be large enough to
+ // hold every pinned register, so cap the occupancy accordingly. This lets a
+ // wide pinned accumulator (e.g. 192 VGPRs) force occupancy down without the
+ // user having to set __launch_bounds__ / amdgpu-waves-per-eu by hand.
+ auto *MFI = MF.getInfo<SIMachineFunctionInfo>();
+ unsigned Req = std::max(ReqVGPRs, ReqAGPRs);
+ if (Req) {
+ // Occupancy achievable while reserving `Req` registers per wave; cap the
+ // waves-per-EU (and hence the RA's VGPR budget) so the pinned range fits.
+ unsigned Occ = ST.getOccupancyWithNumVGPRs(Req);
+ auto WPE = MFI->getWavesPerEU();
+ unsigned NewMax = WPE.second ? std::min(WPE.second, Occ) : Occ;
+ unsigned NewMin = std::min(WPE.first ? WPE.first : NewMax, NewMax);
+ MFI->setWavesPerEU(NewMin, NewMax);
+ MFI->limitOccupancy(NewMax);
+ }
+
+ return true;
+}
diff --git a/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll b/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
index bb7ed3b58f8af..20cbb8873a1cf 100644
--- a/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
+++ b/llvm/test/CodeGen/AMDGPU/llc-pipeline.ll
@@ -113,6 +113,7 @@
; GCN-O0-NEXT: Finalize ISel and expand pseudo-instructions
; GCN-O0-NEXT: Local Stack Slot Allocation
; GCN-O0-NEXT: Register Usage Information Propagation
+; GCN-O0-NEXT: SI pre-color pinned registers
; GCN-O0-NEXT: Eliminate PHI nodes for register allocation
; GCN-O0-NEXT: SI Lower control flow pseudo instructions
; GCN-O0-NEXT: Two-Address instruction pass
@@ -355,6 +356,7 @@
; GCN-O1-NEXT: SI Shrink Instructions
; GCN-O1-NEXT: Register Usage Information Propagation
; GCN-O1-NEXT: AMDGPU Prepare AGPR Alloc
+; GCN-O1-NEXT: SI pre-color pinned registers
; GCN-O1-NEXT: Detect Dead Lanes
; GCN-O1-NEXT: Remove dead machine instructions
; GCN-O1-NEXT: Init Undef Pass
@@ -684,6 +686,7 @@
; GCN-O1-OPTS-NEXT: SI Shrink Instructions
; GCN-O1-OPTS-NEXT: Register Usage Information Propagation
; GCN-O1-OPTS-NEXT: AMDGPU Prepare AGPR Alloc
+; GCN-O1-OPTS-NEXT: SI pre-color pinned registers
; GCN-O1-OPTS-NEXT: Detect Dead Lanes
; GCN-O1-OPTS-NEXT: Remove dead machine instructions
; GCN-O1-OPTS-NEXT: Init Undef Pass
@@ -1017,6 +1020,7 @@
; GCN-O2-NEXT: SI Shrink Instructions
; GCN-O2-NEXT: Register Usage Information Propagation
; GCN-O2-NEXT: AMDGPU Prepare AGPR Alloc
+; GCN-O2-NEXT: SI pre-color pinned registers
; GCN-O2-NEXT: Detect Dead Lanes
; GCN-O2-NEXT: Remove dead machine instructions
; GCN-O2-NEXT: Init Undef Pass
@@ -1366,6 +1370,7 @@
; GCN-O3-NEXT: SI Shrink Instructions
; GCN-O3-NEXT: Register Usage Information Propagation
; GCN-O3-NEXT: AMDGPU Prepare AGPR Alloc
+; GCN-O3-NEXT: SI pre-color pinned registers
; GCN-O3-NEXT: Detect Dead Lanes
; GCN-O3-NEXT: Remove dead machine instructions
; GCN-O3-NEXT: Init Undef Pass
>From 7a1c352e29afd0d9d9716b42e86f95bf28ef6995 Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Thu, 2 Jul 2026 11:47:26 +0000
Subject: [PATCH 02/19] [AMDGPU] Scope register-pinning to the
amdgpu_pin_vgpr/agpr attribute
Public API is the declaration attribute amdgpu_pin_vgpr(N) /
amdgpu_pin_agpr(N); remove the exploratory __builtin_amdgcn_pin_* builtins
(the attribute lowers to the llvm.amdgcn.pin.* intrinsics directly).
Make the attribute robust and correct on real MFMA kernels:
- Front-end chunker: pin arbitrary dword widths using an i32 base
(i32/v2i32/v4i32/v8i32/v16i32 all have selection patterns); the previous
float base needed v1f32/v2f32 which have none and crashed isel on 32/64-bit
values.
- SIFoldOperands: recognize PIN_AGPR as an AGPR terminator so a load feeding
it is folded to an AGPR-born load (buffer/global/ds), no vgpr->agpr copy.
- PIN_AGPR source class widened to AV so the folded AGPR source is legal.
- SIPreColorPins: follow PHI edges when routing the accumulator to VGPR
(keeps a loop-carried accumulator in VGPR); rewrite AGPR const inits
(V_ACCVGPR_WRITE imm) to VGPR V_MOV so clear()==0 needs no agpr copy;
hard-pin REG_SEQUENCE load tuples to fixed AGPRs so A/B placement survives
low register pressure.
- Correctness: a pin whose source is a sub-register of a shared load (e.g.
ds_read2 loading two fragments into one wide reg) is made a no-op instead of
emitting overlapping physreg copies that miscompiled.
Result on gfx950 (MI355X), verified: v_mfma v[C], a[A], a[B] with A/B loaded
directly into AGPR and the accumulator in VGPR, zero v_accvgpr, spill-free;
LDS-staged inputs (global->LDS->register) also land in AGPR correctly.
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 6 -
clang/lib/CodeGen/CGExpr.cpp | 72 ++---
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 16 -
llvm/lib/Target/AMDGPU/SIFoldOperands.cpp | 8 +-
llvm/lib/Target/AMDGPU/SIInstructions.td | 6 +-
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 338 ++++++++++++++++----
6 files changed, 320 insertions(+), 126 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 2853f9084db2e..77b402a8a63b2 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -213,12 +213,6 @@ def __builtin_amdgcn_ds_permute : AMDGPUBuiltin<"int(int, int)", [Const]>;
def __builtin_amdgcn_ds_bpermute : AMDGPUBuiltin<"int(int, int)", [Const]>;
def __builtin_amdgcn_readfirstlane : AMDGPUBuiltin<"int(int)", [Const]>;
def __builtin_amdgcn_readlane : AMDGPUBuiltin<"int(int, int)", [Const]>;
-def __builtin_amdgcn_pin_vgpr : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
-def __builtin_amdgcn_pin_agpr : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
-def __builtin_amdgcn_pin_vgpr_v4f32 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, float>, _Constant int)", [Const]>;
-def __builtin_amdgcn_pin_agpr_v4f32 : AMDGPUBuiltin<"_ExtVector<4, float>(_ExtVector<4, float>, _Constant int)", [Const]>;
-def __builtin_amdgcn_pin_vgpr_v16f32 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, float>, _Constant int)", [Const]>;
-def __builtin_amdgcn_pin_agpr_v16f32 : AMDGPUBuiltin<"_ExtVector<16, float>(_ExtVector<16, float>, _Constant int)", [Const]>;
def __builtin_amdgcn_wave_shuffle : AMDGPUBuiltin<"int(int, int)", [Const]> {
let Documentation = [DocWaveShuffle];
let ArgNames = ["src", "idx"];
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 68c8a10177cf8..ba73d8474b266 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -3070,54 +3070,44 @@ llvm::Value *CodeGenFunction::emitAMDGPUPinnedValue(llvm::Value *V,
llvm::Intrinsic::ID IID =
IsAGPR ? llvm::Intrinsic::amdgcn_pin_agpr : llvm::Intrinsic::amdgcn_pin_vgpr;
- auto *F32 = llvm::Type::getFloatTy(getLLVMContext());
- // Pin a value whose type is exactly W dwords (W in {1,4,8,16}).
- auto pinExact = [&](llvm::Value *Chunk, unsigned RegNo) -> llvm::Value * {
- llvm::Function *Fn = CGM.getIntrinsic(IID, {Chunk->getType()});
- return Builder.CreateCall(
- Fn, {Chunk, llvm::ConstantInt::get(Int32Ty, RegNo)});
- };
+ // Pin selection patterns exist for i32-based widths 1/2/4/8/16 dwords
+ // (i32, v2i32, v4i32, v8i32, v16i32). Use the i32 element type: a float base
+ // would need v1f32 / v2f32, which have no pattern and crash isel. Any dword
+ // count decomposes into a descending sequence of these widths (e.g. 12 -> 8+4,
+ // 2 -> a single v2i32, a 3-dword tail -> 2 + 1).
+ llvm::Type *I32 = Int32Ty;
+ auto *VecI = llvm::FixedVectorType::get(I32, Lanes);
+ llvm::Value *Vec = Builder.CreateBitCast(V, VecI);
- // Single supported-width value: bitcast to <Lanes x float> and pin directly.
- auto pinWidth = [&](llvm::Value *In, unsigned RegNo,
- unsigned W) -> llvm::Value * {
- llvm::Type *VecTy =
- W == 1 ? (llvm::Type *)F32 : llvm::FixedVectorType::get(F32, W);
- llvm::Value *C = Builder.CreateBitCast(In, VecTy);
- C = pinExact(C, RegNo);
- return C;
+ auto pin = [&](llvm::Value *Chunk, unsigned RegNo) -> llvm::Value * {
+ llvm::Function *Fn = CGM.getIntrinsic(IID, {Chunk->getType()});
+ return Builder.CreateCall(Fn,
+ {Chunk, llvm::ConstantInt::get(Int32Ty, RegNo)});
};
-
- // Value fits a single pin (<=16 dwords: 1/4/8/16)?
- auto roundWidth = [](unsigned L) -> unsigned {
- if (L == 1) return 1;
- if (L <= 4) return 4;
- if (L <= 8) return 8;
- return 16;
+ auto floorWidth = [](unsigned L) -> unsigned {
+ if (L >= 16) return 16;
+ if (L >= 8) return 8;
+ if (L >= 4) return 4;
+ if (L >= 2) return 2;
+ return 1;
};
- if (Lanes <= 16 && (Lanes == 1 || Lanes == 4 || Lanes == 8 || Lanes == 16)) {
- llvm::Value *Pinned = pinWidth(V, Reg, Lanes);
- return Builder.CreateBitCast(Pinned, Ty);
- }
-
- // Wide value: chunk into 16-dword pieces (register-resident via
- // llvm.vector.{extract,insert}), pinning each to consecutive registers.
- auto *VecF = llvm::FixedVectorType::get(F32, Lanes);
- llvm::Value *Vec = Builder.CreateBitCast(V, VecF);
- auto *V16 = llvm::FixedVectorType::get(F32, 16);
unsigned Off = 0;
while (Off < Lanes) {
- unsigned W = Lanes - Off >= 16 ? 16 : roundWidth(Lanes - Off);
- if (Off + W > Lanes)
- break; // leave a tiny non-power tail unpinned
- llvm::Value *Idx = llvm::ConstantInt::get(Int64Ty, Off);
- llvm::Type *SubTy =
- W == 16 ? (llvm::Type *)V16 : llvm::FixedVectorType::get(F32, W);
- llvm::Value *Sub = Builder.CreateExtractVector(SubTy, Vec, Idx);
- Sub = pinExact(Sub, Reg + Off);
- Vec = Builder.CreateInsertVector(VecF, Vec, Sub, Idx);
+ unsigned W = floorWidth(Lanes - Off);
+ if (W == 1) {
+ llvm::Value *Idx = llvm::ConstantInt::get(Int32Ty, Off);
+ llvm::Value *Elt = Builder.CreateExtractElement(Vec, Idx);
+ Elt = pin(Elt, Reg + Off);
+ Vec = Builder.CreateInsertElement(Vec, Elt, Idx);
+ } else {
+ llvm::Value *Idx = llvm::ConstantInt::get(Int64Ty, Off);
+ llvm::Value *Sub = Builder.CreateExtractVector(
+ llvm::FixedVectorType::get(I32, W), Vec, Idx);
+ Sub = pin(Sub, Reg + Off);
+ Vec = Builder.CreateInsertVector(VecI, Vec, Sub, Idx);
+ }
Off += W;
}
return Builder.CreateBitCast(Vec, Ty);
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 2b9be8b75afb2..e4be3869700cd 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -695,22 +695,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_readfirstlane:
return emitBuiltinWithOneOverloadedType<1>(*this, E,
Intrinsic::amdgcn_readfirstlane);
- case AMDGPU::BI__builtin_amdgcn_pin_vgpr:
- case AMDGPU::BI__builtin_amdgcn_pin_vgpr_v4f32:
- case AMDGPU::BI__builtin_amdgcn_pin_vgpr_v16f32:
- case AMDGPU::BI__builtin_amdgcn_pin_agpr:
- case AMDGPU::BI__builtin_amdgcn_pin_agpr_v4f32:
- case AMDGPU::BI__builtin_amdgcn_pin_agpr_v16f32: {
- bool IsVGPR = BuiltinID == AMDGPU::BI__builtin_amdgcn_pin_vgpr ||
- BuiltinID == AMDGPU::BI__builtin_amdgcn_pin_vgpr_v4f32 ||
- BuiltinID == AMDGPU::BI__builtin_amdgcn_pin_vgpr_v16f32;
- Intrinsic::ID IID =
- IsVGPR ? Intrinsic::amdgcn_pin_vgpr : Intrinsic::amdgcn_pin_agpr;
- llvm::Value *Val = EmitScalarExpr(E->getArg(0));
- llvm::Value *Reg = EmitScalarExpr(E->getArg(1));
- llvm::Function *F = CGM.getIntrinsic(IID, {Val->getType()});
- return Builder.CreateCall(F, {Val, Reg});
- }
case AMDGPU::BI__builtin_amdgcn_div_fixup:
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
case AMDGPU::BI__builtin_amdgcn_div_fixuph:
diff --git a/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp b/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp
index 04a9ed487d655..857a0f838778e 100644
--- a/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp
+++ b/llvm/lib/Target/AMDGPU/SIFoldOperands.cpp
@@ -2724,9 +2724,15 @@ bool SIFoldOperandsImpl::tryFoldLoad(MachineInstr &MI) {
return false;
// Check that all uses a copy to an agpr or a reg_sequence producing an agpr.
+ // A PIN_AGPR_B* pseudo (from llvm.amdgcn.pin.agpr) is also an AGPR terminator:
+ // its def is already an AGPR tuple, so a load feeding it should be folded into
+ // AGPR (a native buffer_load into AGPR) exactly like a copy-to-AGPR.
+ auto IsPinAgpr = [&](const MachineInstr &MI) {
+ return TII->getName(MI.getOpcode()).starts_with("PIN_AGPR_B");
+ };
while (!Users.empty()) {
const MachineInstr *I = Users.pop_back_val();
- if (!I->isCopy() && !I->isRegSequence())
+ if (!I->isCopy() && !I->isRegSequence() && !IsPinAgpr(*I))
return false;
Register DstReg = I->getOperand(0).getReg();
// Physical registers may have more than one instruction definitions
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 7d9a40a7e1e20..efee1706c6fc5 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -438,8 +438,12 @@ class PinPseudo<RegisterClass DstRC, RegisterClass SrcRC> :
foreach w = [32,64,96,128,160,192,224,256,288,320,352,384,512,1024] in {
defvar VRC = !cast<RegisterClass>(!if(!eq(w,32), "VGPR_32", "VReg_"#w));
defvar ARC = !cast<RegisterClass>(!if(!eq(w,32), "AGPR_32", "AReg_"#w));
+ // AV source: the pinned value arrives in a VGPR, but SIFoldOperands' AGPR load
+ // fold may rewrite it to an AGPR in place (a native buffer_load into AGPR), so
+ // the source must accept either file.
+ defvar AVRC = !cast<RegisterClass>(!if(!eq(w,32), "AV_32", "AV_"#w));
def PIN_VGPR_B#w : PinPseudo<VRC, VRC>;
- def PIN_AGPR_B#w : PinPseudo<ARC, VRC>;
+ def PIN_AGPR_B#w : PinPseudo<ARC, AVRC>;
}
// Map each supported value type to the width-appropriate pseudo via its size.
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index ed65b0781483f..574bf6013ddb6 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -56,6 +56,29 @@ static cl::opt<bool> EnableHardPin(
cl::desc("Use hard register pre-coloring for llvm.amdgcn.pin.* (else soft "
"allocation hints only)"));
+// When an MFMA input is pinned to AGPR, either force the accumulator into VGPR
+// via the mixed vgprcd form (option 1: v[C], a[A], a[B]) or leave the
+// hardware-native all-AGPR form untouched (option 2: a[D], a[A], a[B], a[C]).
+static cl::opt<bool> PinAgprVgprC(
+ "amdgpu-pin-agpr-vgpr-c", cl::init(true), cl::Hidden,
+ cl::desc("For an AGPR-pinned MFMA input, convert the consuming MFMA to the "
+ "vgprcd form so its accumulator stays in VGPR (else keep the "
+ "native all-AGPR form)"));
+
+// Extra VGPRs (beyond the pinned accumulator's own footprint) the occupancy cap
+// must reserve for addressing / load temporaries so the accumulator stays
+// resident. Chosen so both a 64-VGPR (128x128) and a 96-VGPR (192x128) tile stay
+// spill-free without __launch_bounds__.
+// Experimental: when >0, an AGPR-input pin caps occupancy so the vgprcd-pinned
+// accumulator (plus this many VGPRs of headroom) stays resident, avoiding
+// __launch_bounds__. Default 0 (off): auto-driving occupancy from this pass
+// currently perturbs the hard-pinned physreg live ranges and can produce invalid
+// MIR at low occupancy -- use __launch_bounds__ to control occupancy instead.
+static cl::opt<unsigned> PinAccVGPRMargin(
+ "amdgpu-pin-acc-vgpr-margin", cl::init(0), cl::Hidden,
+ cl::desc("If nonzero, VGPRs reserved on top of a vgprcd-pinned accumulator "
+ "so an AGPR-input pin can drive occupancy (experimental)"));
+
namespace {
class SIPreColorPins : public MachineFunctionPass {
@@ -131,6 +154,12 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
DenseSet<MCRegUnit> Claimed;
bool NeedRecomputeLiveIns = false;
unsigned ReqVGPRs = 0, ReqAGPRs = 0; // highest register a pin needs, +1
+ // Accumulator tiles moved to VGPR by the vgprcd conversion (A,B->AGPR pins).
+ // Their total VGPR footprint drives occupancy: moving A/B out of the VGPR file
+ // lets the compiler raise occupancy, shrinking the per-wave VGPR budget until
+ // the (now VGPR) accumulator no longer fits and spills/rotates through AGPRs.
+ // Capping occupancy so the accumulator stays resident avoids that.
+ DenseSet<Register> AccTiles;
for (MachineInstr *Pin : Pins) {
Register Dst = Pin->getOperand(0).getReg();
@@ -157,73 +186,244 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
// is a no-op when the target file is incompatible (e.g. a VGPR load feeding
// an AGPR-pinned input keeps its VGPR def and gets a copy).
{
- DenseSet<Register> Seen;
- SmallVector<Register, 8> WL;
- SmallPtrSet<MachineInstr *, 8> AccMFMAs; // MFMAs whose vdst is pinned
- auto AddC = [&](Register R) {
- if (R.isVirtual() && Seen.insert(R).second)
- WL.push_back(R);
- };
- AddC(Src);
- AddC(Dst);
- for (unsigned I = 0; I < WL.size(); ++I) {
- for (MachineOperand &MO : MRI.reg_operands(WL[I])) {
- MachineInstr *MI = MO.getParent();
- if (MI->isCopy() || MI->isRegSequence()) {
- for (MachineOperand &O : MI->operands())
- if (O.isReg())
- AddC(O.getReg());
- }
- if (MO.isTied())
- AddC(MI->getOperand(MI->findTiedOperandIdx(MO.getOperandNo()))
- .getReg());
- if (TII->isMAI(*MI)) {
- int S2 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
- AMDGPU::OpName::src2);
- if (S2 >= 0) {
- if (MI->getOperand(0).isReg())
- AddC(MI->getOperand(0).getReg());
- if (MI->getOperand(S2).isReg())
- AddC(MI->getOperand(S2).getReg());
+ // Gather the copy/REG_SEQUENCE/tie-connected component of `Seeds` and
+ // constrain every member to the requested register file. MFMA
+ // src2<->vdst accumulator edges are followed only when `FollowAcc` is set.
+ // Otherwise an MFMA that *uses* a component register as src0/src1 is
+ // recorded in `Inputs` and treated as a leaf, so pinning an input to AGPR
+ // does not drag the (large, loop-carried) accumulator into the AGPR file.
+ // `Recompute` re-derives each class from its defs first -- needed after an
+ // opcode conversion, since constrainRegClass cannot cross the disjoint
+ // AGPR/VGPR files.
+ auto constrainComponent = [&](ArrayRef<Register> Seeds, bool AGPRFile,
+ bool FollowAcc, bool Recompute,
+ SmallPtrSetImpl<MachineInstr *> &Inputs) {
+ DenseSet<Register> Seen;
+ SmallVector<Register, 16> WL;
+ auto Add = [&](Register R) {
+ if (R.isVirtual() && Seen.insert(R).second)
+ WL.push_back(R);
+ };
+ for (Register R : Seeds)
+ Add(R);
+ for (unsigned I = 0; I < WL.size(); ++I) {
+ for (MachineOperand &MO : MRI.reg_operands(WL[I])) {
+ MachineInstr *MI = MO.getParent();
+ // Copy / REG_SEQUENCE / PHI all just move the value between vregs;
+ // pull every register operand into the component. PHI matters for
+ // the loop-carried accumulator: without it the carried value stays
+ // in its original file (AGPR) while the vgprcd MFMA computes in VGPR,
+ // forcing an agpr<->vgpr copy every iteration.
+ if (MI->isCopy() || MI->isRegSequence() || MI->isPHI()) {
+ for (MachineOperand &O : MI->operands())
+ if (O.isReg())
+ Add(O.getReg());
+ }
+ if (MO.isTied())
+ Add(MI->getOperand(MI->findTiedOperandIdx(MO.getOperandNo()))
+ .getReg());
+ if (TII->isMAI(*MI)) {
+ int S0 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
+ AMDGPU::OpName::src0);
+ int S1 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
+ AMDGPU::OpName::src1);
+ int S2 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
+ AMDGPU::OpName::src2);
+ unsigned OpNo = MO.getOperandNo();
+ bool IsInput = (S0 >= 0 && OpNo == (unsigned)S0) ||
+ (S1 >= 0 && OpNo == (unsigned)S1);
+ if (IsInput && !FollowAcc) {
+ Inputs.insert(MI);
+ } else if (FollowAcc && S2 >= 0) {
+ if (MI->getOperand(0).isReg())
+ Add(MI->getOperand(0).getReg());
+ if (MI->getOperand(S2).isReg())
+ Add(MI->getOperand(S2).getReg());
+ }
}
- // vdst of this MFMA is in the pinned component.
- if (MO.isDef())
- AccMFMAs.insert(MI);
}
}
- }
- // Pinning an accumulator to VGPR while its MFMA inputs are in AGPR needs
- // the vgprcd MFMA form (VGPR dst/srcC, AGPR-or-VGPR srcA/B). ISel picks
- // the all-AGPR form because the function needs AGPRs; convert the reached
- // accumulator MFMAs to the vgprcd form, then re-derive the component's
- // register classes from the rewritten (VGPR-producing) defs. src0/src1
- // (the AGPR-pinned inputs) stay put -- vgprcd's AVSrc accepts them.
- bool Converted = false;
- if (!WantAGPR) {
- for (MachineInstr *MI : AccMFMAs) {
+ for (Register R : WL) {
+ // A constant accumulator init (e.g. clear()==0) materialized in an
+ // AGPR via V_ACCVGPR_WRITE cannot be constrained to VGPR (its dst is
+ // AGPR-only), so it would stay in AGPR and be copied into the VGPR
+ // accumulator every kernel launch (write-0-to-agpr then read-to-vgpr).
+ // When routing the accumulator to VGPR, rewrite such an init to a
+ // plain VGPR V_MOV so the constant is born in VGPR (no agpr<->vgpr copy).
+ if (!AGPRFile)
+ for (MachineInstr &Def :
+ make_early_inc_range(MRI.def_instructions(R))) {
+ if (Def.getOpcode() == AMDGPU::V_ACCVGPR_WRITE_B32_e64 &&
+ Def.getNumOperands() >= 2 && Def.getOperand(1).isImm())
+ Def.setDesc(TII->get(AMDGPU::V_MOV_B32_e32));
+ }
+ if (Recompute)
+ MRI.recomputeRegClass(R);
+ unsigned Sz = TRI->getRegSizeInBits(*MRI.getRegClass(R));
+ const TargetRegisterClass *Want =
+ AGPRFile ? TRI->getAGPRClassForBitWidth(Sz)
+ : TRI->getVGPRClassForBitWidth(Sz);
+ if (Want)
+ MRI.constrainRegClass(R, Want);
+ }
+ };
+
+ SmallPtrSet<MachineInstr *, 8> InputMFMAs;
+ Register Seeds[] = {Src, Dst};
+ // Constrain the pinned value's own component to its file. For an AGPR
+ // input pin, stop at the MFMAs that consume it (recorded in InputMFMAs).
+ constrainComponent(Seeds, /*AGPRFile=*/WantAGPR, /*FollowAcc=*/!WantAGPR,
+ /*Recompute=*/false, InputMFMAs);
+
+ // An AGPR-pinned MFMA input needs the mixed vgprcd form (VGPR dst/srcC,
+ // AGPR-or-VGPR srcA/B) so the accumulator can stay in VGPR. ISel picks the
+ // all-AGPR form because the function needs AGPRs; convert each consuming
+ // MFMA to vgprcd, then constrain its accumulator (vdst/srcC chain) to VGPR
+ // -- re-deriving classes from the converted, VGPR-producing defs. This
+ // keeps the whole accumulation chain in VGPR without pinning it, so it
+ // stays coalesced (no chunked pins, no agpr<->vgpr shuffle).
+ if (WantAGPR && PinAgprVgprC && !InputMFMAs.empty()) {
+ SmallVector<Register, 8> AccSeeds;
+ for (MachineInstr *MI : InputMFMAs) {
int VOp = AMDGPU::getMFMASrcCVDstVGPROp(MI->getOpcode());
- if (VOp != -1) {
- MI->setDesc(TII->get(VOp));
- Converted = true;
+ if (VOp == -1)
+ continue; // already vgprcd form
+ MI->setDesc(TII->get(VOp));
+ if (MI->getOperand(0).isReg()) {
+ AccSeeds.push_back(MI->getOperand(0).getReg());
+ // Each converted MFMA's vdst is one accumulator tile now living in
+ // VGPR; track distinct tiles for the occupancy cap below.
+ AccTiles.insert(MI->getOperand(0).getReg());
}
+ int S2 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
+ AMDGPU::OpName::src2);
+ if (S2 >= 0 && MI->getOperand(S2).isReg())
+ AccSeeds.push_back(MI->getOperand(S2).getReg());
+ }
+ if (!AccSeeds.empty()) {
+ SmallPtrSet<MachineInstr *, 8> Ignore;
+ constrainComponent(AccSeeds, /*AGPRFile=*/false, /*FollowAcc=*/true,
+ /*Recompute=*/true, Ignore);
}
}
- for (Register R : WL) {
- // constrainRegClass cannot cross register files (AGPR<->VGPR are
- // disjoint); after an opcode conversion the class is re-derived instead.
- if (Converted)
- MRI.recomputeRegClass(R);
- unsigned Sz = TRI->getRegSizeInBits(*MRI.getRegClass(R));
- const TargetRegisterClass *Want =
- WantAGPR ? TRI->getAGPRClassForBitWidth(Sz)
- : TRI->getVGPRClassForBitWidth(Sz);
- if (Want)
- MRI.constrainRegClass(R, Want);
+ }
+
+ // When the pinned source is a *subregister* of a larger value, that register
+ // is shared -- e.g. a combined ds_read2 loads two pinned fragments into one
+ // wide register, each pin taking a sub-slice. Neither a hard pin (rewriting
+ // the whole wide reg to one narrow physreg) nor a soft COPY+hint is safe: the
+ // soft copies read/write overlapping physreg sub-slices and the allocator
+ // clobbers one before the other is read (miscompile). But the shared load was
+ // already class-constrained to the requested file above (and tryFoldLoad put
+ // it in AGPR), so the pin is redundant -- make it a no-op: replace uses of the
+ // pin result with the source (sub)register directly and erase the pin.
+ if (Pin->getOperand(1).getSubReg()) {
+ unsigned SubIdx = Pin->getOperand(1).getSubReg();
+ for (MachineOperand &MO :
+ llvm::make_early_inc_range(MRI.use_operands(Dst))) {
+ MO.setSubReg(TRI->composeSubRegIndices(SubIdx, MO.getSubReg()));
+ MO.setReg(Src);
}
+ Pin->eraseFromParent();
+ continue;
}
bool Hard = EnableHardPin && PR && Src.isVirtual() && Dst.isVirtual();
+ // Deterministic AGPR placement for a load tuple. When an AGPR pin's value is
+ // a REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a
+ // fixed physical AGPR sub-register so A/B are *born* in fixed AGPRs. Without
+ // this the MFMA A/B operands are AV (agpr-or-vgpr) and the coalescer /
+ // allocator moves them back to VGPR whenever pressure is low, making the pin
+ // non-deterministic. The accumulator was already routed to VGPR (vgprcd) by
+ // the file-constraint step above, so this yields v[D], a[A], a[B] with the
+ // accumulator free to occupy the whole VGPR file.
+ if (Hard && WantAGPR) {
+ MachineInstr *RS = MRI.getVRegDef(Src);
+ MachineBasicBlock *PinMBB = Pin->getParent();
+ bool Ok = RS && RS->isRegSequence() && RS->getParent() == PinMBB;
+ for (MCRegUnit U : TRI->regunits(PR))
+ if (Ok && Claimed.contains(U))
+ Ok = false;
+
+ // Collect element (reg, subreg-index) pairs.
+ SmallVector<std::pair<Register, unsigned>, 16> Elems;
+ if (Ok)
+ for (unsigned I = 1; I + 1 < RS->getNumOperands(); I += 2) {
+ const MachineOperand &Reg = RS->getOperand(I);
+ const MachineOperand &Sub = RS->getOperand(I + 1);
+ if (!Reg.isReg() || !Reg.getReg().isVirtual() || Reg.getSubReg() ||
+ !Sub.isImm() || !TRI->getSubReg(PR, Sub.getImm())) {
+ Ok = false;
+ break;
+ }
+ Elems.push_back({Reg.getReg(), (unsigned)Sub.getImm()});
+ }
+
+ // Every use of the pinned result and of each element must legally accept
+ // the physical (sub)register and live in this block.
+ auto LegalHere = [&](MachineOperand &MO, MCRegister T) {
+ if (!T || MO.getParent()->getParent() != PinMBB)
+ return false;
+ const TargetRegisterClass *OpRC =
+ MO.getParent()->getRegClassConstraint(MO.getOperandNo(), TII, TRI);
+ return !OpRC || OpRC->contains(T);
+ };
+ if (Ok)
+ for (MachineOperand &MO : MRI.reg_operands(Dst)) {
+ if (MO.getParent() == Pin)
+ continue;
+ MCRegister T = MO.getSubReg() ? TRI->getSubReg(PR, MO.getSubReg()) : PR;
+ if (!LegalHere(MO, T)) {
+ Ok = false;
+ break;
+ }
+ }
+ if (Ok)
+ for (auto [Elem, SubIdx] : Elems) {
+ MCRegister PhysSub = TRI->getSubReg(PR, SubIdx);
+ for (MachineOperand &MO : MRI.reg_operands(Elem))
+ if (!LegalHere(MO, PhysSub)) {
+ Ok = false;
+ break;
+ }
+ if (!Ok)
+ break;
+ }
+
+ if (Ok) {
+ // Point each element's def/uses at its physical AGPR sub-register.
+ for (auto [Elem, SubIdx] : Elems) {
+ MCRegister PhysSub = TRI->getSubReg(PR, SubIdx);
+ SmallVector<MachineOperand *, 4> Ops;
+ for (MachineOperand &MO : MRI.reg_operands(Elem))
+ Ops.push_back(&MO);
+ for (MachineOperand *MO : Ops) {
+ MO->setReg(PhysSub);
+ MO->setSubReg(0);
+ MO->setIsRenamable(false);
+ }
+ }
+ // Point the pinned-result uses at the physical tuple.
+ SmallVector<MachineOperand *, 16> Ops;
+ for (MachineOperand &MO : MRI.reg_operands(Dst))
+ if (MO.getParent() != Pin)
+ Ops.push_back(&MO);
+ for (MachineOperand *MO : Ops) {
+ MCRegister T = MO->getSubReg() ? TRI->getSubReg(PR, MO->getSubReg()) : PR;
+ MO->setReg(T);
+ MO->setSubReg(0);
+ MO->setIsRenamable(false);
+ }
+ for (MCRegUnit U : TRI->regunits(PR))
+ Claimed.insert(U);
+ RS->eraseFromParent();
+ Pin->eraseFromParent();
+ NeedRecomputeLiveIns = true;
+ continue;
+ }
+ }
+
// Grow the set of virtual registers that must share PR by following tie
// edges (both ends of a tied operand pair must be the same register).
SmallVector<Register, 8> Comp;
@@ -406,19 +606,35 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
fullyRecomputeLiveIns(MBBs);
}
- // Let the pins drive occupancy: the register budget must be large enough to
- // hold every pinned register, so cap the occupancy accordingly. This lets a
- // wide pinned accumulator (e.g. 192 VGPRs) force occupancy down without the
- // user having to set __launch_bounds__ / amdgpu-waves-per-eu by hand.
+ // Let a *VGPR* pin drive occupancy: a wide pinned VGPR value (e.g. a 192-VGPR
+ // accumulator) must fit the per-wave VGPR budget, so cap occupancy to make
+ // room without the user setting __launch_bounds__ / amdgpu-waves-per-eu.
+ // AGPR pins must NOT drive this: AGPRs are a separate file, and feeding an
+ // AGPR count into the VGPR occupancy formula wrongly raises occupancy and
+ // shrinks the VGPR budget, spilling the (VGPR) accumulator into AGPRs.
auto *MFI = MF.getInfo<SIMachineFunctionInfo>();
- unsigned Req = std::max(ReqVGPRs, ReqAGPRs);
+ // Total VGPR footprint of the accumulator tiles routed to VGPR, plus a margin
+ // for addressing/temps. When A/B are pinned to AGPR the accumulator must stay
+ // VGPR-resident; this caps occupancy so its budget is large enough.
+ unsigned AccVGPRs = 0;
+ if (PinAccVGPRMargin) {
+ for (Register R : AccTiles)
+ if (R.isVirtual())
+ AccVGPRs += TRI->getRegSizeInBits(*MRI.getRegClass(R)) / 32;
+ if (AccVGPRs)
+ AccVGPRs += PinAccVGPRMargin;
+ }
+ unsigned Req = std::max(ReqVGPRs, AccVGPRs);
if (Req) {
// Occupancy achievable while reserving `Req` registers per wave; cap the
// waves-per-EU (and hence the RA's VGPR budget) so the pinned range fits.
unsigned Occ = ST.getOccupancyWithNumVGPRs(Req);
auto WPE = MFI->getWavesPerEU();
unsigned NewMax = WPE.second ? std::min(WPE.second, Occ) : Occ;
- unsigned NewMin = std::min(WPE.first ? WPE.first : NewMax, NewMax);
+ // Only cap the *max* occupancy; keep the min low (1 unless the function
+ // already required more). Forcing min==max over-constrains the allocator and
+ // breaks physreg liveness for hard-pinned loop-body tuples at low occupancy.
+ unsigned NewMin = std::min(WPE.first ? WPE.first : 1u, NewMax);
MFI->setWavesPerEU(NewMin, NewMax);
MFI->limitOccupancy(NewMax);
}
>From 3857a42c0bde8fc40a9990a557755d3aed2d9a01 Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Thu, 2 Jul 2026 13:39:09 +0000
Subject: [PATCH 03/19] [AMDGPU] Condense register-pinning comments to LLVM
style
---
clang/lib/CodeGen/CGExpr.cpp | 8 +-
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 123 ++++++++--------------
2 files changed, 48 insertions(+), 83 deletions(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index ba73d8474b266..92e2e9dd8279d 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -3071,11 +3071,9 @@ llvm::Value *CodeGenFunction::emitAMDGPUPinnedValue(llvm::Value *V,
llvm::Intrinsic::ID IID =
IsAGPR ? llvm::Intrinsic::amdgcn_pin_agpr : llvm::Intrinsic::amdgcn_pin_vgpr;
- // Pin selection patterns exist for i32-based widths 1/2/4/8/16 dwords
- // (i32, v2i32, v4i32, v8i32, v16i32). Use the i32 element type: a float base
- // would need v1f32 / v2f32, which have no pattern and crash isel. Any dword
- // count decomposes into a descending sequence of these widths (e.g. 12 -> 8+4,
- // 2 -> a single v2i32, a 3-dword tail -> 2 + 1).
+ // Pin patterns exist for i32-based widths of 1/2/4/8/16 dwords (i32, v2i32,
+ // v4i32, v8i32, v16i32); a float base would need v1f32/v2f32, which have none
+ // and crash isel. Decompose any dword count into these widths (e.g. 12 -> 8+4).
llvm::Type *I32 = Int32Ty;
auto *VecI = llvm::FixedVectorType::get(I32, Lanes);
llvm::Value *Vec = Builder.CreateBitCast(V, VecI);
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 574bf6013ddb6..45423b9779c16 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -56,28 +56,22 @@ static cl::opt<bool> EnableHardPin(
cl::desc("Use hard register pre-coloring for llvm.amdgcn.pin.* (else soft "
"allocation hints only)"));
-// When an MFMA input is pinned to AGPR, either force the accumulator into VGPR
-// via the mixed vgprcd form (option 1: v[C], a[A], a[B]) or leave the
-// hardware-native all-AGPR form untouched (option 2: a[D], a[A], a[B], a[C]).
+// If set, convert an AGPR-pinned input's MFMA to the mixed vgprcd form
+// (v[C], a[A], a[B]) so the accumulator stays in VGPR; else keep the native
+// all-AGPR form (a[D], a[A], a[B], a[C]).
static cl::opt<bool> PinAgprVgprC(
"amdgpu-pin-agpr-vgpr-c", cl::init(true), cl::Hidden,
- cl::desc("For an AGPR-pinned MFMA input, convert the consuming MFMA to the "
- "vgprcd form so its accumulator stays in VGPR (else keep the "
- "native all-AGPR form)"));
-
-// Extra VGPRs (beyond the pinned accumulator's own footprint) the occupancy cap
-// must reserve for addressing / load temporaries so the accumulator stays
-// resident. Chosen so both a 64-VGPR (128x128) and a 96-VGPR (192x128) tile stay
-// spill-free without __launch_bounds__.
-// Experimental: when >0, an AGPR-input pin caps occupancy so the vgprcd-pinned
-// accumulator (plus this many VGPRs of headroom) stays resident, avoiding
-// __launch_bounds__. Default 0 (off): auto-driving occupancy from this pass
-// currently perturbs the hard-pinned physreg live ranges and can produce invalid
-// MIR at low occupancy -- use __launch_bounds__ to control occupancy instead.
+ cl::desc("Convert an AGPR-input MFMA to vgprcd to keep its accumulator in "
+ "VGPR (else keep the native all-AGPR form)"));
+
+// Experimental (default off): if nonzero, an AGPR-input pin caps occupancy so
+// the vgprcd accumulator plus this many VGPRs of headroom stay resident, in
+// place of __launch_bounds__. Driving occupancy here can perturb hard-pinned
+// physreg live ranges at low occupancy, so __launch_bounds__ is preferred.
static cl::opt<unsigned> PinAccVGPRMargin(
"amdgpu-pin-acc-vgpr-margin", cl::init(0), cl::Hidden,
- cl::desc("If nonzero, VGPRs reserved on top of a vgprcd-pinned accumulator "
- "so an AGPR-input pin can drive occupancy (experimental)"));
+ cl::desc("If nonzero, VGPRs reserved above a vgprcd accumulator so an "
+ "AGPR-input pin can drive occupancy (experimental)"));
namespace {
@@ -154,11 +148,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
DenseSet<MCRegUnit> Claimed;
bool NeedRecomputeLiveIns = false;
unsigned ReqVGPRs = 0, ReqAGPRs = 0; // highest register a pin needs, +1
- // Accumulator tiles moved to VGPR by the vgprcd conversion (A,B->AGPR pins).
- // Their total VGPR footprint drives occupancy: moving A/B out of the VGPR file
- // lets the compiler raise occupancy, shrinking the per-wave VGPR budget until
- // the (now VGPR) accumulator no longer fits and spills/rotates through AGPRs.
- // Capping occupancy so the accumulator stays resident avoids that.
+ // Accumulator tiles routed to VGPR by the vgprcd conversion; their footprint
+ // optionally drives the occupancy cap (see PinAccVGPRMargin).
DenseSet<Register> AccTiles;
for (MachineInstr *Pin : Pins) {
@@ -177,24 +168,17 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
else
ReqVGPRs = std::max(ReqVGPRs, RegNo + NumRegs);
- // Constrain the register *file* of the pinned value and every vreg reachable
- // through copies / REG_SEQUENCE / the MFMA accumulator edge to VGPR (for
- // pin_vgpr) or AGPR (for pin_agpr). This keeps a VGPR-pinned accumulator in
- // VGPRs even when its MFMA inputs are pinned to AGPRs (the MFMA then uses the
- // mixed v[D], a[A], a[B] form). Unlike a physreg pin this is just a class
- // narrowing, so it works for loop-carried PHI values too. constrainRegClass
- // is a no-op when the target file is incompatible (e.g. a VGPR load feeding
- // an AGPR-pinned input keeps its VGPR def and gets a copy).
+ // Constrain the pinned value's register file (and connected vregs) to VGPR
+ // or AGPR. This is a class narrowing, not a physreg pin, so it also works
+ // for loop-carried PHI values; it no-ops when the file is incompatible.
{
// Gather the copy/REG_SEQUENCE/tie-connected component of `Seeds` and
- // constrain every member to the requested register file. MFMA
- // src2<->vdst accumulator edges are followed only when `FollowAcc` is set.
- // Otherwise an MFMA that *uses* a component register as src0/src1 is
- // recorded in `Inputs` and treated as a leaf, so pinning an input to AGPR
- // does not drag the (large, loop-carried) accumulator into the AGPR file.
- // `Recompute` re-derives each class from its defs first -- needed after an
- // opcode conversion, since constrainRegClass cannot cross the disjoint
- // AGPR/VGPR files.
+ // constrain each member to the requested file. MFMA src2<->vdst edges are
+ // followed only when `FollowAcc`; otherwise an MFMA *using* a member as
+ // src0/src1 is recorded in `Inputs` as a leaf, so an input pin does not
+ // drag the loop-carried accumulator into the AGPR file. `Recompute`
+ // re-derives classes from defs first (needed after an opcode conversion,
+ // since constrainRegClass cannot cross the disjoint AGPR/VGPR files).
auto constrainComponent = [&](ArrayRef<Register> Seeds, bool AGPRFile,
bool FollowAcc, bool Recompute,
SmallPtrSetImpl<MachineInstr *> &Inputs) {
@@ -244,12 +228,9 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
}
for (Register R : WL) {
- // A constant accumulator init (e.g. clear()==0) materialized in an
- // AGPR via V_ACCVGPR_WRITE cannot be constrained to VGPR (its dst is
- // AGPR-only), so it would stay in AGPR and be copied into the VGPR
- // accumulator every kernel launch (write-0-to-agpr then read-to-vgpr).
- // When routing the accumulator to VGPR, rewrite such an init to a
- // plain VGPR V_MOV so the constant is born in VGPR (no agpr<->vgpr copy).
+ // A constant accumulator init (e.g. clear()==0) placed in an AGPR by
+ // V_ACCVGPR_WRITE can't be constrained to VGPR; rewrite it to V_MOV so
+ // the constant is born in VGPR instead of copied from AGPR each launch.
if (!AGPRFile)
for (MachineInstr &Def :
make_early_inc_range(MRI.def_instructions(R))) {
@@ -275,13 +256,11 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
constrainComponent(Seeds, /*AGPRFile=*/WantAGPR, /*FollowAcc=*/!WantAGPR,
/*Recompute=*/false, InputMFMAs);
- // An AGPR-pinned MFMA input needs the mixed vgprcd form (VGPR dst/srcC,
- // AGPR-or-VGPR srcA/B) so the accumulator can stay in VGPR. ISel picks the
- // all-AGPR form because the function needs AGPRs; convert each consuming
- // MFMA to vgprcd, then constrain its accumulator (vdst/srcC chain) to VGPR
- // -- re-deriving classes from the converted, VGPR-producing defs. This
- // keeps the whole accumulation chain in VGPR without pinning it, so it
- // stays coalesced (no chunked pins, no agpr<->vgpr shuffle).
+ // ISel picks the all-AGPR MFMA form when the function needs AGPRs. To keep
+ // the accumulator in VGPR, convert each consuming MFMA to vgprcd and
+ // constrain its accumulator (vdst/srcC chain) to VGPR, re-deriving classes
+ // from the converted defs. The chain stays coalesced in VGPR (no chunked
+ // pins, no agpr<->vgpr shuffle).
if (WantAGPR && PinAgprVgprC && !InputMFMAs.empty()) {
SmallVector<Register, 8> AccSeeds;
for (MachineInstr *MI : InputMFMAs) {
@@ -308,15 +287,11 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
}
- // When the pinned source is a *subregister* of a larger value, that register
- // is shared -- e.g. a combined ds_read2 loads two pinned fragments into one
- // wide register, each pin taking a sub-slice. Neither a hard pin (rewriting
- // the whole wide reg to one narrow physreg) nor a soft COPY+hint is safe: the
- // soft copies read/write overlapping physreg sub-slices and the allocator
- // clobbers one before the other is read (miscompile). But the shared load was
- // already class-constrained to the requested file above (and tryFoldLoad put
- // it in AGPR), so the pin is redundant -- make it a no-op: replace uses of the
- // pin result with the source (sub)register directly and erase the pin.
+ // A sub-register source means the value is a slice of a shared register
+ // (e.g. one ds_read2 loads two pinned fragments into one wide reg). Pinning
+ // it -- hard or soft -- would move overlapping physreg sub-slices and
+ // miscompile. The shared reg is already in the right file (above), so the pin
+ // is redundant: forward the source (sub)register to the uses and drop it.
if (Pin->getOperand(1).getSubReg()) {
unsigned SubIdx = Pin->getOperand(1).getSubReg();
for (MachineOperand &MO :
@@ -330,14 +305,10 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
bool Hard = EnableHardPin && PR && Src.isVirtual() && Dst.isVirtual();
- // Deterministic AGPR placement for a load tuple. When an AGPR pin's value is
- // a REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a
- // fixed physical AGPR sub-register so A/B are *born* in fixed AGPRs. Without
- // this the MFMA A/B operands are AV (agpr-or-vgpr) and the coalescer /
- // allocator moves them back to VGPR whenever pressure is low, making the pin
- // non-deterministic. The accumulator was already routed to VGPR (vgprcd) by
- // the file-constraint step above, so this yields v[D], a[A], a[B] with the
- // accumulator free to occupy the whole VGPR file.
+ // Deterministic AGPR placement for a load tuple: when the pinned value is a
+ // REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a fixed
+ // physical AGPR sub-register. Otherwise the MFMA A/B operands are AV and the
+ // allocator moves them back to VGPR under low pressure (non-deterministic).
if (Hard && WantAGPR) {
MachineInstr *RS = MRI.getVRegDef(Src);
MachineBasicBlock *PinMBB = Pin->getParent();
@@ -606,16 +577,12 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
fullyRecomputeLiveIns(MBBs);
}
- // Let a *VGPR* pin drive occupancy: a wide pinned VGPR value (e.g. a 192-VGPR
- // accumulator) must fit the per-wave VGPR budget, so cap occupancy to make
- // room without the user setting __launch_bounds__ / amdgpu-waves-per-eu.
- // AGPR pins must NOT drive this: AGPRs are a separate file, and feeding an
- // AGPR count into the VGPR occupancy formula wrongly raises occupancy and
- // shrinks the VGPR budget, spilling the (VGPR) accumulator into AGPRs.
+ // Cap occupancy so a wide VGPR-resident value fits the per-wave budget without
+ // the user setting __launch_bounds__. Only VGPR footprints drive this: AGPRs
+ // are a separate file, so feeding an AGPR count into the VGPR occupancy formula
+ // would wrongly raise occupancy and spill the VGPR accumulator. `AccVGPRs` is
+ // the footprint of the vgprcd accumulator tiles plus PinAccVGPRMargin.
auto *MFI = MF.getInfo<SIMachineFunctionInfo>();
- // Total VGPR footprint of the accumulator tiles routed to VGPR, plus a margin
- // for addressing/temps. When A/B are pinned to AGPR the accumulator must stay
- // VGPR-resident; this caps occupancy so its budget is large enough.
unsigned AccVGPRs = 0;
if (PinAccVGPRMargin) {
for (Register R : AccTiles)
>From ec5433dd188dd19b83c11e042028b9db4d6e60e6 Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Thu, 2 Jul 2026 15:16:02 +0000
Subject: [PATCH 04/19] [AMDGPU] Add tests for the register-pinning
intrinsics/attribute
- llvm/test/CodeGen/AMDGPU/pin-reg.ll: pin.agpr/pin.vgpr lowering (AGPR-born
loads + AGPR MFMA operands, no v_accvgpr), soft fallback
(-amdgpu-hard-pin-regs=0), a shared-load sub-slice regression, and a no-pin
self-containment case (pass is a no-op).
- clang/test/CodeGenHIP/amdgpu-pin-attr.hip: amdgpu_pin_{agpr,vgpr} attribute
lowers stores to llvm.amdgcn.pin.*; constant/template arg accepted.
- clang/test/SemaHIP/amdgpu-pin-attr.hip: attribute arg validation
(non-negative constant; rejects negative/non-constant; Var-only subject).
- mlir/test/Target/LLVMIR/amdgcn-pin.mlir: the intrinsic is reachable from MLIR
via llvm.call_intrinsic (the FlyDSL path), overload mangled from operand type.
All four pass under lit.
---
clang/test/CodeGenHIP/amdgpu-pin-attr.hip | 37 ++++++++++
clang/test/SemaHIP/amdgpu-pin-attr.hip | 25 +++++++
llvm/test/CodeGen/AMDGPU/pin-reg.ll | 90 +++++++++++++++++++++++
mlir/test/Target/LLVMIR/amdgcn-pin.mlir | 21 ++++++
4 files changed, 173 insertions(+)
create mode 100644 clang/test/CodeGenHIP/amdgpu-pin-attr.hip
create mode 100644 clang/test/SemaHIP/amdgpu-pin-attr.hip
create mode 100644 llvm/test/CodeGen/AMDGPU/pin-reg.ll
create mode 100644 mlir/test/Target/LLVMIR/amdgcn-pin.mlir
diff --git a/clang/test/CodeGenHIP/amdgpu-pin-attr.hip b/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
new file mode 100644
index 0000000000000..a64b27806aa3c
--- /dev/null
+++ b/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx950 -x hip \
+// RUN: -fcuda-is-device -emit-llvm -o - %s | FileCheck %s
+
+// The amdgpu_pin_{agpr,vgpr} attribute wraps every store to the local in the
+// corresponding llvm.amdgcn.pin.* intrinsic (chunked to i32-based widths).
+
+typedef float float2 __attribute__((ext_vector_type(2)));
+
+// CHECK-LABEL: define{{.*}}pin_agpr
+// CHECK: call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %{{[0-9]+}}, i32 8)
+__attribute__((device)) void pin_agpr(float2 *out, float2 in) {
+ __attribute__((amdgpu_pin_agpr(8))) float2 x;
+ x = in;
+ *out = x;
+}
+
+// CHECK-LABEL: define{{.*}}pin_vgpr
+// CHECK: call <2 x i32> @llvm.amdgcn.pin.vgpr.v2i32(<2 x i32> %{{[0-9]+}}, i32 4)
+__attribute__((device)) void pin_vgpr(float2 *out, float2 in) {
+ __attribute__((amdgpu_pin_vgpr(4))) float2 x;
+ x = in;
+ *out = x;
+}
+
+// A constant-expression argument (here via a template parameter) is accepted and
+// evaluated at instantiation.
+template <int N>
+__attribute__((device)) void pin_tmpl(float2 *out, float2 in) {
+ __attribute__((amdgpu_pin_agpr(N + 4))) float2 x;
+ x = in;
+ *out = x;
+}
+// CHECK-LABEL: define{{.*}}pin_tmpl
+// CHECK: call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %{{[0-9]+}}, i32 12)
+__attribute__((device)) void use_tmpl(float2 *out, float2 in) {
+ pin_tmpl<8>(out, in);
+}
diff --git a/clang/test/SemaHIP/amdgpu-pin-attr.hip b/clang/test/SemaHIP/amdgpu-pin-attr.hip
new file mode 100644
index 0000000000000..65ef7fdbc616e
--- /dev/null
+++ b/clang/test/SemaHIP/amdgpu-pin-attr.hip
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx950 -x hip \
+// RUN: -fcuda-is-device -fsyntax-only -verify %s
+
+typedef float float2 __attribute__((ext_vector_type(2)));
+
+__attribute__((device)) void ok(void) {
+ __attribute__((amdgpu_pin_agpr(0))) float2 a; // fine
+ __attribute__((amdgpu_pin_vgpr(8))) float2 b; // fine
+ constexpr int base = 16;
+ __attribute__((amdgpu_pin_agpr(base + 4))) float2 c; // constant expr: fine
+ (void)a; (void)b; (void)c;
+}
+
+__attribute__((device)) void bad(int n) { // expected-note {{declared here}}
+ // expected-error at +1 {{'amdgpu_pin_agpr' attribute requires a non-negative integral compile time constant expression}}
+ __attribute__((amdgpu_pin_agpr(-1))) float2 a;
+ // expected-error at +2 {{expression is not an integral constant expression}}
+ // expected-note at +1 {{function parameter 'n' with unknown value cannot be used in a constant expression}}
+ __attribute__((amdgpu_pin_vgpr(n))) float2 b;
+ (void)a; (void)b;
+}
+
+// The attribute only applies to variables.
+// expected-warning at +1 {{'amdgpu_pin_agpr' attribute only applies to variables}}
+__attribute__((device)) __attribute__((amdgpu_pin_agpr(0))) void func(void) {}
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg.ll b/llvm/test/CodeGen/AMDGPU/pin-reg.ll
new file mode 100644
index 0000000000000..9578a0c4c54e4
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg.ll
@@ -0,0 +1,90 @@
+; RUN: llc -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx950 -verify-machineinstrs -amdgpu-hard-pin-regs=0 < %s | FileCheck -check-prefixes=SOFT %s
+
+; Tests for the llvm.amdgcn.pin.{vgpr,agpr} register-pinning intrinsics.
+
+declare i32 @llvm.amdgcn.workitem.id.x()
+declare <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32>, i32 immarg)
+declare <2 x i32> @llvm.amdgcn.pin.vgpr.v2i32(<2 x i32>, i32 immarg)
+declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+; An AGPR pin on the A/B inputs makes the loads AGPR-born and the MFMA read AGPR
+; operands, with no agpr<->vgpr shuffle.
+; CHECK-LABEL: {{^}}pin_agpr_input:
+; CHECK: global_load_{{.*}} a[
+; CHECK: global_load_{{.*}} a[
+; CHECK-NOT: v_accvgpr
+; CHECK: v_mfma_f32_16x16x16_f16 {{[va]}}[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+; The pin is honored even with hard pinning disabled (soft allocation hint).
+; SOFT-LABEL: {{^}}pin_agpr_input:
+; SOFT: v_mfma_f32_16x16x16_f16
+define amdgpu_kernel void @pin_agpr_input(ptr addrspace(1) %pa, ptr addrspace(1) %pb, ptr addrspace(1) %pc) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %ga = getelementptr <4 x half>, ptr addrspace(1) %pa, i32 %tid
+ %gb = getelementptr <4 x half>, ptr addrspace(1) %pb, i32 %tid
+ %gc = getelementptr <4 x float>, ptr addrspace(1) %pc, i32 %tid
+ %a = load <4 x half>, ptr addrspace(1) %ga
+ %b = load <4 x half>, ptr addrspace(1) %gb
+ %ai = bitcast <4 x half> %a to <2 x i32>
+ %bi = bitcast <4 x half> %b to <2 x i32>
+ %ap = call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %ai, i32 0)
+ %bp = call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %bi, i32 8)
+ %af = bitcast <2 x i32> %ap to <4 x half>
+ %bf = bitcast <2 x i32> %bp to <4 x half>
+ %d = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %af, <4 x half> %bf, <4 x float> zeroinitializer, i32 0, i32 0, i32 0)
+ store <4 x float> %d, ptr addrspace(1) %gc
+ ret void
+}
+
+; A VGPR pin keeps its value in VGPRs (identity around the load's natural file).
+; CHECK-LABEL: {{^}}pin_vgpr_value:
+; CHECK: global_load_{{.*}} v[
+; CHECK-NOT: v_accvgpr
+define amdgpu_kernel void @pin_vgpr_value(ptr addrspace(1) %p, ptr addrspace(1) %q) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %gp = getelementptr <2 x i32>, ptr addrspace(1) %p, i32 %tid
+ %gq = getelementptr <2 x i32>, ptr addrspace(1) %q, i32 %tid
+ %v = load <2 x i32>, ptr addrspace(1) %gp
+ %vp = call <2 x i32> @llvm.amdgcn.pin.vgpr.v2i32(<2 x i32> %v, i32 4)
+ store <2 x i32> %vp, ptr addrspace(1) %gq
+ ret void
+}
+
+; Regression: two pins taking sub-slices of ONE wide load must not clobber each
+; other (a naive rewrite of the shared register miscompiled). Both halves must be
+; used; verify-machineinstrs (in the RUN line) also guards liveness.
+; CHECK-LABEL: {{^}}pin_shared_load:
+; CHECK: v_mfma_f32_16x16x16_f16
+define amdgpu_kernel void @pin_shared_load(ptr addrspace(1) %p, ptr addrspace(1) %pc) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %gp = getelementptr <4 x i32>, ptr addrspace(1) %p, i32 %tid
+ %gc = getelementptr <4 x float>, ptr addrspace(1) %pc, i32 %tid
+ %w = load <4 x i32>, ptr addrspace(1) %gp
+ %alo = shufflevector <4 x i32> %w, <4 x i32> poison, <2 x i32> <i32 0, i32 1>
+ %bhi = shufflevector <4 x i32> %w, <4 x i32> poison, <2 x i32> <i32 2, i32 3>
+ %ap = call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %alo, i32 0)
+ %bp = call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %bhi, i32 2)
+ %af = bitcast <2 x i32> %ap to <4 x half>
+ %bf = bitcast <2 x i32> %bp to <4 x half>
+ %d = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %af, <4 x half> %bf, <4 x float> zeroinitializer, i32 0, i32 0, i32 0)
+ store <4 x float> %d, ptr addrspace(1) %gc
+ ret void
+}
+
+; Self-containment: a function with NO pin intrinsic is unaffected by the pass.
+; It gets the target's default MFMA form (accumulator AGPR, inputs VGPR) with no
+; pin-introduced agpr<->vgpr shuffles.
+; CHECK-LABEL: {{^}}no_pin:
+; CHECK: v_mfma_f32_16x16x16_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[
+; CHECK-NOT: v_accvgpr
+define amdgpu_kernel void @no_pin(ptr addrspace(1) %pa, ptr addrspace(1) %pb, ptr addrspace(1) %pc) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %ga = getelementptr <4 x half>, ptr addrspace(1) %pa, i32 %tid
+ %gb = getelementptr <4 x half>, ptr addrspace(1) %pb, i32 %tid
+ %gc = getelementptr <4 x float>, ptr addrspace(1) %pc, i32 %tid
+ %a = load <4 x half>, ptr addrspace(1) %ga
+ %b = load <4 x half>, ptr addrspace(1) %gb
+ %d = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> %a, <4 x half> %b, <4 x float> zeroinitializer, i32 0, i32 0, i32 0)
+ store <4 x float> %d, ptr addrspace(1) %gc
+ ret void
+}
diff --git a/mlir/test/Target/LLVMIR/amdgcn-pin.mlir b/mlir/test/Target/LLVMIR/amdgcn-pin.mlir
new file mode 100644
index 0000000000000..ee3749d2d4e0d
--- /dev/null
+++ b/mlir/test/Target/LLVMIR/amdgcn-pin.mlir
@@ -0,0 +1,21 @@
+// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s
+
+// The register-pinning intrinsics are reachable from MLIR via llvm.call_intrinsic
+// (the mechanism DSLs such as FlyDSL use); the overload is mangled from the
+// operand type. Requires an LLVM that defines llvm.amdgcn.pin.*.
+
+// CHECK-LABEL: define <2 x i32> @pin_agpr
+llvm.func @pin_agpr(%v: vector<2xi32>) -> vector<2xi32> {
+ %r = llvm.mlir.constant(8 : i32) : i32
+ // CHECK: call <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32> %{{[0-9]+}}, i32 8)
+ %p = llvm.call_intrinsic "llvm.amdgcn.pin.agpr"(%v, %r) : (vector<2xi32>, i32) -> vector<2xi32>
+ llvm.return %p : vector<2xi32>
+}
+
+// CHECK-LABEL: define <4 x float> @pin_vgpr
+llvm.func @pin_vgpr(%v: vector<4xf32>) -> vector<4xf32> {
+ %r = llvm.mlir.constant(0 : i32) : i32
+ // CHECK: call <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float> %{{[0-9]+}}, i32 0)
+ %p = llvm.call_intrinsic "llvm.amdgcn.pin.vgpr"(%v, %r) : (vector<4xf32>, i32) -> vector<4xf32>
+ llvm.return %p : vector<4xf32>
+}
>From 95ff9429be6112ac2120d96758f753057407d76c Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Fri, 3 Jul 2026 01:49:22 +0000
Subject: [PATCH 05/19] [AMDGPU] Soft no-op register pins on targets without an
AGPR file
An amdgpu_pin_agpr / llvm.amdgcn.pin.agpr on a subtarget that has no AGPR
register file (e.g. RDNA3/RDNA4, which use WMMA and only have VGPRs) previously
reached register allocation with an AGPR-class destination and failed with
"no registers from class available to allocate".
Guard SIPreColorPins: when the subtarget has no MAI/AGPR support, an AGPR pin is
degraded to a soft no-op -- the source is forwarded to the uses and the pin is
dropped -- so the value stays in its natural VGPR location and the kernel builds
and runs correctly. pin_vgpr is unaffected and continues to place operands in the
requested VGPRs.
Verified on gfx1201 (RX 9070 XT): a WMMA kernel with pin_vgpr places A/B/D in the
requested VGPRs (v[8:11]/v[12:15]/v[20:27], loaded directly), pin_agpr soft
no-ops, and both are bit-identical to the unpinned kernel on the GPU. Adds
llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll; the existing gfx950 test is unchanged
(the guard only triggers on non-AGPR targets).
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 14 ++++++
llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll | 61 +++++++++++++++++++++++
2 files changed, 75 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 45423b9779c16..fcaeab2136804 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -163,6 +163,20 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
// occupancy target (the register budget must cover the pinned range).
unsigned NumRegs = TRI->getRegSizeInBits(*RC) / 32;
bool WantAGPR = TRI->isAGPRClass(RC);
+
+ // Targets without an AGPR file (e.g. RDNA) cannot honor an AGPR pin. Degrade
+ // to a soft no-op -- forward the source to the uses and drop the pin -- so the
+ // value stays in its natural VGPR location instead of failing register
+ // allocation with "no registers from class available".
+ if (WantAGPR && !ST.hasMAIInsts()) {
+ for (MachineOperand &MO :
+ llvm::make_early_inc_range(MRI.use_operands(Dst)))
+ MO.setReg(Src);
+ if (Src.isVirtual())
+ MRI.constrainRegClass(Src, TRI->getEquivalentVGPRClass(RC));
+ Pin->eraseFromParent();
+ continue;
+ }
if (WantAGPR)
ReqAGPRs = std::max(ReqAGPRs, RegNo + NumRegs);
else
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll
new file mode 100644
index 0000000000000..6d3bdec2fd6f9
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll
@@ -0,0 +1,61 @@
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1201 -verify-machineinstrs < %s | FileCheck %s
+
+; RDNA4 (gfx1201) has only a VGPR file and uses WMMA. pin_vgpr must place the
+; operands in the requested VGPRs; pin_agpr has no AGPR file to target and must
+; degrade to a soft no-op (compile cleanly, no AGPRs, correct WMMA) rather than
+; fail register allocation.
+
+; pin_vgpr: A -> v[8:11], B -> v[12:15], D -> v[20:27], loaded straight in.
+; CHECK-LABEL: pin_vgpr_wmma:
+; CHECK: global_load_b128 v[8:11],
+; CHECK: global_load_b128 v[12:15],
+; CHECK: v_wmma_f32_16x16x16_f16 v[20:27], v[8:11], v[12:15]
+define protected amdgpu_kernel void @pin_vgpr_wmma(ptr addrspace(1) nocapture readonly %A, ptr addrspace(1) nocapture readonly %B, ptr addrspace(1) nocapture writeonly %C) {
+entry:
+ %id = tail call i32 @llvm.amdgcn.workitem.id.x()
+ %off = zext i32 %id to i64
+ %pa = getelementptr inbounds <8 x half>, ptr addrspace(1) %A, i64 %off
+ %la = load <4 x i32>, ptr addrspace(1) %pa, align 16
+ %pina = tail call <4 x i32> @llvm.amdgcn.pin.vgpr.v4i32(<4 x i32> %la, i32 8)
+ %a = bitcast <4 x i32> %pina to <8 x half>
+ %pb = getelementptr inbounds <8 x half>, ptr addrspace(1) %B, i64 %off
+ %lb = load <4 x i32>, ptr addrspace(1) %pb, align 16
+ %pinb = tail call <4 x i32> @llvm.amdgcn.pin.vgpr.v4i32(<4 x i32> %lb, i32 12)
+ %b = bitcast <4 x i32> %pinb to <8 x half>
+ %d = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v8f32.v8f16(<8 x half> %a, <8 x half> %b, <8 x float> zeroinitializer)
+ %di = bitcast <8 x float> %d to <8 x i32>
+ %pind = tail call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> %di, i32 20)
+ %pc = getelementptr inbounds <8 x float>, ptr addrspace(1) %C, i64 %off
+ store <8 x i32> %pind, ptr addrspace(1) %pc, align 32
+ ret void
+}
+
+; pin_agpr on gfx1201: soft no-op. Compiles to a plain WMMA, no AGPRs used.
+; CHECK-LABEL: pin_agpr_noop:
+; CHECK-NOT: a[
+; CHECK: v_wmma_f32_16x16x16_f16 v[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}]
+; CHECK-NOT: a[
+; CHECK: .set pin_agpr_noop.num_agpr, 0
+define protected amdgpu_kernel void @pin_agpr_noop(ptr addrspace(1) nocapture readonly %A, ptr addrspace(1) nocapture readonly %B, ptr addrspace(1) nocapture writeonly %C) {
+entry:
+ %id = tail call i32 @llvm.amdgcn.workitem.id.x()
+ %off = zext i32 %id to i64
+ %pa = getelementptr inbounds <8 x half>, ptr addrspace(1) %A, i64 %off
+ %la = load <4 x i32>, ptr addrspace(1) %pa, align 16
+ %pina = tail call <4 x i32> @llvm.amdgcn.pin.agpr.v4i32(<4 x i32> %la, i32 0)
+ %a = bitcast <4 x i32> %pina to <8 x half>
+ %pb = getelementptr inbounds <8 x half>, ptr addrspace(1) %B, i64 %off
+ %lb = load <4 x i32>, ptr addrspace(1) %pb, align 16
+ %pinb = tail call <4 x i32> @llvm.amdgcn.pin.agpr.v4i32(<4 x i32> %lb, i32 4)
+ %b = bitcast <4 x i32> %pinb to <8 x half>
+ %d = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v8f32.v8f16(<8 x half> %a, <8 x half> %b, <8 x float> zeroinitializer)
+ %pc = getelementptr inbounds <8 x float>, ptr addrspace(1) %C, i64 %off
+ store <8 x float> %d, ptr addrspace(1) %pc, align 32
+ ret void
+}
+
+declare i32 @llvm.amdgcn.workitem.id.x()
+declare <8 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v8f32.v8f16(<8 x half>, <8 x half>, <8 x float>)
+declare <4 x i32> @llvm.amdgcn.pin.vgpr.v4i32(<4 x i32>, i32 immarg)
+declare <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32>, i32 immarg)
+declare <4 x i32> @llvm.amdgcn.pin.agpr.v4i32(<4 x i32>, i32 immarg)
>From 6dbaabd4fb4053d557ba080834e49255abbdfdc3 Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Fri, 3 Jul 2026 03:48:08 +0000
Subject: [PATCH 06/19] [AMDGPU] Review cleanup: condense comments, drop dead
code, add docs
Address a thorough review of the register-pinning change:
- Remove the experimental, default-off -amdgpu-pin-acc-vgpr-margin flag and the
dead AccTiles/AccVGPRs occupancy-margin path; the VGPR-footprint occupancy cap
(ReqVGPRs) is unchanged. Drop the now write-only ReqAGPRs.
- Fix a stale PinPseudo comment that described an EmitInstrWithCustomInserter
lowering the pseudos never had (they are lowered by SIPreColorPins).
- Condense the SIPreColorPins file-level and in-body comments and the intrinsic
doc comments to LLVM style; clang-format the touched regions.
- Clang: bail out of emitAMDGPUPinnedValue on non-AMDGCN targets (the pin
intrinsics are AMDGCN-only) instead of emitting invalid IR; simplify the
chunking helper.
- Diagnose amdgpu_pin_{vgpr,agpr} on non-automatic-local variables (globals,
static locals, parameters) with -Wignored-attributes, since CodeGen only pins
stores to an automatic variable's storage.
- Document the attributes in AttrDocs.td and clang ReleaseNotes; assert the pin
pseudo operand shape.
- Tests: add a wide-value chunking case (12 dwords -> v8i32 + v4i32) to the
CodeGen test and non-automatic-local ignored-attribute cases to the Sema test.
Verified on the rebuilt toolchain: gfx950 pin-reg.ll (CHECK + SOFT), gfx1201
pin-reg-gfx12.ll, and the clang Sema/CodeGen tests pass; on a gfx1201 GPU the
pin_vgpr, pin_agpr-no-op, and overlapping-pin kernels are bit-identical to their
unpinned equivalents.
---
clang/include/clang/Basic/Attr.td | 4 +-
clang/include/clang/Basic/AttrDocs.td | 19 +++
clang/lib/CodeGen/CGExpr.cpp | 45 +++--
clang/lib/Sema/SemaAMDGPU.cpp | 23 ++-
clang/test/CodeGenHIP/amdgpu-pin-attr.hip | 12 ++
clang/test/SemaHIP/amdgpu-pin-attr.hip | 11 ++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 11 +-
llvm/lib/Target/AMDGPU/SIInstructions.td | 13 +-
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 198 +++++++++-------------
9 files changed, 178 insertions(+), 158 deletions(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 26251a6d746d0..10bcb1fdba710 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -2528,14 +2528,14 @@ def AMDGPUNumVGPR : InheritableAttr {
def AMDGPUPinVGPR : InheritableAttr {
let Spellings = [Clang<"amdgpu_pin_vgpr", 0>];
let Args = [ExprArgument<"Reg">];
- let Documentation = [Undocumented];
+ let Documentation = [AMDGPUPinRegDocs];
let Subjects = SubjectList<[Var]>;
}
def AMDGPUPinAGPR : InheritableAttr {
let Spellings = [Clang<"amdgpu_pin_agpr", 0>];
let Args = [ExprArgument<"Reg">];
- let Documentation = [Undocumented];
+ let Documentation = [AMDGPUPinRegDocs];
let Subjects = SubjectList<[Var]>;
}
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 3b469437d21e4..a5746903ad046 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -3505,6 +3505,25 @@ An error will be given if:
}];
}
+def AMDGPUPinRegDocs : Documentation {
+ let Category = DocCatAMDGPUAttributes;
+ let Content = [{
+The ``amdgpu_pin_vgpr(N)`` and ``amdgpu_pin_agpr(N)`` attributes request that an
+automatic local variable be placed in the physical VGPR (respectively AGPR) tuple
+starting at register number ``N``. Every store to the variable is lowered through
+the ``llvm.amdgcn.pin.{vgpr,agpr}`` intrinsics, and a value wider than one 32-bit
+register occupies consecutive registers ``N``, ``N+1``, ... .
+
+This is a placement request, not a guarantee: when the requested registers are
+unavailable (for example they conflict with another pinned value that is
+simultaneously live) the allocator falls back to its normal choice, and an AGPR
+pin is ignored on targets that have no AGPR file. The attribute is only
+meaningful on automatic (block-scope, non-static) local variables and is ignored
+elsewhere. It is intended for hand-tuned kernels (e.g. controlling MFMA/WMMA
+operand placement) and requires an AMDGPU target.
+ }];
+}
+
def AMDGPUWavesPerEUDocs : Documentation {
let Category = DocCatAMDGPUAttributes;
let Content = [{
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index 92e2e9dd8279d..fdae64bb60894 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -3060,51 +3060,50 @@ llvm::Value *CodeGenFunction::emitAMDGPUPinnedValue(llvm::Value *V,
auto It = AMDGPUPinnedLocals.find(Addr);
if (It == AMDGPUPinnedLocals.end())
return V;
+ // The pin intrinsics are AMDGCN-only; ignore the attribute on other targets
+ // rather than emit invalid IR.
+ if (!getTarget().getTriple().isAMDGCN())
+ return V;
bool IsAGPR = It->second.first;
unsigned Reg = It->second.second;
llvm::Type *Ty = V->getType();
unsigned Bits = CGM.getDataLayout().getTypeSizeInBits(Ty);
if (Bits == 0 || (Bits % 32) != 0)
- return V; // only whole-dword values are pinnable
+ return V; // Only whole-dword values are pinnable.
unsigned Lanes = Bits / 32;
- llvm::Intrinsic::ID IID =
- IsAGPR ? llvm::Intrinsic::amdgcn_pin_agpr : llvm::Intrinsic::amdgcn_pin_vgpr;
+ llvm::Intrinsic::ID IID = IsAGPR ? llvm::Intrinsic::amdgcn_pin_agpr
+ : llvm::Intrinsic::amdgcn_pin_vgpr;
- // Pin patterns exist for i32-based widths of 1/2/4/8/16 dwords (i32, v2i32,
- // v4i32, v8i32, v16i32); a float base would need v1f32/v2f32, which have none
- // and crash isel. Decompose any dword count into these widths (e.g. 12 -> 8+4).
- llvm::Type *I32 = Int32Ty;
- auto *VecI = llvm::FixedVectorType::get(I32, Lanes);
- llvm::Value *Vec = Builder.CreateBitCast(V, VecI);
+ // Patterns exist only for i32 widths 1/2/4/8/16 (float bases would need
+ // v1f32/v2f32, which have no pattern and crash isel), so bitcast to <N x i32>
+ // and decompose into those widths (e.g. 12 -> 8 + 4).
+ auto *VecTy = llvm::FixedVectorType::get(Int32Ty, Lanes);
+ llvm::Value *Vec = Builder.CreateBitCast(V, VecTy);
- auto pin = [&](llvm::Value *Chunk, unsigned RegNo) -> llvm::Value * {
+ auto Pin = [&](llvm::Value *Chunk, unsigned RegNo) -> llvm::Value * {
llvm::Function *Fn = CGM.getIntrinsic(IID, {Chunk->getType()});
return Builder.CreateCall(Fn,
{Chunk, llvm::ConstantInt::get(Int32Ty, RegNo)});
};
- auto floorWidth = [](unsigned L) -> unsigned {
- if (L >= 16) return 16;
- if (L >= 8) return 8;
- if (L >= 4) return 4;
- if (L >= 2) return 2;
+ auto FloorWidth = [](unsigned L) -> unsigned {
+ for (unsigned W : {16u, 8u, 4u, 2u})
+ if (L >= W)
+ return W;
return 1;
};
- unsigned Off = 0;
- while (Off < Lanes) {
- unsigned W = floorWidth(Lanes - Off);
+ for (unsigned Off = 0; Off < Lanes;) {
+ unsigned W = FloorWidth(Lanes - Off);
if (W == 1) {
llvm::Value *Idx = llvm::ConstantInt::get(Int32Ty, Off);
llvm::Value *Elt = Builder.CreateExtractElement(Vec, Idx);
- Elt = pin(Elt, Reg + Off);
- Vec = Builder.CreateInsertElement(Vec, Elt, Idx);
+ Vec = Builder.CreateInsertElement(Vec, Pin(Elt, Reg + Off), Idx);
} else {
llvm::Value *Idx = llvm::ConstantInt::get(Int64Ty, Off);
llvm::Value *Sub = Builder.CreateExtractVector(
- llvm::FixedVectorType::get(I32, W), Vec, Idx);
- Sub = pin(Sub, Reg + Off);
- Vec = Builder.CreateInsertVector(VecI, Vec, Sub, Idx);
+ llvm::FixedVectorType::get(Int32Ty, W), Vec, Idx);
+ Vec = Builder.CreateInsertVector(VecTy, Vec, Pin(Sub, Reg + Off), Idx);
}
Off += W;
}
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 78768f4608f87..3294c1f19eda2 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -747,21 +747,36 @@ static Expr *checkPinRegArg(Sema &S, const AttributeCommonInfo &CI, Expr *E) {
void SemaAMDGPU::addAMDGPUPinVGPRAttr(Decl *D, const AttributeCommonInfo &CI,
Expr *RegExpr) {
if (Expr *E = checkPinRegArg(SemaRef, CI, RegExpr))
- D->addAttr(::new (getASTContext()) AMDGPUPinVGPRAttr(getASTContext(), CI, E));
+ D->addAttr(::new (getASTContext())
+ AMDGPUPinVGPRAttr(getASTContext(), CI, E));
}
void SemaAMDGPU::addAMDGPUPinAGPRAttr(Decl *D, const AttributeCommonInfo &CI,
Expr *RegExpr) {
if (Expr *E = checkPinRegArg(SemaRef, CI, RegExpr))
- D->addAttr(::new (getASTContext()) AMDGPUPinAGPRAttr(getASTContext(), CI, E));
+ D->addAttr(::new (getASTContext())
+ AMDGPUPinAGPRAttr(getASTContext(), CI, E));
+}
+
+// The pin is applied to stores of an automatic local (see EmitAutoVarAlloca),
+// so it is meaningless on globals, static locals, or parameters; ignore it
+// there.
+static bool isPinnableLocal(Sema &S, Decl *D, const ParsedAttr &AL) {
+ const auto *VD = dyn_cast<VarDecl>(D);
+ if (VD && VD->isLocalVarDecl() && VD->hasLocalStorage())
+ return true;
+ S.Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL;
+ return false;
}
void SemaAMDGPU::handleAMDGPUPinVGPRAttr(Decl *D, const ParsedAttr &AL) {
- addAMDGPUPinVGPRAttr(D, AL, AL.getArgAsExpr(0));
+ if (isPinnableLocal(SemaRef, D, AL))
+ addAMDGPUPinVGPRAttr(D, AL, AL.getArgAsExpr(0));
}
void SemaAMDGPU::handleAMDGPUPinAGPRAttr(Decl *D, const ParsedAttr &AL) {
- addAMDGPUPinAGPRAttr(D, AL, AL.getArgAsExpr(0));
+ if (isPinnableLocal(SemaRef, D, AL))
+ addAMDGPUPinAGPRAttr(D, AL, AL.getArgAsExpr(0));
}
static bool
diff --git a/clang/test/CodeGenHIP/amdgpu-pin-attr.hip b/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
index a64b27806aa3c..654ebe8a433e2 100644
--- a/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
+++ b/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
@@ -22,6 +22,18 @@ __attribute__((device)) void pin_vgpr(float2 *out, float2 in) {
*out = x;
}
+// A value wider than the largest pin width is decomposed into i32 chunks of
+// 16/8/4/2/1 dwords at consecutive register numbers (12 dwords -> 8 + 4).
+typedef float float12 __attribute__((ext_vector_type(12)));
+// CHECK-LABEL: define{{.*}}pin_wide
+// CHECK: call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> %{{[0-9]+}}, i32 0)
+// CHECK: call <4 x i32> @llvm.amdgcn.pin.vgpr.v4i32(<4 x i32> %{{[0-9]+}}, i32 8)
+__attribute__((device)) void pin_wide(float12 *out, float12 in) {
+ __attribute__((amdgpu_pin_vgpr(0))) float12 x;
+ x = in;
+ *out = x;
+}
+
// A constant-expression argument (here via a template parameter) is accepted and
// evaluated at instantiation.
template <int N>
diff --git a/clang/test/SemaHIP/amdgpu-pin-attr.hip b/clang/test/SemaHIP/amdgpu-pin-attr.hip
index 65ef7fdbc616e..3c37be13075ef 100644
--- a/clang/test/SemaHIP/amdgpu-pin-attr.hip
+++ b/clang/test/SemaHIP/amdgpu-pin-attr.hip
@@ -23,3 +23,14 @@ __attribute__((device)) void bad(int n) { // expected-note {{declared here}}
// The attribute only applies to variables.
// expected-warning at +1 {{'amdgpu_pin_agpr' attribute only applies to variables}}
__attribute__((device)) __attribute__((amdgpu_pin_agpr(0))) void func(void) {}
+
+// Only automatic locals are pinnable; the attribute is ignored on globals and
+// static locals (CodeGen pins stores to an automatic variable's storage).
+// expected-warning at +1 {{'amdgpu_pin_vgpr' attribute ignored}}
+__attribute__((device)) __attribute__((amdgpu_pin_vgpr(0))) float2 g_pinned;
+
+__attribute__((device)) void bad_storage(void) {
+ // expected-warning at +1 {{'amdgpu_pin_agpr' attribute ignored}}
+ static __attribute__((amdgpu_pin_agpr(0))) float2 s;
+ (void)s;
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index c57baa8611e56..d5dfba230d8c0 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2558,12 +2558,11 @@ def int_amdgcn_readfirstlane :
Intrinsic<[llvm_any_ty], [LLVMMatchType<0>],
[IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree, IntrNoCreateUndefOrPoison]>;
-// Register-pinning hint. Requests that the value operand be kept in the physical
-// VGPR (int_amdgcn_pin_vgpr) or AGPR (int_amdgcn_pin_agpr) tuple starting at the
-// number given by the second (immediate) operand. Overloaded on the value type:
-// a 32/64/128-bit value pins to 1/2/4 consecutive registers starting at that
-// number. This is a soft register-allocation hint: the allocator prefers those
-// registers when feasible and falls back under pressure. Value passed unchanged.
+// Register-pinning hint. Requests that the (unchanged) value operand be kept in
+// the physical VGPR (int_amdgcn_pin_vgpr) or AGPR (int_amdgcn_pin_agpr) tuple
+// starting at the number given by the immediate second operand. The allocator
+// prefers that placement when feasible and falls back under pressure; on a
+// target with no AGPR file an AGPR pin is a no-op.
def int_amdgcn_pin_vgpr :
Intrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty],
[IntrNoMem, IntrWillReturn, IntrNoCallback, IntrNoFree,
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index efee1706c6fc5..7487c9f9d8271 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -422,12 +422,9 @@ foreach Op = Operations in {
Op.VT, Op.RetReg, Op.Reg>;
}
-// Register-pinning hints. Lowered by EmitInstrWithCustomInserter into a COPY
-// plus a register-allocation hint requesting the numbered VGPR/AGPR tuple.
-// One pseudo per register width; the inserter derives the tuple physreg from the
-// destination register class, so it works for any width/alignment.
-// Expanded by the SIPreColorPins pass (pre-RA, still in SSA form) into either a
-// hard physical-register assignment or a soft COPY + allocation hint.
+// Register-pinning hints, one pseudo per register width. Expanded by the
+// SIPreColorPins pass (pre-RA, in SSA form) into either a hard physical-register
+// assignment or a soft COPY + allocation hint for the numbered VGPR/AGPR tuple.
class PinPseudo<RegisterClass DstRC, RegisterClass SrcRC> :
VPseudoInstSI <(outs DstRC:$vdst), (ins SrcRC:$src, i32imm:$regno), []> {
let hasSideEffects = 0;
@@ -446,7 +443,9 @@ foreach w = [32,64,96,128,160,192,224,256,288,320,352,384,512,1024] in {
def PIN_AGPR_B#w : PinPseudo<ARC, AVRC>;
}
-// Map each supported value type to the width-appropriate pseudo via its size.
+// Selection patterns for the supported value types. Clang decomposes wider or
+// non-i32 values into these i32-based widths (see emitAMDGPUPinnedValue); an
+// unlisted type reaching ISel is unsupported and reports "cannot select".
foreach vt = [i32, f32, v2i32, v4i32, v4f32, v8f16,
v8i32, v8f32, v16i32, v16f32] in {
def : GCNPat<(vt (int_amdgcn_pin_vgpr vt:$s, (i32 timm:$r))),
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index fcaeab2136804..781977c8380b7 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -7,32 +7,19 @@
//===----------------------------------------------------------------------===//
//
/// \file
-/// Lowers the PIN_{VGPR,AGPR}_B* pseudos produced from
-/// llvm.amdgcn.pin.{vgpr,agpr} into a hard register assignment ("pre-coloring").
+/// Lowers the PIN_{VGPR,AGPR}_B* pseudos (from llvm.amdgcn.pin.{vgpr,agpr})
+/// into a hard physical-register assignment ("pre-coloring"): the pinned
+/// value's def and uses are rewritten to reference the requested VGPR/AGPR
+/// tuple directly, so the allocator treats it as fixed interference and cannot
+/// override it (unlike a soft hint). The whole tie-connected component is
+/// rewritten together, so a pin on an MFMA accumulator input also pins its tied
+/// output.
///
-/// The value being pinned is rewritten so that its def and all its uses
-/// reference the requested physical VGPR/AGPR tuple directly. Because the value
-/// is then a physical register in the MIR, the register allocator treats it as
-/// fixed interference and can never place it elsewhere or let another value
-/// clobber it -- unlike the soft allocation hint, this cannot be overridden by
-/// competing coalescer copy-hints (e.g. an MFMA accumulator chain).
-///
-/// Tied operands (e.g. the in-place MFMA accumulator, whose vdst is tied to
-/// src2) require care: both ends of a tie must share the same register. The
-/// pass therefore rewrites the whole *tie-connected component* of virtual
-/// registers, so a pin placed on the accumulator input also pins the tied
-/// output. Subregister references are rewritten to the corresponding physical
-/// subregister.
-///
-/// When hard pinning is not safe (a def in the component is a PHI, REG_SEQUENCE
-/// or IMPLICIT_DEF, the physical (sub)register is not a legal member of some
-/// rewritten operand's register class, or the tuple conflicts with an already
-/// hard-pinned value) the pass falls back to the soft behaviour: a COPY plus a
-/// register-allocation hint. This guarantees the pass never regresses
-/// correctness.
-///
-/// Runs pre-RA while the function is still in SSA form (before PHIElimination /
-/// TwoAddressInstruction), so each value has a single reaching def.
+/// When hard pinning is unsafe (a PHI/REG_SEQUENCE/IMPLICIT_DEF def, a physreg
+/// illegal for some operand's class, or a tuple conflicting with an existing
+/// hard pin) the pass falls back to a COPY plus a soft allocation hint, so it
+/// never regresses correctness. Runs pre-RA in SSA form (before PHIElimination
+/// / TwoAddressInstruction), so each value has a single reaching def.
//
//===----------------------------------------------------------------------===//
@@ -64,15 +51,6 @@ static cl::opt<bool> PinAgprVgprC(
cl::desc("Convert an AGPR-input MFMA to vgprcd to keep its accumulator in "
"VGPR (else keep the native all-AGPR form)"));
-// Experimental (default off): if nonzero, an AGPR-input pin caps occupancy so
-// the vgprcd accumulator plus this many VGPRs of headroom stay resident, in
-// place of __launch_bounds__. Driving occupancy here can perturb hard-pinned
-// physreg live ranges at low occupancy, so __launch_bounds__ is preferred.
-static cl::opt<unsigned> PinAccVGPRMargin(
- "amdgpu-pin-acc-vgpr-margin", cl::init(0), cl::Hidden,
- cl::desc("If nonzero, VGPRs reserved above a vgprcd accumulator so an "
- "AGPR-input pin can drive occupancy (experimental)"));
-
namespace {
class SIPreColorPins : public MachineFunctionPass {
@@ -109,9 +87,9 @@ static bool isPinPseudo(const SIInstrInfo *TII, const MachineInstr &MI) {
return N.starts_with("PIN_VGPR_B") || N.starts_with("PIN_AGPR_B");
}
-// Physical register tuple a pin targets, or 0 if it is not a legal member of the
-// destination register class (e.g. a misaligned start on a target that requires
-// aligned tuples).
+// Physical register tuple a pin targets, or 0 if it is not a legal member of
+// the destination register class (e.g. a misaligned start on a target that
+// requires aligned tuples).
static MCRegister getPinPhysReg(const SIRegisterInfo *TRI,
const TargetRegisterClass *RC, unsigned RegNo) {
unsigned First =
@@ -139,35 +117,31 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (Pins.empty())
return false;
- // Regunits already claimed by a hard pin. A later pin whose tuple overlaps any
- // claimed unit falls back to soft, so two distinct simultaneously-live values
- // can never be forced into the same physical register. (Legitimate reuse of a
- // register by a single value -- e.g. an accumulation chain -- is absorbed by
- // the tie-connected component of the first pin, after which the value is
- // already physical and later pins on it are no-ops.)
+ // Regunits already claimed by a hard pin. A later pin overlapping any claimed
+ // unit falls back to soft, so two distinct live values never share a physreg.
+ // Reuse by a single value (e.g. an accumulation chain) is instead absorbed
+ // into the first pin's tie-connected component, making later pins on it
+ // no-ops.
DenseSet<MCRegUnit> Claimed;
bool NeedRecomputeLiveIns = false;
- unsigned ReqVGPRs = 0, ReqAGPRs = 0; // highest register a pin needs, +1
- // Accumulator tiles routed to VGPR by the vgprcd conversion; their footprint
- // optionally drives the occupancy cap (see PinAccVGPRMargin).
- DenseSet<Register> AccTiles;
-
+ unsigned ReqVGPRs =
+ 0; // highest VGPR a pin needs, +1 (drives the occupancy cap)
for (MachineInstr *Pin : Pins) {
+ assert(Pin->getNumExplicitOperands() == 3 &&
+ "pin pseudo must be (dst, src, regno)");
Register Dst = Pin->getOperand(0).getReg();
Register Src = Pin->getOperand(1).getReg();
unsigned RegNo = Pin->getOperand(2).getImm();
const TargetRegisterClass *RC = MRI.getRegClass(Dst);
MCRegister PR = getPinPhysReg(TRI, RC, RegNo);
- // Record how many registers this pin needs so the pin itself can drive the
- // occupancy target (the register budget must cover the pinned range).
unsigned NumRegs = TRI->getRegSizeInBits(*RC) / 32;
bool WantAGPR = TRI->isAGPRClass(RC);
- // Targets without an AGPR file (e.g. RDNA) cannot honor an AGPR pin. Degrade
- // to a soft no-op -- forward the source to the uses and drop the pin -- so the
- // value stays in its natural VGPR location instead of failing register
- // allocation with "no registers from class available".
+ // Targets without an AGPR file (e.g. RDNA) cannot honor an AGPR pin.
+ // Degrade to a soft no-op -- forward the source to the uses and drop the
+ // pin -- so the value stays in its natural VGPR location instead of failing
+ // register allocation with "no registers from class available".
if (WantAGPR && !ST.hasMAIInsts()) {
for (MachineOperand &MO :
llvm::make_early_inc_range(MRI.use_operands(Dst)))
@@ -177,22 +151,22 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
Pin->eraseFromParent();
continue;
}
- if (WantAGPR)
- ReqAGPRs = std::max(ReqAGPRs, RegNo + NumRegs);
- else
+ // Only VGPR pins drive the occupancy cap (see below); AGPRs are a separate
+ // file that does not affect the VGPR budget.
+ if (!WantAGPR)
ReqVGPRs = std::max(ReqVGPRs, RegNo + NumRegs);
- // Constrain the pinned value's register file (and connected vregs) to VGPR
- // or AGPR. This is a class narrowing, not a physreg pin, so it also works
- // for loop-carried PHI values; it no-ops when the file is incompatible.
+ // Narrow the pinned value's register file to VGPR or AGPR (a class
+ // narrowing, not a physreg pin, so it also works for loop-carried PHIs and
+ // no-ops when the file is incompatible).
{
- // Gather the copy/REG_SEQUENCE/tie-connected component of `Seeds` and
- // constrain each member to the requested file. MFMA src2<->vdst edges are
- // followed only when `FollowAcc`; otherwise an MFMA *using* a member as
- // src0/src1 is recorded in `Inputs` as a leaf, so an input pin does not
- // drag the loop-carried accumulator into the AGPR file. `Recompute`
- // re-derives classes from defs first (needed after an opcode conversion,
- // since constrainRegClass cannot cross the disjoint AGPR/VGPR files).
+ // Constrain the copy/REG_SEQUENCE/PHI/tie-connected component of `Seeds`.
+ // MFMA src2<->vdst edges are followed only when `FollowAcc`; otherwise an
+ // MFMA using a member as src0/src1 is recorded in `Inputs` as a leaf, so
+ // an input pin does not drag the loop-carried accumulator into the AGPR
+ // file. `Recompute` re-derives classes from defs first (needed after an
+ // opcode conversion, since constrainRegClass cannot cross the AGPR/VGPR
+ // files).
auto constrainComponent = [&](ArrayRef<Register> Seeds, bool AGPRFile,
bool FollowAcc, bool Recompute,
SmallPtrSetImpl<MachineInstr *> &Inputs) {
@@ -207,11 +181,9 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (unsigned I = 0; I < WL.size(); ++I) {
for (MachineOperand &MO : MRI.reg_operands(WL[I])) {
MachineInstr *MI = MO.getParent();
- // Copy / REG_SEQUENCE / PHI all just move the value between vregs;
- // pull every register operand into the component. PHI matters for
- // the loop-carried accumulator: without it the carried value stays
- // in its original file (AGPR) while the vgprcd MFMA computes in VGPR,
- // forcing an agpr<->vgpr copy every iteration.
+ // Copy/REG_SEQUENCE/PHI just move the value between vregs; pull in
+ // every register operand. PHI keeps a loop-carried accumulator in
+ // one file (else it needs an agpr<->vgpr copy each iteration).
if (MI->isCopy() || MI->isRegSequence() || MI->isPHI()) {
for (MachineOperand &O : MI->operands())
if (O.isReg())
@@ -243,8 +215,9 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
for (Register R : WL) {
// A constant accumulator init (e.g. clear()==0) placed in an AGPR by
- // V_ACCVGPR_WRITE can't be constrained to VGPR; rewrite it to V_MOV so
- // the constant is born in VGPR instead of copied from AGPR each launch.
+ // V_ACCVGPR_WRITE can't be constrained to VGPR; rewrite it to V_MOV
+ // so the constant is born in VGPR instead of copied from AGPR each
+ // launch.
if (!AGPRFile)
for (MachineInstr &Def :
make_early_inc_range(MRI.def_instructions(R))) {
@@ -270,11 +243,11 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
constrainComponent(Seeds, /*AGPRFile=*/WantAGPR, /*FollowAcc=*/!WantAGPR,
/*Recompute=*/false, InputMFMAs);
- // ISel picks the all-AGPR MFMA form when the function needs AGPRs. To keep
- // the accumulator in VGPR, convert each consuming MFMA to vgprcd and
- // constrain its accumulator (vdst/srcC chain) to VGPR, re-deriving classes
- // from the converted defs. The chain stays coalesced in VGPR (no chunked
- // pins, no agpr<->vgpr shuffle).
+ // ISel picks the all-AGPR MFMA form when the function needs AGPRs. To
+ // keep the accumulator in VGPR, convert each consuming MFMA to vgprcd and
+ // constrain its accumulator (vdst/srcC chain) to VGPR, re-deriving
+ // classes from the converted defs. The chain stays coalesced in VGPR (no
+ // chunked pins, no agpr<->vgpr shuffle).
if (WantAGPR && PinAgprVgprC && !InputMFMAs.empty()) {
SmallVector<Register, 8> AccSeeds;
for (MachineInstr *MI : InputMFMAs) {
@@ -282,14 +255,10 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (VOp == -1)
continue; // already vgprcd form
MI->setDesc(TII->get(VOp));
- if (MI->getOperand(0).isReg()) {
+ if (MI->getOperand(0).isReg())
AccSeeds.push_back(MI->getOperand(0).getReg());
- // Each converted MFMA's vdst is one accumulator tile now living in
- // VGPR; track distinct tiles for the occupancy cap below.
- AccTiles.insert(MI->getOperand(0).getReg());
- }
- int S2 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
- AMDGPU::OpName::src2);
+ int S2 =
+ AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::src2);
if (S2 >= 0 && MI->getOperand(S2).isReg())
AccSeeds.push_back(MI->getOperand(S2).getReg());
}
@@ -304,8 +273,9 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
// A sub-register source means the value is a slice of a shared register
// (e.g. one ds_read2 loads two pinned fragments into one wide reg). Pinning
// it -- hard or soft -- would move overlapping physreg sub-slices and
- // miscompile. The shared reg is already in the right file (above), so the pin
- // is redundant: forward the source (sub)register to the uses and drop it.
+ // miscompile. The shared reg is already in the right file (above), so the
+ // pin is redundant: forward the source (sub)register to the uses and drop
+ // it.
if (Pin->getOperand(1).getSubReg()) {
unsigned SubIdx = Pin->getOperand(1).getSubReg();
for (MachineOperand &MO :
@@ -320,9 +290,10 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
bool Hard = EnableHardPin && PR && Src.isVirtual() && Dst.isVirtual();
// Deterministic AGPR placement for a load tuple: when the pinned value is a
- // REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a fixed
- // physical AGPR sub-register. Otherwise the MFMA A/B operands are AV and the
- // allocator moves them back to VGPR under low pressure (non-deterministic).
+ // REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a
+ // fixed physical AGPR sub-register. Otherwise the MFMA A/B operands are AV
+ // and the allocator moves them back to VGPR under low pressure
+ // (non-deterministic).
if (Hard && WantAGPR) {
MachineInstr *RS = MRI.getVRegDef(Src);
MachineBasicBlock *PinMBB = Pin->getParent();
@@ -358,7 +329,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MachineOperand &MO : MRI.reg_operands(Dst)) {
if (MO.getParent() == Pin)
continue;
- MCRegister T = MO.getSubReg() ? TRI->getSubReg(PR, MO.getSubReg()) : PR;
+ MCRegister T =
+ MO.getSubReg() ? TRI->getSubReg(PR, MO.getSubReg()) : PR;
if (!LegalHere(MO, T)) {
Ok = false;
break;
@@ -395,7 +367,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (MO.getParent() != Pin)
Ops.push_back(&MO);
for (MachineOperand *MO : Ops) {
- MCRegister T = MO->getSubReg() ? TRI->getSubReg(PR, MO->getSubReg()) : PR;
+ MCRegister T =
+ MO->getSubReg() ? TRI->getSubReg(PR, MO->getSubReg()) : PR;
MO->setReg(T);
MO->setSubReg(0);
MO->setIsRenamable(false);
@@ -432,8 +405,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
// MFMA form is 3-address, so an accumulation chain is connected by
// src2->vdst def-use rather than ties; pin the whole chain as a unit.
if (TII->isMAI(*MI)) {
- int Src2 =
- AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::src2);
+ int Src2 = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
+ AMDGPU::OpName::src2);
if (Src2 >= 0) {
const MachineOperand &V2 = MI->getOperand(Src2);
const MachineOperand &VD = MI->getOperand(0);
@@ -469,8 +442,9 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
Hard = false;
break;
}
- const TargetRegisterClass *OpRC = MO.getParent()->getRegClassConstraint(
- MO.getOperandNo(), TII, TRI);
+ const TargetRegisterClass *OpRC =
+ MO.getParent()->getRegClassConstraint(MO.getOperandNo(), TII,
+ TRI);
if (OpRC && !OpRC->contains(Tgt)) {
Hard = false;
break;
@@ -507,8 +481,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
}
- // Partition operands. Non-tied *subregister uses* (e.g. the per-lane reads a
- // wide accumulator feeds into stores) are not rewritten to physical
+ // Partition operands. Non-tied *subregister uses* (e.g. the per-lane reads
+ // a wide accumulator feeds into stores) are not rewritten to physical
// subregisters -- that yields fragile physical-subreg live ranges. Instead
// they read a virtual copy-out of the whole tuple.
SmallVector<MachineOperand *, 16> DirectOps, SubUses;
@@ -591,30 +565,22 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
fullyRecomputeLiveIns(MBBs);
}
- // Cap occupancy so a wide VGPR-resident value fits the per-wave budget without
- // the user setting __launch_bounds__. Only VGPR footprints drive this: AGPRs
- // are a separate file, so feeding an AGPR count into the VGPR occupancy formula
- // would wrongly raise occupancy and spill the VGPR accumulator. `AccVGPRs` is
- // the footprint of the vgprcd accumulator tiles plus PinAccVGPRMargin.
+ // Cap occupancy so a wide VGPR-resident pinned value fits the per-wave budget
+ // without the user setting __launch_bounds__. Only VGPR footprints drive
+ // this: AGPRs are a separate file, so feeding an AGPR count into the VGPR
+ // occupancy formula would wrongly raise occupancy and spill the VGPR
+ // accumulator.
auto *MFI = MF.getInfo<SIMachineFunctionInfo>();
- unsigned AccVGPRs = 0;
- if (PinAccVGPRMargin) {
- for (Register R : AccTiles)
- if (R.isVirtual())
- AccVGPRs += TRI->getRegSizeInBits(*MRI.getRegClass(R)) / 32;
- if (AccVGPRs)
- AccVGPRs += PinAccVGPRMargin;
- }
- unsigned Req = std::max(ReqVGPRs, AccVGPRs);
- if (Req) {
+ if (unsigned Req = ReqVGPRs) {
// Occupancy achievable while reserving `Req` registers per wave; cap the
// waves-per-EU (and hence the RA's VGPR budget) so the pinned range fits.
unsigned Occ = ST.getOccupancyWithNumVGPRs(Req);
auto WPE = MFI->getWavesPerEU();
unsigned NewMax = WPE.second ? std::min(WPE.second, Occ) : Occ;
// Only cap the *max* occupancy; keep the min low (1 unless the function
- // already required more). Forcing min==max over-constrains the allocator and
- // breaks physreg liveness for hard-pinned loop-body tuples at low occupancy.
+ // already required more). Forcing min==max over-constrains the allocator
+ // and breaks physreg liveness for hard-pinned loop-body tuples at low
+ // occupancy.
unsigned NewMin = std::min(WPE.first ? WPE.first : 1u, NewMax);
MFI->setWavesPerEU(NewMin, NewMax);
MFI->limitOccupancy(NewMax);
>From 2dbe55f2861e949a94f76fdf941884d64430701e Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Fri, 3 Jul 2026 05:29:10 +0000
Subject: [PATCH 07/19] [AMDGPU] Fix crash pinning a wide AGPR load-tuple of
subregister slices
The hard-pin load-tuple fast path in SIPreColorPins retargets each REG_SEQUENCE
element's def to a fixed physical AGPR sub-register. It assumed every element is
defined directly by a memory load, but for a wide value assembled from
subregister slices of wider loads (e.g. an 8-dword fp8/fp4 MFMA operand built
from two dwordx4 loads split into dword lanes) the elements are copies, not
loads. Retargeting those produced malformed physreg liveness and crashed the
backend with "Use not jointly dominated by defs".
Require each element to be load-defined; otherwise bail so the general path falls
back to soft (which folds the loads into AGPRs and places them correctly anyway).
Verified on gfx950: the fp8 (2-dword) hard pin is unchanged; an fp4/fp8 scaled
f8f6f4 MFMA with 8-dword AGPR-pinned inputs now compiles verify-clean and still
emits v_mfma_f32_16x16x128_f8f6f4 v[C], a[A], a[B] with the inputs in the
requested AGPRs. Adds the pin_agpr_wide case to pin-reg.ll.
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 13 ++++++++++--
llvm/test/CodeGen/AMDGPU/pin-reg.ll | 24 +++++++++++++++++++++++
2 files changed, 35 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 781977c8380b7..1e1aa9f10caad 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -302,13 +302,22 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (Ok && Claimed.contains(U))
Ok = false;
- // Collect element (reg, subreg-index) pairs.
+ // Collect element (reg, subreg-index) pairs. Each element must be defined
+ // directly by a memory load: this path retargets those load defs to fixed
+ // physical AGPR sub-registers. If an element is instead a subregister copy
+ // of a wider load (e.g. a dwordx4 load split into dword lanes), retargeting
+ // it produces malformed physreg liveness, so bail and let the general path
+ // fall back to soft.
SmallVector<std::pair<Register, unsigned>, 16> Elems;
if (Ok)
for (unsigned I = 1; I + 1 < RS->getNumOperands(); I += 2) {
const MachineOperand &Reg = RS->getOperand(I);
const MachineOperand &Sub = RS->getOperand(I + 1);
- if (!Reg.isReg() || !Reg.getReg().isVirtual() || Reg.getSubReg() ||
+ MachineInstr *ElemDef =
+ Reg.isReg() && Reg.getReg().isVirtual()
+ ? MRI.getVRegDef(Reg.getReg())
+ : nullptr;
+ if (!ElemDef || !ElemDef->mayLoad() || Reg.getSubReg() ||
!Sub.isImm() || !TRI->getSubReg(PR, Sub.getImm())) {
Ok = false;
break;
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg.ll b/llvm/test/CodeGen/AMDGPU/pin-reg.ll
index 9578a0c4c54e4..32179bdc1a83d 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg.ll
@@ -7,6 +7,8 @@ declare i32 @llvm.amdgcn.workitem.id.x()
declare <2 x i32> @llvm.amdgcn.pin.agpr.v2i32(<2 x i32>, i32 immarg)
declare <2 x i32> @llvm.amdgcn.pin.vgpr.v2i32(<2 x i32>, i32 immarg)
declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+declare <8 x i32> @llvm.amdgcn.pin.agpr.v8i32(<8 x i32>, i32 immarg)
+declare <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32>, <8 x i32>, <4 x float>, i32 immarg, i32 immarg, i32 immarg, i32, i32 immarg, i32)
; An AGPR pin on the A/B inputs makes the loads AGPR-born and the MFMA read AGPR
; operands, with no agpr<->vgpr shuffle.
@@ -71,6 +73,28 @@ define amdgpu_kernel void @pin_shared_load(ptr addrspace(1) %p, ptr addrspace(1)
ret void
}
+; A wide (8-dword) AGPR pin whose value is a REG_SEQUENCE of subregister slices
+; of wider loads must not crash: the hard-pin load-tuple fast path bails and the
+; pass falls back to soft, still placing the inputs in AGPRs (checked here via
+; the scaled f8f6f4 MFMA, whose fp8/fp4 A/B are eight dwords). verify-machineinstrs
+; in the RUN line guards against malformed liveness.
+; CHECK-LABEL: {{^}}pin_agpr_wide:
+; CHECK: global_load_{{.*}} a[
+; CHECK: v_mfma_f32_16x16x128_f8f6f4 v[{{[0-9:]+}}], a[{{[0-9:]+}}], a[
+define amdgpu_kernel void @pin_agpr_wide(ptr addrspace(1) %pa, ptr addrspace(1) %pb, ptr addrspace(1) %pc) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %ga = getelementptr <8 x i32>, ptr addrspace(1) %pa, i32 %tid
+ %gb = getelementptr <8 x i32>, ptr addrspace(1) %pb, i32 %tid
+ %gc = getelementptr <4 x float>, ptr addrspace(1) %pc, i32 %tid
+ %a = load <8 x i32>, ptr addrspace(1) %ga
+ %b = load <8 x i32>, ptr addrspace(1) %gb
+ %ap = call <8 x i32> @llvm.amdgcn.pin.agpr.v8i32(<8 x i32> %a, i32 0)
+ %bp = call <8 x i32> @llvm.amdgcn.pin.agpr.v8i32(<8 x i32> %b, i32 8)
+ %d = call <4 x float> @llvm.amdgcn.mfma.scale.f32.16x16x128.f8f6f4.v8i32.v8i32(<8 x i32> %ap, <8 x i32> %bp, <4 x float> zeroinitializer, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0)
+ store <4 x float> %d, ptr addrspace(1) %gc
+ ret void
+}
+
; Self-containment: a function with NO pin intrinsic is unaffected by the pass.
; It gets the target's default MFMA form (accumulator AGPR, inputs VGPR) with no
; pin-introduced agpr<->vgpr shuffles.
>From d06e39cc3743a83046def6429bdbf8bd103ea29e Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang at amd.com>
Date: Fri, 3 Jul 2026 14:20:26 +0000
Subject: [PATCH 08/19] [AMDGPU] Fix block_v2 AGPR-pin regression; narrow the
F8F6F4 crash guard
The earlier fp4 crash fix guarded the load-tuple hard-pin by requiring each
REG_SEQUENCE element to be defined directly by a load (mayLoad). That was too
broad: a real tiled MFMA kernel (opus block_v2, 192x128) assembles its A/B
fragments from buffer_load_dwordx2 pieces via subregister copies, so the guard
bailed it to soft and the A/B inputs stayed in VGPR.
Revert to the plain element check so those load tuples hard-pin again, and guard
the actual offender narrowly: the scaled f8f6f4 MFMA (V_MFMA_*_F8F6F4) consuming
a wide AGPR tuple hits a machine-scheduler "Use not jointly dominated by defs"
error under the direct physical rewrite. Detect it by walking the pinned value's
uses (through copy/reg_sequence/subreg ops) and, if one is reached, leave that
pin to the soft path (which still places the inputs in AGPRs).
Verified on gfx950 (MI355) and gfx942: opus block_v2 192x128 pins A/B into AGPR
(24/24 v_mfma_f32_16x16x16_f16 v[C], a[A], a[B], 10 buffer_load a[..], 0
v_accvgpr) and runs valid (nrms 2.4e-4); fp4/f8f6f4 no longer crashes and still
places inputs in AGPRs via the soft path; fp8 hard-pin and pin-reg.ll /
pin-reg-gfx12.ll lit tests unchanged.
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 27 ++++++++++++++++++-----
1 file changed, 22 insertions(+), 5 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 1e1aa9f10caad..982a2300fc7e5 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -298,6 +298,27 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
MachineInstr *RS = MRI.getVRegDef(Src);
MachineBasicBlock *PinMBB = Pin->getParent();
bool Ok = RS && RS->isRegSequence() && RS->getParent() == PinMBB;
+ // A scaled MFMA (mfma_scale_*, f8f6f4) consuming a wide AGPR tuple hits a
+ // machine-scheduler liveness error under the direct physical rewrite; leave
+ // those to the soft path (which still places the inputs in AGPRs). Walk the
+ // pinned value's uses (through copy/reg_sequence/subreg ops) for one.
+ if (Ok) {
+ SmallVector<Register, 8> WL{Dst};
+ DenseSet<Register> WSeen{Dst};
+ for (unsigned I = 0; I < WL.size() && Ok; ++I)
+ for (MachineInstr &U : MRI.use_nodbg_instructions(WL[I])) {
+ if (TII->getName(U.getOpcode()).contains("F8F6F4")) {
+ Ok = false;
+ break;
+ }
+ if (U.isCopy() || U.isRegSequence() || U.isPHI() ||
+ U.getOpcode() == TargetOpcode::INSERT_SUBREG ||
+ U.getOpcode() == TargetOpcode::EXTRACT_SUBREG)
+ for (const MachineOperand &D : U.defs())
+ if (D.getReg().isVirtual() && WSeen.insert(D.getReg()).second)
+ WL.push_back(D.getReg());
+ }
+ }
for (MCRegUnit U : TRI->regunits(PR))
if (Ok && Claimed.contains(U))
Ok = false;
@@ -313,11 +334,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (unsigned I = 1; I + 1 < RS->getNumOperands(); I += 2) {
const MachineOperand &Reg = RS->getOperand(I);
const MachineOperand &Sub = RS->getOperand(I + 1);
- MachineInstr *ElemDef =
- Reg.isReg() && Reg.getReg().isVirtual()
- ? MRI.getVRegDef(Reg.getReg())
- : nullptr;
- if (!ElemDef || !ElemDef->mayLoad() || Reg.getSubReg() ||
+ if (!Reg.isReg() || !Reg.getReg().isVirtual() || Reg.getSubReg() ||
!Sub.isImm() || !TRI->getSubReg(PR, Sub.getImm())) {
Ok = false;
break;
>From 67a147b832de9658eee70de6381c84019e038728 Mon Sep 17 00:00:00 2001
From: carhuang <carhuang at amd.com>
Date: Mon, 6 Jul 2026 05:56:07 +0000
Subject: [PATCH 09/19] [AMDGPU] Adapt pin-reg to amd-staging: 2-arg
getOccupancyWithNumVGPRs (dynamic VGPR block size); match InstrMapping
signature (int32_t/uint32_t) for getMFMASrcCVDstVGPROp
---
llvm/lib/Target/AMDGPU/SIInstrInfo.h | 2 +-
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 3 ++-
2 files changed, 3 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 9d2f58fcc0371..988f8d2de2baf 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -1906,7 +1906,7 @@ namespace AMDGPU {
/// \returns the VGPR (vgprcd) form of an MFMA that uses AGPRs for srcC/vdst,
/// or -1. Lets an accumulator be pinned into VGPRs with AGPR inputs.
LLVM_READONLY
- int getMFMASrcCVDstVGPROp(uint16_t Opcode);
+ int32_t getMFMASrcCVDstVGPROp(uint32_t Opcode);
/// \returns v_cmpx version of a v_cmp instruction.
LLVM_READONLY
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 982a2300fc7e5..697728cddd015 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -600,7 +600,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (unsigned Req = ReqVGPRs) {
// Occupancy achievable while reserving `Req` registers per wave; cap the
// waves-per-EU (and hence the RA's VGPR budget) so the pinned range fits.
- unsigned Occ = ST.getOccupancyWithNumVGPRs(Req);
+ unsigned Occ =
+ ST.getOccupancyWithNumVGPRs(Req, MFI->getDynamicVGPRBlockSize());
auto WPE = MFI->getWavesPerEU();
unsigned NewMax = WPE.second ? std::min(WPE.second, Occ) : Occ;
// Only cap the *max* occupancy; keep the min low (1 unless the function
>From 99e9f43bcd62cd71a3357e3dfc66b0240962726a Mon Sep 17 00:00:00 2001
From: carhuang <carhuang at amd.com>
Date: Mon, 6 Jul 2026 06:05:28 +0000
Subject: [PATCH 10/19] [AMDGPU] Update pin-reg test expectations for
amd-staging codegen (VGPR-default MFMA accumulator; .L-prefixed resource
symbols)
---
llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll | 2 +-
llvm/test/CodeGen/AMDGPU/pin-reg.ll | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll
index 6d3bdec2fd6f9..1ba10642676ba 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx12.ll
@@ -35,7 +35,7 @@ entry:
; CHECK-NOT: a[
; CHECK: v_wmma_f32_16x16x16_f16 v[{{[0-9:]+}}], v[{{[0-9:]+}}], v[{{[0-9:]+}}]
; CHECK-NOT: a[
-; CHECK: .set pin_agpr_noop.num_agpr, 0
+; CHECK: .set {{\.?L?}}pin_agpr_noop.num_agpr, 0
define protected amdgpu_kernel void @pin_agpr_noop(ptr addrspace(1) nocapture readonly %A, ptr addrspace(1) nocapture readonly %B, ptr addrspace(1) nocapture writeonly %C) {
entry:
%id = tail call i32 @llvm.amdgcn.workitem.id.x()
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg.ll b/llvm/test/CodeGen/AMDGPU/pin-reg.ll
index 32179bdc1a83d..48dba0f1e94e9 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg.ll
@@ -99,7 +99,7 @@ define amdgpu_kernel void @pin_agpr_wide(ptr addrspace(1) %pa, ptr addrspace(1)
; It gets the target's default MFMA form (accumulator AGPR, inputs VGPR) with no
; pin-introduced agpr<->vgpr shuffles.
; CHECK-LABEL: {{^}}no_pin:
-; CHECK: v_mfma_f32_16x16x16_f16 a[{{[0-9:]+}}], v[{{[0-9:]+}}], v[
+; CHECK: v_mfma_f32_16x16x16_f16 v[{{[0-9:]+}}], v[{{[0-9:]+}}], v[
; CHECK-NOT: v_accvgpr
define amdgpu_kernel void @no_pin(ptr addrspace(1) %pa, ptr addrspace(1) %pb, ptr addrspace(1) %pc) {
%tid = call i32 @llvm.amdgcn.workitem.id.x()
>From 59755f0dd68de3ca8225e9273f9d8bce50ddd5b0 Mon Sep 17 00:00:00 2001
From: carhuang <carhuang at amd.com>
Date: Mon, 6 Jul 2026 06:16:32 +0000
Subject: [PATCH 11/19] [AMDGPU] gfx1250: enable VGPR>256 register pinning
G1: add v256-v1023 to clang GCCRegNames so register numbers >255 are
accepted for amdgpu_pin_{vgpr,agpr}(N) and inline asm.
G2: opt-in -amdgpu-enable-high-vgpr-inline-asm to bind v256-v1023 in
inline-asm 'v' physreg constraints on 1024-addressable-VGPR targets.
---
clang/lib/Basic/Targets/AMDGPU.cpp | 89 ++++++++++++++++++++++-
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 17 ++++-
2 files changed, 104 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index c15ac0462918d..b6d20faba89cc 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -105,7 +105,94 @@ const char *const AMDGPUTargetInfo::GCCRegNames[] = {
"v225", "v226", "v227", "v228", "v229", "v230", "v231", "v232", "v233",
"v234", "v235", "v236", "v237", "v238", "v239", "v240", "v241", "v242",
"v243", "v244", "v245", "v246", "v247", "v248", "v249", "v250", "v251",
- "v252", "v253", "v254", "v255", "s0", "s1", "s2", "s3", "s4",
+ "v252", "v253", "v254", "v255",
+ "v256", "v257", "v258", "v259", "v260", "v261", "v262", "v263", "v264",
+ "v265", "v266", "v267", "v268", "v269", "v270", "v271", "v272", "v273",
+ "v274", "v275", "v276", "v277", "v278", "v279", "v280", "v281", "v282",
+ "v283", "v284", "v285", "v286", "v287", "v288", "v289", "v290", "v291",
+ "v292", "v293", "v294", "v295", "v296", "v297", "v298", "v299", "v300",
+ "v301", "v302", "v303", "v304", "v305", "v306", "v307", "v308", "v309",
+ "v310", "v311", "v312", "v313", "v314", "v315", "v316", "v317", "v318",
+ "v319", "v320", "v321", "v322", "v323", "v324", "v325", "v326", "v327",
+ "v328", "v329", "v330", "v331", "v332", "v333", "v334", "v335", "v336",
+ "v337", "v338", "v339", "v340", "v341", "v342", "v343", "v344", "v345",
+ "v346", "v347", "v348", "v349", "v350", "v351", "v352", "v353", "v354",
+ "v355", "v356", "v357", "v358", "v359", "v360", "v361", "v362", "v363",
+ "v364", "v365", "v366", "v367", "v368", "v369", "v370", "v371", "v372",
+ "v373", "v374", "v375", "v376", "v377", "v378", "v379", "v380", "v381",
+ "v382", "v383", "v384", "v385", "v386", "v387", "v388", "v389", "v390",
+ "v391", "v392", "v393", "v394", "v395", "v396", "v397", "v398", "v399",
+ "v400", "v401", "v402", "v403", "v404", "v405", "v406", "v407", "v408",
+ "v409", "v410", "v411", "v412", "v413", "v414", "v415", "v416", "v417",
+ "v418", "v419", "v420", "v421", "v422", "v423", "v424", "v425", "v426",
+ "v427", "v428", "v429", "v430", "v431", "v432", "v433", "v434", "v435",
+ "v436", "v437", "v438", "v439", "v440", "v441", "v442", "v443", "v444",
+ "v445", "v446", "v447", "v448", "v449", "v450", "v451", "v452", "v453",
+ "v454", "v455", "v456", "v457", "v458", "v459", "v460", "v461", "v462",
+ "v463", "v464", "v465", "v466", "v467", "v468", "v469", "v470", "v471",
+ "v472", "v473", "v474", "v475", "v476", "v477", "v478", "v479", "v480",
+ "v481", "v482", "v483", "v484", "v485", "v486", "v487", "v488", "v489",
+ "v490", "v491", "v492", "v493", "v494", "v495", "v496", "v497", "v498",
+ "v499", "v500", "v501", "v502", "v503", "v504", "v505", "v506", "v507",
+ "v508", "v509", "v510", "v511", "v512", "v513", "v514", "v515", "v516",
+ "v517", "v518", "v519", "v520", "v521", "v522", "v523", "v524", "v525",
+ "v526", "v527", "v528", "v529", "v530", "v531", "v532", "v533", "v534",
+ "v535", "v536", "v537", "v538", "v539", "v540", "v541", "v542", "v543",
+ "v544", "v545", "v546", "v547", "v548", "v549", "v550", "v551", "v552",
+ "v553", "v554", "v555", "v556", "v557", "v558", "v559", "v560", "v561",
+ "v562", "v563", "v564", "v565", "v566", "v567", "v568", "v569", "v570",
+ "v571", "v572", "v573", "v574", "v575", "v576", "v577", "v578", "v579",
+ "v580", "v581", "v582", "v583", "v584", "v585", "v586", "v587", "v588",
+ "v589", "v590", "v591", "v592", "v593", "v594", "v595", "v596", "v597",
+ "v598", "v599", "v600", "v601", "v602", "v603", "v604", "v605", "v606",
+ "v607", "v608", "v609", "v610", "v611", "v612", "v613", "v614", "v615",
+ "v616", "v617", "v618", "v619", "v620", "v621", "v622", "v623", "v624",
+ "v625", "v626", "v627", "v628", "v629", "v630", "v631", "v632", "v633",
+ "v634", "v635", "v636", "v637", "v638", "v639", "v640", "v641", "v642",
+ "v643", "v644", "v645", "v646", "v647", "v648", "v649", "v650", "v651",
+ "v652", "v653", "v654", "v655", "v656", "v657", "v658", "v659", "v660",
+ "v661", "v662", "v663", "v664", "v665", "v666", "v667", "v668", "v669",
+ "v670", "v671", "v672", "v673", "v674", "v675", "v676", "v677", "v678",
+ "v679", "v680", "v681", "v682", "v683", "v684", "v685", "v686", "v687",
+ "v688", "v689", "v690", "v691", "v692", "v693", "v694", "v695", "v696",
+ "v697", "v698", "v699", "v700", "v701", "v702", "v703", "v704", "v705",
+ "v706", "v707", "v708", "v709", "v710", "v711", "v712", "v713", "v714",
+ "v715", "v716", "v717", "v718", "v719", "v720", "v721", "v722", "v723",
+ "v724", "v725", "v726", "v727", "v728", "v729", "v730", "v731", "v732",
+ "v733", "v734", "v735", "v736", "v737", "v738", "v739", "v740", "v741",
+ "v742", "v743", "v744", "v745", "v746", "v747", "v748", "v749", "v750",
+ "v751", "v752", "v753", "v754", "v755", "v756", "v757", "v758", "v759",
+ "v760", "v761", "v762", "v763", "v764", "v765", "v766", "v767", "v768",
+ "v769", "v770", "v771", "v772", "v773", "v774", "v775", "v776", "v777",
+ "v778", "v779", "v780", "v781", "v782", "v783", "v784", "v785", "v786",
+ "v787", "v788", "v789", "v790", "v791", "v792", "v793", "v794", "v795",
+ "v796", "v797", "v798", "v799", "v800", "v801", "v802", "v803", "v804",
+ "v805", "v806", "v807", "v808", "v809", "v810", "v811", "v812", "v813",
+ "v814", "v815", "v816", "v817", "v818", "v819", "v820", "v821", "v822",
+ "v823", "v824", "v825", "v826", "v827", "v828", "v829", "v830", "v831",
+ "v832", "v833", "v834", "v835", "v836", "v837", "v838", "v839", "v840",
+ "v841", "v842", "v843", "v844", "v845", "v846", "v847", "v848", "v849",
+ "v850", "v851", "v852", "v853", "v854", "v855", "v856", "v857", "v858",
+ "v859", "v860", "v861", "v862", "v863", "v864", "v865", "v866", "v867",
+ "v868", "v869", "v870", "v871", "v872", "v873", "v874", "v875", "v876",
+ "v877", "v878", "v879", "v880", "v881", "v882", "v883", "v884", "v885",
+ "v886", "v887", "v888", "v889", "v890", "v891", "v892", "v893", "v894",
+ "v895", "v896", "v897", "v898", "v899", "v900", "v901", "v902", "v903",
+ "v904", "v905", "v906", "v907", "v908", "v909", "v910", "v911", "v912",
+ "v913", "v914", "v915", "v916", "v917", "v918", "v919", "v920", "v921",
+ "v922", "v923", "v924", "v925", "v926", "v927", "v928", "v929", "v930",
+ "v931", "v932", "v933", "v934", "v935", "v936", "v937", "v938", "v939",
+ "v940", "v941", "v942", "v943", "v944", "v945", "v946", "v947", "v948",
+ "v949", "v950", "v951", "v952", "v953", "v954", "v955", "v956", "v957",
+ "v958", "v959", "v960", "v961", "v962", "v963", "v964", "v965", "v966",
+ "v967", "v968", "v969", "v970", "v971", "v972", "v973", "v974", "v975",
+ "v976", "v977", "v978", "v979", "v980", "v981", "v982", "v983", "v984",
+ "v985", "v986", "v987", "v988", "v989", "v990", "v991", "v992", "v993",
+ "v994", "v995", "v996", "v997", "v998", "v999", "v1000", "v1001", "v1002",
+ "v1003", "v1004", "v1005", "v1006", "v1007", "v1008", "v1009", "v1010", "v1011",
+ "v1012", "v1013", "v1014", "v1015", "v1016", "v1017", "v1018", "v1019", "v1020",
+ "v1021", "v1022", "v1023",
+ "s0", "s1", "s2", "s3", "s4",
"s5", "s6", "s7", "s8", "s9", "s10", "s11", "s12", "s13",
"s14", "s15", "s16", "s17", "s18", "s19", "s20", "s21", "s22",
"s23", "s24", "s25", "s26", "s27", "s28", "s29", "s30", "s31",
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 1d165910adeba..a1b771f13613f 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -73,6 +73,19 @@ static DenormalFPEnv getDenormalFPEnv(const MachineFunction &MF) {
return MF.getInfo<SIMachineFunctionInfo>()->getMode().getDenormalFPEnv();
}
+// On gfx1250 the architectural VGPR window is v0-v255; VGPRs v256-v1023 are
+// reachable only via the S_SET_VGPR_MSB addressing mode, which the
+// AMDGPULowerVGPREncoding pass cannot manage across an opaque inline-asm body.
+// Naming v256+ in an inline-asm physical-register constraint is therefore
+// off by default (and treated as out-of-bounds, matching non-gfx1250 targets).
+// This opt-in flag lets a caller bind v256+ as a cross-boundary asm operand
+// when they take responsibility for any in-asm S_SET_VGPR_MSB themselves.
+static cl::opt<bool> EnableHighVGPRInlineAsm(
+ "amdgpu-enable-high-vgpr-inline-asm", cl::Hidden,
+ cl::desc("Allow naming VGPRs v256-v1023 in inline-asm constraints on "
+ "targets with 1024 addressable VGPRs (gfx1250+)"),
+ cl::init(false));
+
static bool denormalModeIsFlushAllF32(const MachineFunction &MF) {
const SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
return Info->getMode().FP32Denormals == DenormalMode::getPreserveSign();
@@ -19697,7 +19710,9 @@ SITargetLowering::getRegForInlineAsmConstraint(const TargetRegisterInfo *TRI_,
auto [Kind, Idx, NumRegs] = AMDGPU::parseAsmConstraintPhysReg(Constraint);
if (Kind != '\0') {
if (Kind == 'v') {
- RC = &AMDGPU::VGPR_32_Lo256RegClass;
+ RC = (Subtarget->has1024AddressableVGPRs() && EnableHighVGPRInlineAsm)
+ ? &AMDGPU::VGPR_32RegClass
+ : &AMDGPU::VGPR_32_Lo256RegClass;
} else if (Kind == 's') {
RC = &AMDGPU::SGPR_32RegClass;
} else if (Kind == 'a') {
>From 87daa43af4e906c199ed324764516784f26ad76e Mon Sep 17 00:00:00 2001
From: carhuang <carhuang at amd.com>
Date: Mon, 6 Jul 2026 06:18:30 +0000
Subject: [PATCH 12/19] [AMDGPU] Add gfx1250 VGPR>256 pin smoke test (pin to
v300, S_SET_VGPR_MSB, num_vgpr=304)
---
llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll | 19 +++++++++++++++++++
1 file changed, 19 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
new file mode 100644
index 0000000000000..f6cd994e99a1f
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
@@ -0,0 +1,19 @@
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1250 -verify-machineinstrs < %s | FileCheck -check-prefixes=CHECK %s
+
+; gfx1250 has 1024 addressable VGPRs. A pin to a VGPR index >= 256 must be
+; honored: the value is placed in the requested high tuple (reachable via the
+; S_SET_VGPR_MSB addressing mode) and the VGPR count / occupancy cap grows to
+; cover the pinned range.
+
+declare <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float>, i32 immarg)
+
+; CHECK-LABEL: {{^}}pin_high_vgpr:
+; CHECK: s_set_vgpr_msb
+; CHECK: v[{{[0-9:]+}}] /*v[300:303]*/
+; CHECK: .set .Lpin_high_vgpr.num_vgpr, 304
+define amdgpu_kernel void @pin_high_vgpr(ptr addrspace(1) %p) {
+ %v = load <4 x float>, ptr addrspace(1) %p
+ %pv = call <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float> %v, i32 300)
+ store <4 x float> %pv, ptr addrspace(1) %p
+ ret void
+}
>From 3d86996cbe71d7ae42ccda6c1fea3eabf58f691f Mon Sep 17 00:00:00 2001
From: carhuang <carhuang at amd.com>
Date: Mon, 6 Jul 2026 06:30:54 +0000
Subject: [PATCH 13/19] [AMDGPU] Gate v256-v1023 register names to gfx1250
(1024 addressable VGPRs)
The v256-v1023 GCC register names (G1) were unconditionally visible on all
AMDGPU targets, so inline asm on gfx942/etc (256 VGPRs) wrongly accepted names
like v300. Move the high VGPR names to the end of GCCRegNames and drop them in
getGCCRegNames() unless the target has 1024 addressable VGPRs (gfx1250). Add a
Sema test asserting v300 is accepted on gfx1250 and rejected on gfx942.
---
clang/lib/Basic/Targets/AMDGPU.cpp | 108 +++++++++++----------
clang/lib/Basic/Targets/AMDGPU.h | 7 ++
clang/test/Sema/amdgpu-high-vgpr-regname.c | 15 +++
3 files changed, 81 insertions(+), 49 deletions(-)
create mode 100644 clang/test/Sema/amdgpu-high-vgpr-regname.c
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index b6d20faba89cc..783bfac5859e5 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -105,7 +105,54 @@ const char *const AMDGPUTargetInfo::GCCRegNames[] = {
"v225", "v226", "v227", "v228", "v229", "v230", "v231", "v232", "v233",
"v234", "v235", "v236", "v237", "v238", "v239", "v240", "v241", "v242",
"v243", "v244", "v245", "v246", "v247", "v248", "v249", "v250", "v251",
- "v252", "v253", "v254", "v255",
+ "v252", "v253", "v254", "v255", "s0", "s1", "s2", "s3", "s4",
+ "s5", "s6", "s7", "s8", "s9", "s10", "s11", "s12", "s13",
+ "s14", "s15", "s16", "s17", "s18", "s19", "s20", "s21", "s22",
+ "s23", "s24", "s25", "s26", "s27", "s28", "s29", "s30", "s31",
+ "s32", "s33", "s34", "s35", "s36", "s37", "s38", "s39", "s40",
+ "s41", "s42", "s43", "s44", "s45", "s46", "s47", "s48", "s49",
+ "s50", "s51", "s52", "s53", "s54", "s55", "s56", "s57", "s58",
+ "s59", "s60", "s61", "s62", "s63", "s64", "s65", "s66", "s67",
+ "s68", "s69", "s70", "s71", "s72", "s73", "s74", "s75", "s76",
+ "s77", "s78", "s79", "s80", "s81", "s82", "s83", "s84", "s85",
+ "s86", "s87", "s88", "s89", "s90", "s91", "s92", "s93", "s94",
+ "s95", "s96", "s97", "s98", "s99", "s100", "s101", "s102", "s103",
+ "s104", "s105", "s106", "s107", "s108", "s109", "s110", "s111", "s112",
+ "s113", "s114", "s115", "s116", "s117", "s118", "s119", "s120", "s121",
+ "s122", "s123", "s124", "s125", "s126", "s127", "exec", "vcc", "scc",
+ "m0", "flat_scratch", "exec_lo", "exec_hi", "vcc_lo", "vcc_hi",
+ "flat_scratch_lo", "flat_scratch_hi",
+ "a0", "a1", "a2", "a3", "a4", "a5", "a6", "a7", "a8",
+ "a9", "a10", "a11", "a12", "a13", "a14", "a15", "a16", "a17",
+ "a18", "a19", "a20", "a21", "a22", "a23", "a24", "a25", "a26",
+ "a27", "a28", "a29", "a30", "a31", "a32", "a33", "a34", "a35",
+ "a36", "a37", "a38", "a39", "a40", "a41", "a42", "a43", "a44",
+ "a45", "a46", "a47", "a48", "a49", "a50", "a51", "a52", "a53",
+ "a54", "a55", "a56", "a57", "a58", "a59", "a60", "a61", "a62",
+ "a63", "a64", "a65", "a66", "a67", "a68", "a69", "a70", "a71",
+ "a72", "a73", "a74", "a75", "a76", "a77", "a78", "a79", "a80",
+ "a81", "a82", "a83", "a84", "a85", "a86", "a87", "a88", "a89",
+ "a90", "a91", "a92", "a93", "a94", "a95", "a96", "a97", "a98",
+ "a99", "a100", "a101", "a102", "a103", "a104", "a105", "a106", "a107",
+ "a108", "a109", "a110", "a111", "a112", "a113", "a114", "a115", "a116",
+ "a117", "a118", "a119", "a120", "a121", "a122", "a123", "a124", "a125",
+ "a126", "a127", "a128", "a129", "a130", "a131", "a132", "a133", "a134",
+ "a135", "a136", "a137", "a138", "a139", "a140", "a141", "a142", "a143",
+ "a144", "a145", "a146", "a147", "a148", "a149", "a150", "a151", "a152",
+ "a153", "a154", "a155", "a156", "a157", "a158", "a159", "a160", "a161",
+ "a162", "a163", "a164", "a165", "a166", "a167", "a168", "a169", "a170",
+ "a171", "a172", "a173", "a174", "a175", "a176", "a177", "a178", "a179",
+ "a180", "a181", "a182", "a183", "a184", "a185", "a186", "a187", "a188",
+ "a189", "a190", "a191", "a192", "a193", "a194", "a195", "a196", "a197",
+ "a198", "a199", "a200", "a201", "a202", "a203", "a204", "a205", "a206",
+ "a207", "a208", "a209", "a210", "a211", "a212", "a213", "a214", "a215",
+ "a216", "a217", "a218", "a219", "a220", "a221", "a222", "a223", "a224",
+ "a225", "a226", "a227", "a228", "a229", "a230", "a231", "a232", "a233",
+ "a234", "a235", "a236", "a237", "a238", "a239", "a240", "a241", "a242",
+ "a243", "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",
+ "a252", "a253", "a254", "a255",
+ // High VGPRs v256-v1023, addressable only on gfx1250+ (1024-addressable-vgprs).
+ // Kept last so getGCCRegNames() can drop them on targets without that feature.
"v256", "v257", "v258", "v259", "v260", "v261", "v262", "v263", "v264",
"v265", "v266", "v267", "v268", "v269", "v270", "v271", "v272", "v273",
"v274", "v275", "v276", "v277", "v278", "v279", "v280", "v281", "v282",
@@ -191,57 +238,20 @@ const char *const AMDGPUTargetInfo::GCCRegNames[] = {
"v994", "v995", "v996", "v997", "v998", "v999", "v1000", "v1001", "v1002",
"v1003", "v1004", "v1005", "v1006", "v1007", "v1008", "v1009", "v1010", "v1011",
"v1012", "v1013", "v1014", "v1015", "v1016", "v1017", "v1018", "v1019", "v1020",
- "v1021", "v1022", "v1023",
- "s0", "s1", "s2", "s3", "s4",
- "s5", "s6", "s7", "s8", "s9", "s10", "s11", "s12", "s13",
- "s14", "s15", "s16", "s17", "s18", "s19", "s20", "s21", "s22",
- "s23", "s24", "s25", "s26", "s27", "s28", "s29", "s30", "s31",
- "s32", "s33", "s34", "s35", "s36", "s37", "s38", "s39", "s40",
- "s41", "s42", "s43", "s44", "s45", "s46", "s47", "s48", "s49",
- "s50", "s51", "s52", "s53", "s54", "s55", "s56", "s57", "s58",
- "s59", "s60", "s61", "s62", "s63", "s64", "s65", "s66", "s67",
- "s68", "s69", "s70", "s71", "s72", "s73", "s74", "s75", "s76",
- "s77", "s78", "s79", "s80", "s81", "s82", "s83", "s84", "s85",
- "s86", "s87", "s88", "s89", "s90", "s91", "s92", "s93", "s94",
- "s95", "s96", "s97", "s98", "s99", "s100", "s101", "s102", "s103",
- "s104", "s105", "s106", "s107", "s108", "s109", "s110", "s111", "s112",
- "s113", "s114", "s115", "s116", "s117", "s118", "s119", "s120", "s121",
- "s122", "s123", "s124", "s125", "s126", "s127", "exec", "vcc", "scc",
- "m0", "flat_scratch", "exec_lo", "exec_hi", "vcc_lo", "vcc_hi",
- "flat_scratch_lo", "flat_scratch_hi",
- "a0", "a1", "a2", "a3", "a4", "a5", "a6", "a7", "a8",
- "a9", "a10", "a11", "a12", "a13", "a14", "a15", "a16", "a17",
- "a18", "a19", "a20", "a21", "a22", "a23", "a24", "a25", "a26",
- "a27", "a28", "a29", "a30", "a31", "a32", "a33", "a34", "a35",
- "a36", "a37", "a38", "a39", "a40", "a41", "a42", "a43", "a44",
- "a45", "a46", "a47", "a48", "a49", "a50", "a51", "a52", "a53",
- "a54", "a55", "a56", "a57", "a58", "a59", "a60", "a61", "a62",
- "a63", "a64", "a65", "a66", "a67", "a68", "a69", "a70", "a71",
- "a72", "a73", "a74", "a75", "a76", "a77", "a78", "a79", "a80",
- "a81", "a82", "a83", "a84", "a85", "a86", "a87", "a88", "a89",
- "a90", "a91", "a92", "a93", "a94", "a95", "a96", "a97", "a98",
- "a99", "a100", "a101", "a102", "a103", "a104", "a105", "a106", "a107",
- "a108", "a109", "a110", "a111", "a112", "a113", "a114", "a115", "a116",
- "a117", "a118", "a119", "a120", "a121", "a122", "a123", "a124", "a125",
- "a126", "a127", "a128", "a129", "a130", "a131", "a132", "a133", "a134",
- "a135", "a136", "a137", "a138", "a139", "a140", "a141", "a142", "a143",
- "a144", "a145", "a146", "a147", "a148", "a149", "a150", "a151", "a152",
- "a153", "a154", "a155", "a156", "a157", "a158", "a159", "a160", "a161",
- "a162", "a163", "a164", "a165", "a166", "a167", "a168", "a169", "a170",
- "a171", "a172", "a173", "a174", "a175", "a176", "a177", "a178", "a179",
- "a180", "a181", "a182", "a183", "a184", "a185", "a186", "a187", "a188",
- "a189", "a190", "a191", "a192", "a193", "a194", "a195", "a196", "a197",
- "a198", "a199", "a200", "a201", "a202", "a203", "a204", "a205", "a206",
- "a207", "a208", "a209", "a210", "a211", "a212", "a213", "a214", "a215",
- "a216", "a217", "a218", "a219", "a220", "a221", "a222", "a223", "a224",
- "a225", "a226", "a227", "a228", "a229", "a230", "a231", "a232", "a233",
- "a234", "a235", "a236", "a237", "a238", "a239", "a240", "a241", "a242",
- "a243", "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",
- "a252", "a253", "a254", "a255"
+ "v1021", "v1022", "v1023"
};
+// Number of trailing high-VGPR names (v256-v1023) in GCCRegNames that are only
+// valid on gfx1250+ (targets with the 1024-addressable-vgprs feature).
+static constexpr size_t NumHighVGPRRegNames = 1024 - 256;
+
ArrayRef<const char *> AMDGPUTargetInfo::getGCCRegNames() const {
- return llvm::ArrayRef(GCCRegNames);
+ ArrayRef<const char *> Names(GCCRegNames);
+ // v256-v1023 are addressable only on gfx1250+; hide them elsewhere so they
+ // cannot be named in inline asm / clobbers on targets with just 256 VGPRs.
+ if (!has1024AddressableVGPRs())
+ Names = Names.drop_back(NumHighVGPRRegNames);
+ return Names;
}
bool AMDGPUTargetInfo::initFeatureMap(
diff --git a/clang/lib/Basic/Targets/AMDGPU.h b/clang/lib/Basic/Targets/AMDGPU.h
index f8933ebee8ffd..f5134d43557a0 100644
--- a/clang/lib/Basic/Targets/AMDGPU.h
+++ b/clang/lib/Basic/Targets/AMDGPU.h
@@ -42,6 +42,13 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
/// Whether having image instructions.
bool HasImage = false;
+ /// Whether the target has 1024 addressable VGPRs (v0-v1023). When false, only
+ /// v0-v255 are valid register names. Gates the high VGPR names in
+ /// getGCCRegNames() so v256-v1023 cannot be named on targets with 256 VGPRs.
+ bool has1024AddressableVGPRs() const {
+ return GPUKind == llvm::AMDGPU::GK_GFX1250;
+ }
+
/// Target ID is device name followed by optional feature name postfixed
/// by plus or minus sign delimitted by colon, e.g. gfx908:xnack+:sramecc-.
/// If the target ID contains feature+, map it to true.
diff --git a/clang/test/Sema/amdgpu-high-vgpr-regname.c b/clang/test/Sema/amdgpu-high-vgpr-regname.c
new file mode 100644
index 0000000000000..20f018f691f1a
--- /dev/null
+++ b/clang/test/Sema/amdgpu-high-vgpr-regname.c
@@ -0,0 +1,15 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -fsyntax-only -verify=gfx1250 %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -fsyntax-only -verify=gfx942 %s
+
+// VGPRs v256-v1023 are addressable only on gfx1250+ (1024 addressable VGPRs);
+// naming them in inline asm must be rejected on targets with only 256 VGPRs.
+
+// gfx1250-no-diagnostics
+
+void low_vgpr(void) {
+ __asm__ volatile("" ::: "v100");
+}
+
+void high_vgpr(void) {
+ __asm__ volatile("" ::: "v300"); // gfx942-error {{unknown register name 'v300' in asm}}
+}
>From f30e4933735bd1e0baa0cd07f5c6b3a1632d4515 Mon Sep 17 00:00:00 2001
From: demonsan <1462264754 at qq.com>
Date: Wed, 5 Aug 2026 11:05:48 +0000
Subject: [PATCH 14/19] [AMDGPU] Keep register pins alive through a
by-reference lambda capture
The amdgpu_pin_{vgpr,agpr} attribute was recognised only when the store's
destination pointer was the variable's own alloca, which EmitAutoVarAlloca
registers in AMDGPUPinnedLocals. A variable written from inside a lambda that
captured it by reference is reached through the capture field instead, in a
different function whose AMDGPUPinnedLocals is empty, so the store address never
matched and the attribute was silently dropped -- no llvm.amdgcn.pin.* was
emitted at all.
Register the capture-field storage when EmitDeclRefLValue forms an lvalue for a
captured variable that carries the attribute. Nothing has to be threaded across
functions: inside the lambda body the DeclRefExpr still names the original
VarDecl, so the register number is read straight off the attribute.
This is what real tile kernels need -- accumulator and fragment writes there go
through always_inline helper lambdas rather than being spelled out at the
assignment site. On a gfx1250 bf16 GEMM the attribute now reaches IR (218 pin
intrinsics after inlining, previously none). A store through an explicit
reference variable (T &r = v; r = x;) is a separate binding and still not
tracked.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
clang/lib/CodeGen/CGExpr.cpp | 20 ++++++++++++++++++--
clang/lib/CodeGen/CodeGenFunction.h | 7 +++++++
clang/test/CodeGenHIP/amdgpu-pin-attr.hip | 12 ++++++++++++
3 files changed, 37 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index fdae64bb60894..024d8cabf0e65 100644
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -3055,6 +3055,19 @@ void CodeGenFunction::EmitStoreThroughLValue(RValue Src, LValue Dst,
EmitStoreOfScalar(SV, Dst, isInit);
}
+void CodeGenFunction::tryTrackAMDGPUPinnedCapture(const VarDecl *VD,
+ const LValue &LV) {
+ bool IsAGPR = VD->hasAttr<AMDGPUPinAGPRAttr>();
+ if (!IsAGPR && !VD->hasAttr<AMDGPUPinVGPRAttr>())
+ return;
+ if (!getTarget().getTriple().isAMDGCN() || !LV.isSimple())
+ return;
+ const Expr *RegE = IsAGPR ? VD->getAttr<AMDGPUPinAGPRAttr>()->getReg()
+ : VD->getAttr<AMDGPUPinVGPRAttr>()->getReg();
+ unsigned Reg = RegE->EvaluateKnownConstInt(getContext()).getZExtValue();
+ AMDGPUPinnedLocals[LV.getPointer(*this)] = {IsAGPR, Reg};
+}
+
llvm::Value *CodeGenFunction::emitAMDGPUPinnedValue(llvm::Value *V,
llvm::Value *Addr) {
auto It = AMDGPUPinnedLocals.find(Addr);
@@ -3727,8 +3740,11 @@ LValue CodeGenFunction::EmitDeclRefLValue(const DeclRefExpr *E) {
// Check for captured variables.
if (E->refersToEnclosingVariableOrCapture()) {
VD = VD->getCanonicalDecl();
- if (auto *FD = LambdaCaptureFields.lookup(VD))
- return EmitCapturedFieldLValue(*this, FD, CXXABIThisValue);
+ if (auto *FD = LambdaCaptureFields.lookup(VD)) {
+ LValue CapLVal = EmitCapturedFieldLValue(*this, FD, CXXABIThisValue);
+ tryTrackAMDGPUPinnedCapture(VD, CapLVal);
+ return CapLVal;
+ }
if (CapturedStmtInfo) {
auto I = LocalDeclMap.find(VD);
if (I != LocalDeclMap.end()) {
diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index b469840fc7f26..f85dbe8186269 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -1565,6 +1565,13 @@ class CodeGenFunction : public CodeGenTypeCache {
/// otherwise return \p V unchanged.
llvm::Value *emitAMDGPUPinnedValue(llvm::Value *V, llvm::Value *Addr);
+ /// Record \p LV's storage as pinned storage when \p VD carries an
+ /// amdgpu_pin_{vgpr,agpr} attribute. The variable's own declaration registers
+ /// its alloca (\sa EmitAutoVarAlloca), but a by-reference capture reaches it
+ /// through the capture field of a different function, so the store address
+ /// never matches the alloca and the pin would be silently dropped.
+ void tryTrackAMDGPUPinnedCapture(const VarDecl *VD, const LValue &LV);
+
// Keep track of the cleanups for callee-destructed parameters pushed to the
// cleanup stack so that they can be deactivated later.
llvm::DenseMap<const ParmVarDecl *, EHScopeStack::stable_iterator>
diff --git a/clang/test/CodeGenHIP/amdgpu-pin-attr.hip b/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
index 654ebe8a433e2..721ca4ad60b64 100644
--- a/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
+++ b/clang/test/CodeGenHIP/amdgpu-pin-attr.hip
@@ -34,6 +34,18 @@ __attribute__((device)) void pin_wide(float12 *out, float12 in) {
*out = x;
}
+// A store from inside a lambda that captured the variable by reference reaches
+// it through the capture field of a different function, not through its alloca,
+// so the pin has to be recognised from the referenced declaration.
+// CHECK-LABEL: define{{.*}}pin_captured
+// CHECK: call <2 x i32> @llvm.amdgcn.pin.vgpr.v2i32(<2 x i32> %{{[0-9]+}}, i32 6)
+__attribute__((device)) void pin_captured(float2 *out, float2 in) {
+ __attribute__((amdgpu_pin_vgpr(6))) float2 x;
+ auto set = [&](float2 v) { x = v; };
+ set(in);
+ *out = x;
+}
+
// A constant-expression argument (here via a template parameter) is accepted and
// evaluated at instantiation.
template <int N>
>From ed1fc3dc72addd903760b95374d4fab314cc68bf Mon Sep 17 00:00:00 2001
From: demonsan <1462264754 at qq.com>
Date: Wed, 5 Aug 2026 11:06:03 +0000
Subject: [PATCH 15/19] [AMDGPU] Don't pre-color a PHI operand; report why a
pin was not pre-colored
SIPreColorPins collected every operand referencing the pinned component and
substituted the physical tuple into all of them, PHI operands included. That
crashes before PHIElimination can lower the PHI: LiveVariables walks PHI sources
through getVarInfo(), which asserts the register is virtual. It is reached
whenever a pinned value defined outside a loop flows into a loop-carried PHI --
several accumulators initialised ahead of the loop is enough. Leave such a pin
to the soft path.
The pass was also completely opaque: a pin that quietly degraded looked exactly
like one that was honoured, since the VGPR-footprint occupancy cap is raised
either way. Add hard/soft/no-op statistics and record which condition rejected
the pin, so -stats and -debug-only=si-pre-color-pins answer that directly. On a
gfx1250 GEMM this immediately shows 2 hard pins against 214 soft ones, 176 of
them because the pinned component is defined by a REG_SEQUENCE -- the shape a
loop-carried wide accumulator always takes, since values crossing a block
boundary are passed per dword and reassembled.
The new test crashes without the operand check and is unchanged by the
diagnostics.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 31 +++++++++++++++++++++
llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll | 31 +++++++++++++++++++++
2 files changed, 62 insertions(+)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 697728cddd015..8fda103ed5270 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -29,6 +29,7 @@
#include "SIRegisterInfo.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
+#include "llvm/ADT/Statistic.h"
#include "llvm/CodeGen/LivePhysRegs.h"
#include "llvm/CodeGen/MachineFunctionPass.h"
#include "llvm/CodeGen/MachineRegisterInfo.h"
@@ -38,6 +39,10 @@ using namespace llvm;
#define DEBUG_TYPE "si-pre-color-pins"
+STATISTIC(NumHardPins, "Number of values pre-colored to the requested register");
+STATISTIC(NumSoftPins, "Number of pins degraded to a soft allocation hint");
+STATISTIC(NumNoOpPins, "Number of pins dropped as a no-op");
+
static cl::opt<bool> EnableHardPin(
"amdgpu-hard-pin-regs", cl::init(true), cl::Hidden,
cl::desc("Use hard register pre-coloring for llvm.amdgcn.pin.* (else soft "
@@ -149,6 +154,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (Src.isVirtual())
MRI.constrainRegClass(Src, TRI->getEquivalentVGPRClass(RC));
Pin->eraseFromParent();
+ ++NumNoOpPins;
continue;
}
// Only VGPR pins drive the occupancy cap (see below); AGPRs are a separate
@@ -288,6 +294,8 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
bool Hard = EnableHardPin && PR && Src.isVirtual() && Dst.isVirtual();
+ // Why a pin could not be pre-colored, for -debug-only=si-pre-color-pins.
+ const char *SoftWhy = Hard ? "?" : "disabled or non-virtual";
// Deterministic AGPR placement for a load tuple: when the pinned value is a
// REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a
@@ -404,6 +412,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
RS->eraseFromParent();
Pin->eraseFromParent();
NeedRecomputeLiveIns = true;
+ ++NumHardPins;
continue;
}
}
@@ -454,6 +463,9 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MachineInstr &DefMI : MRI.def_instructions(R)) {
if (DefMI.isPHI() || DefMI.isRegSequence() || DefMI.isImplicitDef()) {
Hard = false;
+ SoftWhy = DefMI.isPHI() ? "PHI def"
+ : DefMI.isRegSequence() ? "REG_SEQUENCE def"
+ : "IMPLICIT_DEF";
break;
}
}
@@ -462,10 +474,22 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MachineOperand &MO : MRI.reg_operands(R)) {
if (MO.getParent() == Pin)
continue; // the pin itself is erased
+ // A PHI operand must stay virtual: LiveVariables walks PHI sources
+ // through getVarInfo(), which only accepts virtual registers, so a
+ // physreg there crashes before PHIElimination can lower it. This is
+ // reached when a pinned value defined in one block flows into a
+ // loop-carried PHI (several pinned accumulators initialised outside
+ // the loop). Leave the whole pin to the soft path.
+ if (MO.getParent()->isPHI()) {
+ Hard = false;
+ SoftWhy = "PHI use";
+ break;
+ }
MCRegister Tgt =
MO.getSubReg() ? TRI->getSubReg(PR, MO.getSubReg()) : PR;
if (!Tgt) {
Hard = false;
+ SoftWhy = "no such subregister";
break;
}
const TargetRegisterClass *OpRC =
@@ -473,6 +497,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
TRI);
if (OpRC && !OpRC->contains(Tgt)) {
Hard = false;
+ SoftWhy = "operand class rejects the physreg";
break;
}
ToRewrite.push_back(&MO);
@@ -502,6 +527,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MCRegUnit U : TRI->regunits(PR)) {
if (Claimed.contains(U)) {
Hard = false;
+ SoftWhy = "overlaps an earlier hard pin";
break;
}
}
@@ -530,6 +556,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MachineOperand *MO : SubUses)
if (MO->getParent()->getParent() != CopyMBB) {
Hard = false;
+ SoftWhy = "subregister uses span blocks";
break;
}
if (Hard) {
@@ -565,9 +592,13 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MCRegUnit U : TRI->regunits(PR))
Claimed.insert(U);
Pin->eraseFromParent();
+ ++NumHardPins;
continue;
}
+ ++NumSoftPins;
+ LLVM_DEBUG(dbgs() << "pin to " << (WantAGPR ? 'a' : 'v') << RegNo
+ << " not pre-colored: " << SoftWhy << '\n');
// Soft fallback: COPY + register-allocation hint (a no-op hint if the
// physical tuple was illegal).
BuildMI(*Pin->getParent(), Pin, Pin->getDebugLoc(),
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
index f6cd994e99a1f..0b33da82dad68 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
@@ -6,6 +6,8 @@
; cover the pinned range.
declare <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float>, i32 immarg)
+declare <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32>, i32 immarg)
+declare <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.bf16(i1, <16 x bfloat>, i1, <16 x bfloat>, i16, <8 x float>, i1, i1)
; CHECK-LABEL: {{^}}pin_high_vgpr:
; CHECK: s_set_vgpr_msb
@@ -17,3 +19,32 @@ define amdgpu_kernel void @pin_high_vgpr(ptr addrspace(1) %p) {
store <4 x float> %pv, ptr addrspace(1) %p
ret void
}
+
+; A pinned value defined outside a loop and carried into it reaches a PHI. The
+; physical tuple must not be substituted into the PHI operand: LiveVariables
+; walks PHI sources through getVarInfo(), which asserts on a physical register,
+; so the pin falls back to the soft path instead. Two pinned accumulators keep
+; both PHIs live across the back edge.
+; CHECK-LABEL: {{^}}pin_into_loop_phi:
+; CHECK: v_wmma_f32_16x16x32_bf16
+; CHECK: s_endpgm
+define amdgpu_kernel void @pin_into_loop_phi(ptr addrspace(1) %o) {
+entry:
+ %i0 = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> zeroinitializer, i32 100)
+ %i1 = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> zeroinitializer, i32 108)
+ br label %loop
+
+loop:
+ %a1 = phi <8 x i32> [ %i1, %entry ], [ %d.i, %loop ]
+ %a0 = phi <8 x i32> [ %i0, %entry ], [ zeroinitializer, %loop ]
+ %c = bitcast <8 x i32> %a1 to <8 x float>
+ %d = call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.bf16(i1 false, <16 x bfloat> zeroinitializer, i1 false, <16 x bfloat> zeroinitializer, i16 0, <8 x float> %c, i1 false, i1 false)
+ %d.i = bitcast <8 x float> %d to <8 x i32>
+ br i1 false, label %exit, label %loop
+
+exit:
+ %o1 = getelementptr <8 x i32>, ptr addrspace(1) %o, i32 1
+ store <8 x i32> %a0, ptr addrspace(1) %o
+ store <8 x i32> %a1, ptr addrspace(1) %o1
+ ret void
+}
>From fd426b442576aeb178a607544875b7485200af23 Mon Sep 17 00:00:00 2001
From: demonsan <1462264754 at qq.com>
Date: Wed, 5 Aug 2026 11:16:25 +0000
Subject: [PATCH 16/19] [AMDGPU] Take the load-tuple pre-coloring path for VGPR
pins too
A pinned value that a REG_SEQUENCE defines is rejected by the general path,
which needs a single reaching def. The load-tuple path handles exactly that
shape by retargeting each element's def to a fixed physical sub-register, but it
only ran for AGPR pins -- so on a target without an AGPR file every such pin was
dropped. That covers a WMMA B operand read out of LDS by two ds_reads, which is
how a tile kernel feeds the matrix pipe on gfx1250.
Run the path for VGPR pins as well, with two guards for shapes the AGPR side
cannot reach because it only ever pins MFMA A/B inputs: a tied (two-address) use
is a WMMA accumulator and TwoAddressInstruction requires both ends of a tie to
be virtual, and a PHI operand has to stay virtual for LiveVariables. Both fall
back to the general path.
The new test's B operand lands in the requested v[128:135] instead of wherever
the allocator put it; without the change the pin is silently dropped.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 32 +++++++++++++++++----
llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll | 21 ++++++++++++++
2 files changed, 47 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 8fda103ed5270..02905bcd3b4c6 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -297,12 +297,15 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
// Why a pin could not be pre-colored, for -debug-only=si-pre-color-pins.
const char *SoftWhy = Hard ? "?" : "disabled or non-virtual";
- // Deterministic AGPR placement for a load tuple: when the pinned value is a
- // REG_SEQUENCE of (folded) AGPR loads, rewrite each element's def to a
- // fixed physical AGPR sub-register. Otherwise the MFMA A/B operands are AV
- // and the allocator moves them back to VGPR under low pressure
- // (non-deterministic).
- if (Hard && WantAGPR) {
+ // Deterministic placement for a load tuple: when the pinned value is a
+ // REG_SEQUENCE of loads, rewrite each element's def to a fixed physical
+ // sub-register. For an AGPR pin this stops the allocator moving MFMA A/B
+ // operands back to VGPR under low pressure (they are AV-classed, so the
+ // placement is otherwise non-deterministic). For a VGPR pin it is the only
+ // way to place a value the hardware builds from more than one load, such as
+ // a WMMA B operand that two ds_reads assemble out of LDS -- the general path
+ // below rejects any REG_SEQUENCE def.
+ if (Hard) {
MachineInstr *RS = MRI.getVRegDef(Src);
MachineBasicBlock *PinMBB = Pin->getParent();
bool Ok = RS && RS->isRegSequence() && RS->getParent() == PinMBB;
@@ -319,6 +322,23 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
Ok = false;
break;
}
+ // A VGPR pin reaches shapes the AGPR path never sees, because the
+ // AGPR path only ever pins MFMA A/B inputs. A tied (two-address)
+ // use is a WMMA/MFMA accumulator, and TwoAddressInstruction
+ // requires both ends of a tie to be virtual; a PHI operand must
+ // stay virtual as well, since LiveVariables walks PHI sources with
+ // getVarInfo(). Leave both to the general path.
+ if (!WantAGPR) {
+ if (U.isPHI()) {
+ Ok = false;
+ break;
+ }
+ for (const MachineOperand &O : U.operands())
+ if (O.isReg() && O.isUse() && O.isTied() && O.getReg() == WL[I])
+ Ok = false;
+ if (!Ok)
+ break;
+ }
if (U.isCopy() || U.isRegSequence() || U.isPHI() ||
U.getOpcode() == TargetOpcode::INSERT_SUBREG ||
U.getOpcode() == TargetOpcode::EXTRACT_SUBREG)
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
index 0b33da82dad68..fc45f18247d90 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
@@ -20,6 +20,27 @@ define amdgpu_kernel void @pin_high_vgpr(ptr addrspace(1) %p) {
ret void
}
+; A WMMA B operand that two loads assemble is defined by a REG_SEQUENCE, which
+; the general path rejects outright. The load-tuple path now takes VGPR pins as
+; well, so the operand lands in the requested tuple instead of wherever the
+; allocator puts it. Only the AGPR form of this path ran before, and gfx1250 has
+; no AGPR file, so such a pin was always dropped here.
+; CHECK-LABEL: {{^}}pin_two_load_tuple:
+; CHECK: v_wmma_f32_16x16x32_bf16 v[{{[0-9:]+}}], v[128:135],
+define amdgpu_kernel void @pin_two_load_tuple(ptr addrspace(1) %o, ptr addrspace(1) %pa, ptr addrspace(1) %pb) {
+ %p1 = getelementptr <8 x bfloat>, ptr addrspace(1) %pb, i64 1
+ %b0 = load <8 x bfloat>, ptr addrspace(1) %pb, align 16
+ %b1 = load <8 x bfloat>, ptr addrspace(1) %p1, align 16
+ %b = shufflevector <8 x bfloat> %b0, <8 x bfloat> %b1, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15>
+ %b.i = bitcast <16 x bfloat> %b to <8 x i32>
+ %p = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> %b.i, i32 128)
+ %pb.v = bitcast <8 x i32> %p to <16 x bfloat>
+ %a = load <16 x bfloat>, ptr addrspace(1) %pa, align 32
+ %d = call <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.bf16(i1 false, <16 x bfloat> %pb.v, i1 false, <16 x bfloat> %a, i16 0, <8 x float> zeroinitializer, i1 false, i1 false)
+ store <8 x float> %d, ptr addrspace(1) %o
+ ret void
+}
+
; A pinned value defined outside a loop and carried into it reaches a PHI. The
; physical tuple must not be substituted into the PHI operand: LiveVariables
; walks PHI sources through getVarInfo(), which asserts on a physical register,
>From 2f5cac793bea655d170ba6c7afd66ef7bec74e6a Mon Sep 17 00:00:00 2001
From: demonsan <1462264754 at qq.com>
Date: Mon, 10 Aug 2026 02:25:58 +0000
Subject: [PATCH 17/19] [AMDGPU] Pin a load tuple that arrives as slices of
wider loads
The load-tuple path places a value the hardware builds from more than one load
by retargeting each REG_SEQUENCE element's def to a fixed physical
sub-register. It insisted that every element be a whole register and bailed as
soon as one carried a subregister index. But that is the ordinary shape of a
wide load off a divergent address: a 32-byte load is selected as two dwordx4
loads whose lanes reach the REG_SEQUENCE as %wide.subN, so the pin was dropped
and the value went wherever the allocator liked. Nothing about it is specific
to high registers -- a pin to v8 was lost the same way as one to v1016.
A slice cannot be retargeted on its own without stranding the rest of its load,
so place the wider def as a whole instead: getMatchingSuperReg turns "this lane
belongs at that lane of PR" into the tuple the def must occupy, all slices of
one def have to agree on it, and it has to stay inside PR. A permuted,
misaligned or out-of-range layout fails one of those and still falls back to
soft.
Also stop charging occupancy for a pin that was not pre-colored. The cap was
raised from the requested range before the pin's fate was known, so a soft hint
the allocator went on to ignore still cost waves: the float8 above capped a
kernel to one wave (next_free_vgpr 513) while its metadata vgpr_count said 12.
Record the footprint at the two pre-coloring sites instead.
On gfx1250 an <8 x float> pinned to v1016 now occupies v[1016:1023] across both
of its loads, and the reported VGPR count follows the placement rather than the
request.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 78 ++++++++++++++-------
llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll | 24 ++++++-
2 files changed, 76 insertions(+), 26 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index 02905bcd3b4c6..d1f538c8a7f73 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -157,10 +157,15 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
++NumNoOpPins;
continue;
}
- // Only VGPR pins drive the occupancy cap (see below); AGPRs are a separate
+ // Only a VGPR pin that is actually honored drives the occupancy cap (see
+ // below), so this is recorded at the pre-coloring sites rather than here:
+ // a soft hint is free to go unused, and paying occupancy for a hint the
+ // allocator then ignores costs waves for nothing. AGPRs are a separate
// file that does not affect the VGPR budget.
- if (!WantAGPR)
- ReqVGPRs = std::max(ReqVGPRs, RegNo + NumRegs);
+ auto RecordVGPRFootprint = [&] {
+ if (!WantAGPR)
+ ReqVGPRs = std::max(ReqVGPRs, RegNo + NumRegs);
+ };
// Narrow the pinned value's register file to VGPR or AGPR (a class
// narrowing, not a physreg pin, so it also works for loop-carried PHIs and
@@ -351,23 +356,41 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
if (Ok && Claimed.contains(U))
Ok = false;
- // Collect element (reg, subreg-index) pairs. Each element must be defined
- // directly by a memory load: this path retargets those load defs to fixed
- // physical AGPR sub-registers. If an element is instead a subregister copy
- // of a wider load (e.g. a dwordx4 load split into dword lanes), retargeting
- // it produces malformed physreg liveness, so bail and let the general path
- // fall back to soft.
- SmallVector<std::pair<Register, unsigned>, 16> Elems;
+ // Map each element's defining register onto a physical (sub)register of
+ // PR. An element either covers its REG_SEQUENCE slot outright, or is a
+ // subregister slice of a wider def: a 32-byte load, for instance, is
+ // selected as two dwordx4 loads whose lanes reach the REG_SEQUENCE as
+ // %wide.subN. Retargeting such a lane on its own would leave the rest of
+ // the wider def behind, so the whole def is placed instead --
+ // getMatchingSuperReg derives the tuple it must occupy and rejects a
+ // permuted or misaligned layout, and the tuple has to stay inside PR.
+ SmallVector<std::pair<Register, MCRegister>, 16> Elems;
if (Ok)
for (unsigned I = 1; I + 1 < RS->getNumOperands(); I += 2) {
const MachineOperand &Reg = RS->getOperand(I);
const MachineOperand &Sub = RS->getOperand(I + 1);
- if (!Reg.isReg() || !Reg.getReg().isVirtual() || Reg.getSubReg() ||
- !Sub.isImm() || !TRI->getSubReg(PR, Sub.getImm())) {
+ if (!Reg.isReg() || !Reg.getReg().isVirtual() || !Sub.isImm()) {
+ Ok = false;
+ break;
+ }
+ MCRegister Tgt = TRI->getSubReg(PR, Sub.getImm());
+ if (Tgt && Reg.getSubReg())
+ Tgt = TRI->getMatchingSuperReg(Tgt, Reg.getSubReg(),
+ MRI.getRegClass(Reg.getReg()));
+ if (!Tgt || !TRI->isSubRegisterEq(PR, Tgt)) {
+ Ok = false;
+ break;
+ }
+ auto *Prev =
+ find_if(Elems, [&](const std::pair<Register, MCRegister> &E) {
+ return E.first == Reg.getReg();
+ });
+ if (Prev == Elems.end())
+ Elems.emplace_back(Reg.getReg(), Tgt);
+ else if (Prev->second != Tgt) {
Ok = false;
break;
}
- Elems.push_back({Reg.getReg(), (unsigned)Sub.getImm()});
}
// Every use of the pinned result and of each element must legally accept
@@ -391,26 +414,29 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
}
if (Ok)
- for (auto [Elem, SubIdx] : Elems) {
- MCRegister PhysSub = TRI->getSubReg(PR, SubIdx);
- for (MachineOperand &MO : MRI.reg_operands(Elem))
- if (!LegalHere(MO, PhysSub)) {
+ for (auto [Elem, Phys] : Elems) {
+ for (MachineOperand &MO : MRI.reg_operands(Elem)) {
+ MCRegister T =
+ MO.getSubReg() ? TRI->getSubReg(Phys, MO.getSubReg()) : Phys;
+ if (!LegalHere(MO, T)) {
Ok = false;
break;
}
+ }
if (!Ok)
break;
}
if (Ok) {
- // Point each element's def/uses at its physical AGPR sub-register.
- for (auto [Elem, SubIdx] : Elems) {
- MCRegister PhysSub = TRI->getSubReg(PR, SubIdx);
+ // Point each element's def/uses at its physical (sub)register.
+ for (auto [Elem, Phys] : Elems) {
SmallVector<MachineOperand *, 4> Ops;
for (MachineOperand &MO : MRI.reg_operands(Elem))
Ops.push_back(&MO);
for (MachineOperand *MO : Ops) {
- MO->setReg(PhysSub);
+ MCRegister T =
+ MO->getSubReg() ? TRI->getSubReg(Phys, MO->getSubReg()) : Phys;
+ MO->setReg(T);
MO->setSubReg(0);
MO->setIsRenamable(false);
}
@@ -432,6 +458,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
RS->eraseFromParent();
Pin->eraseFromParent();
NeedRecomputeLiveIns = true;
+ RecordVGPRFootprint();
++NumHardPins;
continue;
}
@@ -612,6 +639,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
for (MCRegUnit U : TRI->regunits(PR))
Claimed.insert(U);
Pin->eraseFromParent();
+ RecordVGPRFootprint();
++NumHardPins;
continue;
}
@@ -643,10 +671,10 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
// Cap occupancy so a wide VGPR-resident pinned value fits the per-wave budget
- // without the user setting __launch_bounds__. Only VGPR footprints drive
- // this: AGPRs are a separate file, so feeding an AGPR count into the VGPR
- // occupancy formula would wrongly raise occupancy and spill the VGPR
- // accumulator.
+ // without the user setting __launch_bounds__. Only the VGPR footprint of a
+ // pre-colored pin drives this: AGPRs are a separate file, so feeding an AGPR
+ // count into the VGPR occupancy formula would wrongly raise occupancy and
+ // spill the VGPR accumulator.
auto *MFI = MF.getInfo<SIMachineFunctionInfo>();
if (unsigned Req = ReqVGPRs) {
// Occupancy achievable while reserving `Req` registers per wave; cap the
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
index fc45f18247d90..e5e162ba444f0 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
@@ -8,6 +8,7 @@
declare <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float>, i32 immarg)
declare <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32>, i32 immarg)
declare <8 x float> @llvm.amdgcn.wmma.f32.16x16x32.bf16(i1, <16 x bfloat>, i1, <16 x bfloat>, i16, <8 x float>, i1, i1)
+declare i32 @llvm.amdgcn.workitem.id.x()
; CHECK-LABEL: {{^}}pin_high_vgpr:
; CHECK: s_set_vgpr_msb
@@ -41,14 +42,35 @@ define amdgpu_kernel void @pin_two_load_tuple(ptr addrspace(1) %o, ptr addrspace
ret void
}
+; A 32-byte load off a divergent address is selected as two dwordx4 loads whose
+; lanes reach the REG_SEQUENCE as subregister slices of the wider load. Placing
+; a lane on its own would strand the rest of its load, so each load is placed as
+; a whole onto its half of the pinned tuple.
+; CHECK-LABEL: {{^}}pin_split_wide_load:
+; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[304:307]*/
+; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[300:303]*/
+; CHECK: .set .Lpin_split_wide_load.num_vgpr, 308
+define amdgpu_kernel void @pin_split_wide_load(ptr addrspace(1) %in, ptr addrspace(1) %out) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %idx = sext i32 %tid to i64
+ %a = getelementptr inbounds <8 x i32>, ptr addrspace(1) %in, i64 %idx
+ %v = load <8 x i32>, ptr addrspace(1) %a, align 32
+ %pv = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> %v, i32 300)
+ %o = getelementptr inbounds <8 x i32>, ptr addrspace(1) %out, i64 %idx
+ store <8 x i32> %pv, ptr addrspace(1) %o, align 32
+ ret void
+}
+
; A pinned value defined outside a loop and carried into it reaches a PHI. The
; physical tuple must not be substituted into the PHI operand: LiveVariables
; walks PHI sources through getVarInfo(), which asserts on a physical register,
; so the pin falls back to the soft path instead. Two pinned accumulators keep
-; both PHIs live across the back edge.
+; both PHIs live across the back edge. A soft hint costs no occupancy: the VGPR
+; count reflects what the allocator used, not the range the pins asked for.
; CHECK-LABEL: {{^}}pin_into_loop_phi:
; CHECK: v_wmma_f32_16x16x32_bf16
; CHECK: s_endpgm
+; CHECK: .set .Lpin_into_loop_phi.num_vgpr, 16
define amdgpu_kernel void @pin_into_loop_phi(ptr addrspace(1) %o) {
entry:
%i0 = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> zeroinitializer, i32 100)
>From ad54993aac47893553d1d732f47bd22dc1397d33 Mon Sep 17 00:00:00 2001
From: demonsan <1462264754 at qq.com>
Date: Mon, 10 Aug 2026 02:28:34 +0000
Subject: [PATCH 18/19] [AMDGPU] Fold away a REG_SEQUENCE that only rebuilds a
pinned tuple
Pre-coloring rewrites every operand of the pinned value, including those inside
a REG_SEQUENCE that a consumer needs as one wider register -- the two halves a
32-byte store reads, for instance. What is left is a REG_SEQUENCE whose sources
are all physical and consecutive, which the allocator can only satisfy by
copying each lane into a fresh virtual tuple. The placement is undone right
after it is made: the value is loaded into v[300:307] and then moved down to
v[0:7] by eight v_dual_mov before the stores read it.
Such a REG_SEQUENCE already names a physical tuple, so forward that tuple to
the uses and erase it. The first lane fixes the candidate and the rest must
agree, which rejects a permuted, gapped or misaligned run; uses whose operand
class will not take the tuple leave it alone. Folding one can expose another
when a wide value is reassembled in stages, so iterate, and only bother in
functions that actually pre-colored something.
The eight copies disappear from the test, the two stores read v[304:307] and
v[300:303] directly, and one s_set_vgpr_msb now covers the whole body instead
of the msb being toggled around the copies.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 72 +++++++++++++++++++++
llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll | 7 +-
2 files changed, 78 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index d1f538c8a7f73..c82276753102a 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -107,6 +107,63 @@ static MCRegister getPinPhysReg(const SIRegisterInfo *TRI,
return MCRegister();
}
+// Rewriting a pinned value's operands can leave a REG_SEQUENCE that does
+// nothing but reassemble a run of the pinned tuple back into a virtual
+// register -- the two halves a 32-byte store reads, say. The allocator
+// materializes that as one copy per lane, undoing the placement. If the run
+// names a physical tuple outright, forward it to the uses and drop the
+// REG_SEQUENCE. Returns true if it was folded away.
+static bool foldPhysRegSequence(const SIInstrInfo *TII,
+ const SIRegisterInfo *TRI,
+ MachineRegisterInfo &MRI, MachineInstr &RS) {
+ Register Def = RS.getOperand(0).getReg();
+ if (!Def.isVirtual() || RS.getNumOperands() < 3)
+ return false;
+
+ MCRegister Tuple;
+ for (unsigned I = 1; I + 1 < RS.getNumOperands(); I += 2) {
+ const MachineOperand &Src = RS.getOperand(I);
+ const MachineOperand &Sub = RS.getOperand(I + 1);
+ if (!Src.isReg() || !Src.getReg().isPhysical() || Src.getSubReg() ||
+ !Sub.isImm())
+ return false;
+ MCRegister Phys = Src.getReg().asMCReg();
+ // The first lane fixes the candidate tuple; the rest must agree with it,
+ // which rejects a permuted, gapped or misaligned run.
+ if (!Tuple)
+ Tuple =
+ TRI->getMatchingSuperReg(Phys, Sub.getImm(), MRI.getRegClass(Def));
+ else if (TRI->getSubReg(Tuple, Sub.getImm()) != Phys)
+ return false;
+ if (!Tuple)
+ return false;
+ }
+
+ SmallVector<MachineOperand *, 8> Uses;
+ for (MachineOperand &MO : MRI.reg_operands(Def)) {
+ if (MO.getParent() == &RS)
+ continue;
+ MCRegister T = MO.getSubReg() ? TRI->getSubReg(Tuple, MO.getSubReg())
+ : MCRegister(Tuple);
+ if (MO.isDef() || !T)
+ return false;
+ const TargetRegisterClass *OpRC =
+ MO.getParent()->getRegClassConstraint(MO.getOperandNo(), TII, TRI);
+ if (OpRC && !OpRC->contains(T))
+ return false;
+ Uses.push_back(&MO);
+ }
+
+ for (MachineOperand *MO : Uses) {
+ MO->setReg(MO->getSubReg() ? TRI->getSubReg(Tuple, MO->getSubReg())
+ : MCRegister(Tuple));
+ MO->setSubReg(0);
+ MO->setIsRenamable(false);
+ }
+ RS.eraseFromParent();
+ return true;
+}
+
bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
const SIInstrInfo *TII = ST.getInstrInfo();
@@ -129,6 +186,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
// no-ops.
DenseSet<MCRegUnit> Claimed;
bool NeedRecomputeLiveIns = false;
+ bool AnyHardPin = false;
unsigned ReqVGPRs =
0; // highest VGPR a pin needs, +1 (drives the occupancy cap)
for (MachineInstr *Pin : Pins) {
@@ -459,6 +517,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
Pin->eraseFromParent();
NeedRecomputeLiveIns = true;
RecordVGPRFootprint();
+ AnyHardPin = true;
++NumHardPins;
continue;
}
@@ -640,6 +699,7 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
Claimed.insert(U);
Pin->eraseFromParent();
RecordVGPRFootprint();
+ AnyHardPin = true;
++NumHardPins;
continue;
}
@@ -660,6 +720,18 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
Pin->eraseFromParent();
}
+ // Clean up the REG_SEQUENCEs the rewrite left behind. Folding one can expose
+ // another (a wide tuple reassembled in stages), so iterate to a fixpoint.
+ for (bool Folded = AnyHardPin; Folded;) {
+ Folded = false;
+ for (MachineBasicBlock &MBB : MF)
+ for (MachineInstr &MI : make_early_inc_range(MBB))
+ if (MI.isRegSequence() && foldPhysRegSequence(TII, TRI, MRI, MI)) {
+ Folded = true;
+ NeedRecomputeLiveIns = true;
+ }
+ }
+
// Cross-BB hard pins introduce physical registers that are live across basic
// block boundaries; recompute physreg live-in lists so the verifier and the
// allocator see correct liveness.
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
index e5e162ba444f0..cebc74f38ca12 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
@@ -45,10 +45,15 @@ define amdgpu_kernel void @pin_two_load_tuple(ptr addrspace(1) %o, ptr addrspace
; A 32-byte load off a divergent address is selected as two dwordx4 loads whose
; lanes reach the REG_SEQUENCE as subregister slices of the wider load. Placing
; a lane on its own would strand the rest of its load, so each load is placed as
-; a whole onto its half of the pinned tuple.
+; a whole onto its half of the pinned tuple. The stores read the pinned tuple
+; directly: the REG_SEQUENCEs that reassemble each half are folded away rather
+; than materialised as a copy per lane.
; CHECK-LABEL: {{^}}pin_split_wide_load:
; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[304:307]*/
; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[300:303]*/
+; CHECK-NOT: v_mov
+; CHECK: global_store_b128 v{{[0-9]+}}, v[{{[0-9:]+}}] /*v[304:307]*/
+; CHECK: global_store_b128 v{{[0-9]+}}, v[{{[0-9:]+}}] /*v[300:303]*/
; CHECK: .set .Lpin_split_wide_load.num_vgpr, 308
define amdgpu_kernel void @pin_split_wide_load(ptr addrspace(1) %in, ptr addrspace(1) %out) {
%tid = call i32 @llvm.amdgcn.workitem.id.x()
>From ab44fa23291a37df61380c1d02ec35905731d147 Mon Sep 17 00:00:00 2001
From: demonsan <1462264754 at qq.com>
Date: Mon, 10 Aug 2026 02:30:54 +0000
Subject: [PATCH 19/19] [AMDGPU] Let a pin take over a tuple whose occupant is
already dead
Any overlap with a tuple an earlier pin had claimed sent the later pin to the
soft path. That is too blunt for the way the attribute is lowered: Clang wraps
every store to a pinned variable in its own pin, so `x = load; x = f(x);`
arrives as two pins on one tuple, and only the first was honoured. The
variable's later value -- the one the kernel goes on to use -- ended up
wherever the allocator put it, which is precisely what the user asked to avoid.
Sharing a tuple is safe when the values do not overlap, so decide that instead
of assuming the worst. Judge it per lane, because the halves of a tuple are
independent registers updated at different points, and within one block, so
ordering is an index comparison. A lane may be rewritten by the very
instruction that last reads it, since sources are read before the result is
written; anything touching a lane after its rewrite, an access from another
block, or a call clobbering the tuple rejects the takeover.
Reassigning a pinned <8 x float> now stays in v[300:307] end to end: the loads
fill the tuple and four v_pk_fma update it in place. Two values that really are
live at once still get separate registers, and the soft-path reason for that is
now reported before the operand walk so an occupied tuple reads as
"overlaps a live earlier hard pin" rather than as whatever else the value's
shape trips over.
Co-authored-by: Cursor <cursoragent at cursor.com>
---
llvm/lib/Target/AMDGPU/SIPreColorPins.cpp | 112 ++++++++++++++++----
llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll | 43 ++++++++
2 files changed, 136 insertions(+), 19 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
index c82276753102a..978f2786c4142 100644
--- a/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
+++ b/llvm/lib/Target/AMDGPU/SIPreColorPins.cpp
@@ -180,15 +180,73 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
return false;
// Regunits already claimed by a hard pin. A later pin overlapping any claimed
- // unit falls back to soft, so two distinct live values never share a physreg.
- // Reuse by a single value (e.g. an accumulation chain) is instead absorbed
- // into the first pin's tie-connected component, making later pins on it
- // no-ops.
+ // unit may only take the tuple over once the earlier occupant is dead there
+ // (see canTakeOverClaim); otherwise it falls back to soft, so two live values
+ // never share a physreg. Reuse by a single value (e.g. an accumulation chain)
+ // is instead absorbed into the first pin's tie-connected component, making
+ // later pins on it no-ops.
DenseSet<MCRegUnit> Claimed;
bool NeedRecomputeLiveIns = false;
bool AnyHardPin = false;
unsigned ReqVGPRs =
0; // highest VGPR a pin needs, +1 (drives the occupancy cap)
+
+ // Whether a new value, defined by `Defs` as (instruction, physical target)
+ // pairs covering `PR`, can take `PR` over from whatever holds it now.
+ //
+ // Clang wraps every store to a pinned variable in its own pin, so
+ // `x = load; x = f(x);` arrives as two pins on one tuple. Refusing the second
+ // would leave the variable's later value wherever the allocator likes, which
+ // defeats the point of pinning it. The takeover is safe exactly when no lane
+ // of PR is touched after the instruction that rewrites it: judge that per
+ // lane, since each half of a tuple can be updated at a different point, and
+ // within one block, so instruction order is a plain index comparison. A lane
+ // may be rewritten by the very instruction that last reads it -- an in-place
+ // update reads its sources before writing its result.
+ auto canTakeOverClaim =
+ [&](MCRegister PR, MachineBasicBlock *MBB,
+ ArrayRef<std::pair<MachineInstr *, MCRegister>> Defs) {
+ DenseMap<MachineInstr *, unsigned> Order;
+ for (MachineInstr &MI : *MBB)
+ Order.insert({&MI, Order.size()});
+
+ DenseMap<MCRegUnit, unsigned> DefAt;
+ for (auto [DefMI, Phys] : Defs) {
+ if (!DefMI || DefMI->getParent() != MBB)
+ return false;
+ for (MCRegUnit U : TRI->regunits(Phys)) {
+ auto [It, New] = DefAt.try_emplace(U, Order[DefMI]);
+ if (!New)
+ It->second = std::min(It->second, Order[DefMI]);
+ }
+ }
+ // A lane the new value never writes would keep the old one alive under
+ // it, with no def to order the accesses against.
+ for (MCRegUnit U : TRI->regunits(PR))
+ if (!DefAt.contains(U))
+ return false;
+
+ for (MachineBasicBlock &B : MF)
+ for (MachineInstr &MI : B)
+ for (const MachineOperand &MO : MI.operands()) {
+ if (MO.isRegMask() && MO.clobbersPhysReg(PR))
+ return false;
+ if (!MO.isReg() || !MO.getReg().isPhysical())
+ continue;
+ for (MCRegUnit U : TRI->regunits(MO.getReg().asMCReg())) {
+ auto It = DefAt.find(U);
+ if (It == DefAt.end())
+ continue;
+ if (&B != MBB)
+ return false;
+ unsigned At = Order[&MI];
+ if (At > It->second || (At == It->second && MO.isDef()))
+ return false;
+ }
+ }
+ return true;
+ };
+
for (MachineInstr *Pin : Pins) {
assert(Pin->getNumExplicitOperands() == 3 &&
"pin pseudo must be (dst, src, regno)");
@@ -410,10 +468,6 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
WL.push_back(D.getReg());
}
}
- for (MCRegUnit U : TRI->regunits(PR))
- if (Ok && Claimed.contains(U))
- Ok = false;
-
// Map each element's defining register onto a physical (sub)register of
// PR. An element either covers its REG_SEQUENCE slot outright, or is a
// subregister slice of a wider def: a 32-byte load, for instance, is
@@ -485,6 +539,16 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
break;
}
+ // Each element's def is where its share of the tuple is written, which
+ // is what decides whether an earlier occupant can be displaced.
+ if (Ok && any_of(TRI->regunits(PR),
+ [&](MCRegUnit U) { return Claimed.contains(U); })) {
+ SmallVector<std::pair<MachineInstr *, MCRegister>, 16> Defs;
+ for (auto [Elem, Phys] : Elems)
+ Defs.emplace_back(MRI.getVRegDef(Elem), Phys);
+ Ok = canTakeOverClaim(PR, PinMBB, Defs);
+ }
+
if (Ok) {
// Point each element's def/uses at its physical (sub)register.
for (auto [Elem, Phys] : Elems) {
@@ -562,6 +626,27 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
}
+ // Conflict with an existing hard pin on overlapping regunits? Checked
+ // before the operand walk below, so an occupied tuple is reported as such
+ // rather than as whatever else the value's shape happens to trip over.
+ if (Hard && any_of(TRI->regunits(PR),
+ [&](MCRegUnit U) { return Claimed.contains(U); })) {
+ SmallVector<std::pair<MachineInstr *, MCRegister>, 8> Defs;
+ for (Register R : Comp)
+ for (MachineInstr &D : MRI.def_instructions(R)) {
+ if (&D == Pin)
+ continue; // erased below; Dst is written by the component's def
+ for (const MachineOperand &MO : D.defs())
+ if (MO.isReg() && MO.getReg() == R)
+ Defs.emplace_back(
+ &D, MO.getSubReg() ? TRI->getSubReg(PR, MO.getSubReg()) : PR);
+ }
+ if (!canTakeOverClaim(PR, Pin->getParent(), Defs)) {
+ Hard = false;
+ SoftWhy = "overlaps a live earlier hard pin";
+ }
+ }
+
// Collect and validate every operand referencing a component register.
SmallVector<MachineOperand *, 16> ToRewrite;
if (Hard) {
@@ -628,17 +713,6 @@ bool SIPreColorPins::runOnMachineFunction(MachineFunction &MF) {
}
}
- // Conflict with an existing hard pin on overlapping regunits?
- if (Hard) {
- for (MCRegUnit U : TRI->regunits(PR)) {
- if (Claimed.contains(U)) {
- Hard = false;
- SoftWhy = "overlaps an earlier hard pin";
- break;
- }
- }
- }
-
// Partition operands. Non-tied *subregister uses* (e.g. the per-lane reads
// a wide accumulator feeds into stores) are not rewritten to physical
// subregisters -- that yields fragile physical-subreg live ranges. Instead
diff --git a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
index cebc74f38ca12..8090902917560 100644
--- a/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
+++ b/llvm/test/CodeGen/AMDGPU/pin-reg-gfx1250.ll
@@ -66,6 +66,49 @@ define amdgpu_kernel void @pin_split_wide_load(ptr addrspace(1) %in, ptr addrspa
ret void
}
+; Clang pins every store to a pinned variable, so reassigning one yields two
+; pins on the same tuple. The second takes the tuple over -- the update is
+; in-place, reading and writing the same registers -- instead of leaving the
+; variable's later value wherever the allocator puts it.
+; CHECK-LABEL: {{^}}pin_reassigned_variable:
+; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[304:307]*/
+; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[300:303]*/
+; CHECK: v_pk_fma_f32 v[{{[0-9:]+}}] /*v[306:307]*/, {{.*}}v[{{[0-9:]+}}] /*v[306:307]*/
+; CHECK: v_pk_fma_f32 v[{{[0-9:]+}}] /*v[300:301]*/, {{.*}}v[{{[0-9:]+}}] /*v[300:301]*/
+; CHECK: .set .Lpin_reassigned_variable.num_vgpr, 308
+define amdgpu_kernel void @pin_reassigned_variable(ptr addrspace(1) %p) {
+ %tid = call i32 @llvm.amdgcn.workitem.id.x()
+ %idx = sext i32 %tid to i64
+ %a = getelementptr inbounds <8 x float>, ptr addrspace(1) %p, i64 %idx
+ %v = load <8 x float>, ptr addrspace(1) %a, align 32
+ %v.i = bitcast <8 x float> %v to <8 x i32>
+ %p1 = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> %v.i, i32 300)
+ %f = bitcast <8 x i32> %p1 to <8 x float>
+ %m = fmul contract <8 x float> %f, splat (float 3.000000e+00)
+ %s = fadd contract <8 x float> %m, splat (float -1.000000e+00)
+ %s.i = bitcast <8 x float> %s to <8 x i32>
+ %p2 = call <8 x i32> @llvm.amdgcn.pin.vgpr.v8i32(<8 x i32> %s.i, i32 300)
+ store <8 x i32> %p2, ptr addrspace(1) %a, align 32
+ ret void
+}
+
+; Two distinct values whose live ranges overlap must not share the tuple, even
+; though both ask for it: the second one is still loaded when the first is read.
+; CHECK-LABEL: {{^}}pin_two_live_values:
+; CHECK: global_load_b128 v[{{[0-9:]+}}] /*v[300:303]*/
+; CHECK: global_load_b128 v[0:3],
+; CHECK: global_store_b128 v{{[0-9]+}}, v[{{[0-9:]+}}] /*v[300:303]*/
+; CHECK: global_store_b128 v{{[0-9]+}}, v[0:3],
+define amdgpu_kernel void @pin_two_live_values(ptr addrspace(1) %p, ptr addrspace(1) %q) {
+ %a = load volatile <4 x float>, ptr addrspace(1) %p
+ %pa = call <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float> %a, i32 300)
+ %b = load volatile <4 x float>, ptr addrspace(1) %q
+ %pb = call <4 x float> @llvm.amdgcn.pin.vgpr.v4f32(<4 x float> %b, i32 300)
+ store volatile <4 x float> %pa, ptr addrspace(1) %p
+ store volatile <4 x float> %pb, ptr addrspace(1) %q
+ ret void
+}
+
; A pinned value defined outside a loop and carried into it reaches a PHI. The
; physical tuple must not be substituted into the PHI operand: LiveVariables
; walks PHI sources through getVarInfo(), which asserts on a physical register,
More information about the llvm-commits
mailing list