[llvm] [LLVM][AutoUpgrade] Support default args on undeclared multi-call upgrades (PR #216246)
Dharuni R Acharya via llvm-commits
llvm-commits at lists.llvm.org
Mon Aug 24 22:49:04 PDT 2026
https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/216246
>From 313cbb5bd5f3fa9f9f0e642d2478074aeca97d62 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Fri, 14 Aug 2026 04:58:08 +0000
Subject: [PATCH 1/6] [LLVM][Auto-Upgrade] Support default args on overloaded
intrinsics and undeclared multi-call upgrades
This patch extends intrinsic DefaultValue auto-upgrade to overloaded intrinsics and to calls that omit an explicit declare.
Multiple undeclared calls can leave uniquified temps; this patch recovers the intrinsic via prefix matching in upgradeIntrinsicWithDefaultArgs,
then validate the partial signature before filling default ImmArgs.
Signed-off-by: DharuniRAcharya <dharunira at nvidia.com>
---
llvm/include/llvm/IR/Intrinsics.h | 8 ++
llvm/lib/IR/AutoUpgrade.cpp | 78 +++++++++++++++----
llvm/lib/IR/Intrinsics.cpp | 39 ++++++++--
llvm/test/TableGen/intrinsic-default-args.td | 25 +++++-
.../TableGen/Basic/CodeGenIntrinsics.cpp | 10 ---
5 files changed, 124 insertions(+), 36 deletions(-)
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index a4799751832bf..c63a04ce06565 100644
--- a/llvm/include/llvm/IR/Intrinsics.h
+++ b/llvm/include/llvm/IR/Intrinsics.h
@@ -308,6 +308,14 @@ LLVM_ABI bool isSignatureValid(Function *F,
SmallVectorImpl<Type *> &OverloadTys,
raw_ostream &OS = nulls());
+/// Same as previous, but \p FT may omit exactly \p NumMissingTrailingParams
+/// trailing parameters. The omitted parameters must have concrete integer type
+/// so that all overload types can be resolved from the provided signature.
+LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT,
+ SmallVectorImpl<Type *> &OverloadTys,
+ unsigned NumMissingTrailingParams,
+ raw_ostream &OS = nulls());
+
// Checks if the intrinsic name matches with its signature and if not
// returns the declaration with the same signature and remangled name.
// An existing GlobalValue with the wanted name but with a wrong prototype
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 67f6e55eb500f..86f9fdbbe5cb0 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1467,25 +1467,14 @@ static bool convertIntrinsicValidType(StringRef Name,
return false;
}
-static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn) {
- Intrinsic::ID IID = Intrinsic::lookupIntrinsicID(F->getName());
- if (IID == Intrinsic::not_intrinsic)
- return false;
-
+static bool getDefaultArgUpgradeInfo(
+ Function *F, Intrinsic::ID IID, SmallVectorImpl<Type *> &OverloadTys) {
auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
if (Defaults.empty())
return false;
- // Overloaded intrinsics are out of scope for the default-arg feature
- // and will be supported in a follow-up.
- if (Intrinsic::isOverloaded(IID))
- return false;
-
- // Get the canonical full declaration for this intrinsic.
- Function *FullDecl = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
-
- // If the existing declaration already has all args, nothing to upgrade
- if (F->arg_size() >= FullDecl->arg_size())
+ unsigned FullArgCount = FirstDefault + Defaults.size();
+ if (F->arg_size() >= FullArgCount)
return false;
// Defaults are a contiguous trailing block, so checking the first missing
@@ -1493,7 +1482,62 @@ static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn) {
if (F->arg_size() < FirstDefault)
return false;
- NewFn = FullDecl;
+ unsigned NumMissingTrailingParams = FullArgCount - F->arg_size();
+ if (!Intrinsic::isSignatureValid(IID, F->getFunctionType(), OverloadTys,
+ NumMissingTrailingParams))
+ return false;
+
+ return true;
+}
+
+static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn) {
+ Intrinsic::ID IID = F->getIntrinsicID();
+ SmallVector<Type *, 4> OverloadTys;
+
+ if (IID != Intrinsic::not_intrinsic) {
+ if (!getDefaultArgUpgradeInfo(F, IID, OverloadTys))
+ return false;
+ } else {
+ Function *BestMatch = nullptr;
+ SmallVector<Type *, 4> BestOverloadTys;
+ for (Function &Candidate : *F->getParent()) {
+ Intrinsic::ID CandidateIID = Candidate.getIntrinsicID();
+ if (CandidateIID == Intrinsic::not_intrinsic)
+ continue;
+
+ StringRef CandidateName = Candidate.getName();
+ StringRef Suffix = F->getName();
+ if (!Suffix.consume_front(CandidateName) ||
+ !Suffix.consume_front(".") || Suffix.empty())
+ continue;
+
+ unsigned UniqueID;
+ if (Suffix.getAsInteger(10, UniqueID))
+ continue;
+
+ SmallVector<Type *, 4> CandidateOverloadTys;
+ if (!getDefaultArgUpgradeInfo(F, CandidateIID, CandidateOverloadTys))
+ continue;
+
+ if (!BestMatch ||
+ CandidateName.size() > BestMatch->getName().size()) {
+ BestMatch = &Candidate;
+ IID = CandidateIID;
+ BestOverloadTys = std::move(CandidateOverloadTys);
+ }
+ }
+
+ if (!BestMatch)
+ return false;
+ OverloadTys = std::move(BestOverloadTys);
+ }
+
+ auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
+ unsigned FullArgCount = FirstDefault + Defaults.size();
+ NewFn =
+ Intrinsic::getOrInsertDeclaration(F->getParent(), IID, OverloadTys);
+ assert(NewFn->arg_size() == FullArgCount &&
+ "default argument table does not match intrinsic signature");
return true;
}
@@ -2179,7 +2223,7 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
// to both detect an intrinsic which needs upgrading, and to provide the
// upgraded form of the intrinsic. We should perhaps have two separate
// functions for this.
- if (upgradeIntrinsicDeclWithDefaultArgs(F, NewFn))
+ if (upgradeIntrinsicWithDefaultArgs(F, NewFn))
return true;
return false;
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index 266af8e06a230..f976f05fe9b80 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -44,7 +44,8 @@ static bool isSignatureValid(FunctionType *FTy,
ArrayRef<Intrinsic::IITDescriptor> &Infos,
unsigned NumArgs, bool IsVarArg,
SmallVectorImpl<Type *> &OverloadTys,
- raw_ostream &OS);
+ raw_ostream &OS,
+ unsigned NumMissingTrailingParams = 0);
/// Table of string intrinsic names indexed by enum value.
#define GET_INTRINSIC_NAME_TABLE
@@ -1333,13 +1334,18 @@ matchIntrinsicType(Type *Ty, ArrayRef<Intrinsic::IITDescriptor> &Infos,
/// \p IsVarArg. The overloaded types for the intrinsic are pushed to the
/// \p OverloadTys vector.
///
+/// If \p NumMissingTrailingParams is non-zero, \p FTy may omit exactly that
+/// many trailing parameters. Omitted parameters must have concrete integer
+/// types and therefore cannot contribute an unresolved overload type.
+///
/// If the type is not valid, returns false and prints an error message to
/// \p OS.
static bool isSignatureValid(FunctionType *FTy,
ArrayRef<Intrinsic::IITDescriptor> &Infos,
unsigned NumArgs, bool IsVarArg,
SmallVectorImpl<Type *> &OverloadTys,
- raw_ostream &OS) {
+ raw_ostream &OS,
+ unsigned NumMissingTrailingParams) {
SmallVector<DeferredIntrinsicMatchInfo, 2> DeferredChecks;
assert(!Infos.empty() && "Table consistency error");
@@ -1352,9 +1358,10 @@ static bool isSignatureValid(FunctionType *FTy,
DeferredChecks, false, OS))
return false;
- if (FTy->getNumParams() != NumArgs) {
+ unsigned ProvidedArgs = FTy->getNumParams();
+ if (ProvidedArgs + NumMissingTrailingParams != NumArgs) {
OS << "intrinsic has incorrect number of args. Expected " << NumArgs
- << ", but got " << FTy->getNumParams();
+ << ", but got " << ProvidedArgs;
return false;
}
@@ -1373,6 +1380,19 @@ static bool isSignatureValid(FunctionType *FTy,
return false;
}
+ if (NumMissingTrailingParams) {
+ // Default arguments are materialized as ConstantInt values, requiring one
+ // concrete integer descriptor per omitted parameter.
+ if (Infos.size() != NumMissingTrailingParams ||
+ llvm::any_of(Infos, [](Intrinsic::IITDescriptor D) {
+ return D.Kind != Intrinsic::IITDescriptor::Integer;
+ })) {
+ OS << "intrinsic has unresolved trailing argument types!";
+ return false;
+ }
+ Infos = {};
+ }
+
if (!Infos.empty()) {
OS << "intrinsic has too few arguments!";
return false;
@@ -1399,13 +1419,22 @@ bool Intrinsic::hasStructReturnType(ID id) {
bool Intrinsic::isSignatureValid(Intrinsic::ID ID, FunctionType *FT,
SmallVectorImpl<Type *> &OverloadTys,
raw_ostream &OS) {
+ return isSignatureValid(ID, FT, OverloadTys,
+ /*NumMissingTrailingParams=*/0, OS);
+}
+
+bool Intrinsic::isSignatureValid(Intrinsic::ID ID, FunctionType *FT,
+ SmallVectorImpl<Type *> &OverloadTys,
+ unsigned NumMissingTrailingParams,
+ raw_ostream &OS) {
if (!ID)
return false;
SmallVector<Intrinsic::IITDescriptor, 8> Table;
auto [TableRef, NumArgs, IsVarArg] = getIntrinsicInfoTableEntries(ID, Table);
- return ::isSignatureValid(FT, TableRef, NumArgs, IsVarArg, OverloadTys, OS);
+ return ::isSignatureValid(FT, TableRef, NumArgs, IsVarArg, OverloadTys, OS,
+ NumMissingTrailingParams);
}
bool Intrinsic::isSignatureValid(Function *F,
diff --git a/llvm/test/TableGen/intrinsic-default-args.td b/llvm/test/TableGen/intrinsic-default-args.td
index e372a870880e7..1afffc964360b 100644
--- a/llvm/test/TableGen/intrinsic-default-args.td
+++ b/llvm/test/TableGen/intrinsic-default-args.td
@@ -6,7 +6,7 @@
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_NEGATIVE 2>&1 | FileCheck %s --check-prefix=ERR-NEGATIVE
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_RANGE 2>&1 | FileCheck %s --check-prefix=ERR-RANGE
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_NONINT 2>&1 | FileCheck %s --check-prefix=ERR-NONINT
-// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_OVERLOADED 2>&1 | FileCheck %s --check-prefix=ERR-OVERLOADED
+// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DTEST_OVERLOADED | FileCheck %s --check-prefix=OVERLOADED
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_GAP 2>&1 | FileCheck %s --check-prefix=ERR-GAP
include "llvm/IR/Intrinsics.td"
@@ -77,12 +77,29 @@ def int_test_nonint :
[ImmArg<ArgIndex<1>, DefaultValue<1>>]>;
#endif
-// A default on an overloaded intrinsic is rejected (not yet supported).
-#ifdef ERROR_OVERLOADED
-// ERR-OVERLOADED: error: default argument values are not supported for overloaded intrinsics
+#ifdef TEST_OVERLOADED
def int_test_overloaded :
Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, llvm_i32_ty],
[ImmArg<ArgIndex<1>, DefaultValue<5>>]>;
+
+def int_test_overloaded_two_defaults :
+ Intrinsic<[llvm_anyint_ty],
+ [LLVMMatchType<0>, llvm_i1_ty, llvm_i32_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<0>>,
+ ImmArg<ArgIndex<2>, DefaultValue<7>>]>;
+
+def int_test_overloaded_ptr :
+ Intrinsic<[llvm_anyptr_ty], [LLVMMatchType<0>, llvm_i32_ty],
+ [ImmArg<ArgIndex<1>, DefaultValue<1>>]>;
+
+// Header (1 << 32) | 1 = 4294967297 for single trailing defaults.
+// Header (2 << 32) | 1 = 8589934593 for two trailing defaults.
+// OVERLOADED: static constexpr uint64_t DefaultArgValuesTable[] = {
+// OVERLOADED-NEXT: 0, // offset 0: sentinel for intrinsics without defaults
+// OVERLOADED-DAG: {{.*}}4294967297,{{.*}}5,{{.*}}0,
+// OVERLOADED-DAG: {{.*}}8589934593,{{.*}}0,{{.*}}7,{{.*}}0,
+// OVERLOADED-DAG: {{.*}}4294967297,{{.*}}1,{{.*}}0,
+// OVERLOADED: Intrinsic::getAllDefaultArgValues(ID IID) {
#endif
// Defaults must form a contiguous trailing block.
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 76fc2a4f38811..7bcfcc6d0f87f 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -388,16 +388,6 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
for (auto &Attrs : ArgumentAttributes)
llvm::sort(Attrs);
- // Default values are not yet supported for overloaded intrinsics
- // (overloaded support will come in a follow-up).
- if (isOverloaded &&
- llvm::any_of(ParamDefaultValues, [](const std::optional<uint64_t> &DV) {
- return DV.has_value();
- }))
- PrintFatalError(TheDef->getLoc(),
- "default argument values are not supported for "
- "overloaded intrinsics");
-
// Validate: defaults must form a contiguous trailing block ending at
// the last parameter (mirrors C++ default-argument rules).
unsigned NumParams = IS.ParamTys.size();
>From 4746c3467b6ab2ef5a95c196c844c5e9efddbcc7 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Fri, 14 Aug 2026 05:03:51 +0000
Subject: [PATCH 2/6] Fix formatting
---
llvm/lib/IR/AutoUpgrade.cpp | 14 ++++++--------
1 file changed, 6 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 86f9fdbbe5cb0..03979cf4047fa 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1467,8 +1467,8 @@ static bool convertIntrinsicValidType(StringRef Name,
return false;
}
-static bool getDefaultArgUpgradeInfo(
- Function *F, Intrinsic::ID IID, SmallVectorImpl<Type *> &OverloadTys) {
+static bool getDefaultArgUpgradeInfo(Function *F, Intrinsic::ID IID,
+ SmallVectorImpl<Type *> &OverloadTys) {
auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
if (Defaults.empty())
return false;
@@ -1507,8 +1507,8 @@ static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn) {
StringRef CandidateName = Candidate.getName();
StringRef Suffix = F->getName();
- if (!Suffix.consume_front(CandidateName) ||
- !Suffix.consume_front(".") || Suffix.empty())
+ if (!Suffix.consume_front(CandidateName) || !Suffix.consume_front(".") ||
+ Suffix.empty())
continue;
unsigned UniqueID;
@@ -1519,8 +1519,7 @@ static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn) {
if (!getDefaultArgUpgradeInfo(F, CandidateIID, CandidateOverloadTys))
continue;
- if (!BestMatch ||
- CandidateName.size() > BestMatch->getName().size()) {
+ if (!BestMatch || CandidateName.size() > BestMatch->getName().size()) {
BestMatch = &Candidate;
IID = CandidateIID;
BestOverloadTys = std::move(CandidateOverloadTys);
@@ -1534,8 +1533,7 @@ static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn) {
auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
unsigned FullArgCount = FirstDefault + Defaults.size();
- NewFn =
- Intrinsic::getOrInsertDeclaration(F->getParent(), IID, OverloadTys);
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID, OverloadTys);
assert(NewFn->arg_size() == FullArgCount &&
"default argument table does not match intrinsic signature");
return true;
>From 2fbd3f46a7e146001a8b267fe507ca42ae89ef0a Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Thu, 20 Aug 2026 10:48:27 +0000
Subject: [PATCH 3/6] Address comments
---
llvm/include/llvm/IR/Intrinsics.h | 8 --
llvm/include/llvm/IR/Intrinsics.td | 4 +-
llvm/lib/IR/AutoUpgrade.cpp | 87 +++++--------------
llvm/lib/IR/Intrinsics.cpp | 39 ++-------
llvm/test/TableGen/intrinsic-default-args.td | 25 +-----
.../TableGen/Basic/CodeGenIntrinsics.cpp | 16 ++--
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 21 +----
7 files changed, 41 insertions(+), 159 deletions(-)
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index c63a04ce06565..a4799751832bf 100644
--- a/llvm/include/llvm/IR/Intrinsics.h
+++ b/llvm/include/llvm/IR/Intrinsics.h
@@ -308,14 +308,6 @@ LLVM_ABI bool isSignatureValid(Function *F,
SmallVectorImpl<Type *> &OverloadTys,
raw_ostream &OS = nulls());
-/// Same as previous, but \p FT may omit exactly \p NumMissingTrailingParams
-/// trailing parameters. The omitted parameters must have concrete integer type
-/// so that all overload types can be resolved from the provided signature.
-LLVM_ABI bool isSignatureValid(Intrinsic::ID ID, FunctionType *FT,
- SmallVectorImpl<Type *> &OverloadTys,
- unsigned NumMissingTrailingParams,
- raw_ostream &OS = nulls());
-
// Checks if the intrinsic name matches with its signature and if not
// returns the declaration with the same signature and remangled name.
// An existing GlobalValue with the wanted name but with a wrong prototype
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 083b933e77d16..1d97eaa26bc64 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -136,9 +136,7 @@ class Returned<ArgIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
-// DefaultValue - A value carrier paired with ImmArg to declare a default for
-// a missing trailing argument. AutoUpgrade fills the default when an old
-// .bc / .ll file is loaded.
+// Default value for a trailing ImmArg, materialized by AutoUpgrade.
class DefaultValue<int val> {
int Value = val;
}
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 03979cf4047fa..b03826fd6a57d 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1468,23 +1468,18 @@ static bool convertIntrinsicValidType(StringRef Name,
}
static bool getDefaultArgUpgradeInfo(Function *F, Intrinsic::ID IID,
- SmallVectorImpl<Type *> &OverloadTys) {
+ unsigned &FullArgCount) {
auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
if (Defaults.empty())
return false;
- unsigned FullArgCount = FirstDefault + Defaults.size();
- if (F->arg_size() >= FullArgCount)
+ if (Intrinsic::isOverloaded(IID))
return false;
- // Defaults are a contiguous trailing block, so checking the first missing
- // argument is enough.
- if (F->arg_size() < FirstDefault)
- return false;
+ FullArgCount = FirstDefault + Defaults.size();
- unsigned NumMissingTrailingParams = FullArgCount - F->arg_size();
- if (!Intrinsic::isSignatureValid(IID, F->getFunctionType(), OverloadTys,
- NumMissingTrailingParams))
+ // Only trailing default arguments can be missing.
+ if (F->arg_size() < FirstDefault || F->arg_size() >= FullArgCount)
return false;
return true;
@@ -1492,48 +1487,13 @@ static bool getDefaultArgUpgradeInfo(Function *F, Intrinsic::ID IID,
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn) {
Intrinsic::ID IID = F->getIntrinsicID();
- SmallVector<Type *, 4> OverloadTys;
-
- if (IID != Intrinsic::not_intrinsic) {
- if (!getDefaultArgUpgradeInfo(F, IID, OverloadTys))
- return false;
- } else {
- Function *BestMatch = nullptr;
- SmallVector<Type *, 4> BestOverloadTys;
- for (Function &Candidate : *F->getParent()) {
- Intrinsic::ID CandidateIID = Candidate.getIntrinsicID();
- if (CandidateIID == Intrinsic::not_intrinsic)
- continue;
-
- StringRef CandidateName = Candidate.getName();
- StringRef Suffix = F->getName();
- if (!Suffix.consume_front(CandidateName) || !Suffix.consume_front(".") ||
- Suffix.empty())
- continue;
-
- unsigned UniqueID;
- if (Suffix.getAsInteger(10, UniqueID))
- continue;
-
- SmallVector<Type *, 4> CandidateOverloadTys;
- if (!getDefaultArgUpgradeInfo(F, CandidateIID, CandidateOverloadTys))
- continue;
-
- if (!BestMatch || CandidateName.size() > BestMatch->getName().size()) {
- BestMatch = &Candidate;
- IID = CandidateIID;
- BestOverloadTys = std::move(CandidateOverloadTys);
- }
- }
- if (!BestMatch)
- return false;
- OverloadTys = std::move(BestOverloadTys);
- }
+ unsigned FullArgCount;
+ if (!getDefaultArgUpgradeInfo(F, IID, FullArgCount))
+ return false;
- auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
- unsigned FullArgCount = FirstDefault + Defaults.size();
- NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID, OverloadTys);
+ rename(F);
+ NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
assert(NewFn->arg_size() == FullArgCount &&
"default argument table does not match intrinsic signature");
return true;
@@ -2217,12 +2177,13 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
+ if (upgradeIntrinsicWithDefaultArgs(F, NewFn))
+ return true;
+
// This may not belong here. This function is effectively being overloaded
// to both detect an intrinsic which needs upgrading, and to provide the
// upgraded form of the intrinsic. We should perhaps have two separate
// functions for this.
- if (upgradeIntrinsicWithDefaultArgs(F, NewFn))
- return true;
return false;
}
@@ -5436,33 +5397,32 @@ static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
unsigned OldArgCount = CI->arg_size();
unsigned NewArgCount = NewFn->arg_size();
- // If the caller already supplied all arguments (or more), nothing to do.
- // This mirrors C++ semantics: an explicitly-passed value is never overridden.
- if (OldArgCount >= NewArgCount)
+ if (OldArgCount > NewArgCount)
return false;
- // Start with the existing arguments from the old call.
- SmallVector<Value *, 8> NewArgs(CI->args());
+ if (OldArgCount == NewArgCount) {
+ if (CI->getFunctionType() != NewFn->getFunctionType())
+ return false;
+ CI->setCalledFunction(NewFn);
+ return true;
+ }
- // Defaults are a contiguous trailing block, so checking the first missing
- // argument is enough.
if (OldArgCount < FirstDefault)
return false;
- // Fill in each missing trailing argument from the table.
+ SmallVector<Value *, 8> NewArgs(CI->args());
+
FunctionType *NewFT = NewFn->getFunctionType();
for (unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
assert(Idx >= FirstDefault && Idx - FirstDefault < Defaults.size() &&
"missing argument outside the default range");
Type *ParamTy = NewFT->getParamType(Idx);
- // Only integer types are supported (i1, i8, i16, i32, i64).
if (!ParamTy->isIntegerTy())
return false;
NewArgs.push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
}
- // Preserve operand bundles by creating the call with them.
SmallVector<OperandBundleDef, 1> OpBundles;
CI->getOperandBundlesAsDefs(OpBundles);
CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
@@ -5588,9 +5548,6 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
CallInst *NewCall = nullptr;
switch (NewFn->getIntrinsicID()) {
default: {
- // Last resort: try the data-driven default-arg upgrade.
- // Handles any intrinsic annotated with ImmArg<..., DefaultValue<...>>
- // in its .td definition, without needing a dedicated case.
if (upgradeIntrinsicCallWithDefaultArgs(CI, NewFn, Builder))
return;
DefaultCase();
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index f976f05fe9b80..266af8e06a230 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -44,8 +44,7 @@ static bool isSignatureValid(FunctionType *FTy,
ArrayRef<Intrinsic::IITDescriptor> &Infos,
unsigned NumArgs, bool IsVarArg,
SmallVectorImpl<Type *> &OverloadTys,
- raw_ostream &OS,
- unsigned NumMissingTrailingParams = 0);
+ raw_ostream &OS);
/// Table of string intrinsic names indexed by enum value.
#define GET_INTRINSIC_NAME_TABLE
@@ -1334,18 +1333,13 @@ matchIntrinsicType(Type *Ty, ArrayRef<Intrinsic::IITDescriptor> &Infos,
/// \p IsVarArg. The overloaded types for the intrinsic are pushed to the
/// \p OverloadTys vector.
///
-/// If \p NumMissingTrailingParams is non-zero, \p FTy may omit exactly that
-/// many trailing parameters. Omitted parameters must have concrete integer
-/// types and therefore cannot contribute an unresolved overload type.
-///
/// If the type is not valid, returns false and prints an error message to
/// \p OS.
static bool isSignatureValid(FunctionType *FTy,
ArrayRef<Intrinsic::IITDescriptor> &Infos,
unsigned NumArgs, bool IsVarArg,
SmallVectorImpl<Type *> &OverloadTys,
- raw_ostream &OS,
- unsigned NumMissingTrailingParams) {
+ raw_ostream &OS) {
SmallVector<DeferredIntrinsicMatchInfo, 2> DeferredChecks;
assert(!Infos.empty() && "Table consistency error");
@@ -1358,10 +1352,9 @@ static bool isSignatureValid(FunctionType *FTy,
DeferredChecks, false, OS))
return false;
- unsigned ProvidedArgs = FTy->getNumParams();
- if (ProvidedArgs + NumMissingTrailingParams != NumArgs) {
+ if (FTy->getNumParams() != NumArgs) {
OS << "intrinsic has incorrect number of args. Expected " << NumArgs
- << ", but got " << ProvidedArgs;
+ << ", but got " << FTy->getNumParams();
return false;
}
@@ -1380,19 +1373,6 @@ static bool isSignatureValid(FunctionType *FTy,
return false;
}
- if (NumMissingTrailingParams) {
- // Default arguments are materialized as ConstantInt values, requiring one
- // concrete integer descriptor per omitted parameter.
- if (Infos.size() != NumMissingTrailingParams ||
- llvm::any_of(Infos, [](Intrinsic::IITDescriptor D) {
- return D.Kind != Intrinsic::IITDescriptor::Integer;
- })) {
- OS << "intrinsic has unresolved trailing argument types!";
- return false;
- }
- Infos = {};
- }
-
if (!Infos.empty()) {
OS << "intrinsic has too few arguments!";
return false;
@@ -1419,22 +1399,13 @@ bool Intrinsic::hasStructReturnType(ID id) {
bool Intrinsic::isSignatureValid(Intrinsic::ID ID, FunctionType *FT,
SmallVectorImpl<Type *> &OverloadTys,
raw_ostream &OS) {
- return isSignatureValid(ID, FT, OverloadTys,
- /*NumMissingTrailingParams=*/0, OS);
-}
-
-bool Intrinsic::isSignatureValid(Intrinsic::ID ID, FunctionType *FT,
- SmallVectorImpl<Type *> &OverloadTys,
- unsigned NumMissingTrailingParams,
- raw_ostream &OS) {
if (!ID)
return false;
SmallVector<Intrinsic::IITDescriptor, 8> Table;
auto [TableRef, NumArgs, IsVarArg] = getIntrinsicInfoTableEntries(ID, Table);
- return ::isSignatureValid(FT, TableRef, NumArgs, IsVarArg, OverloadTys, OS,
- NumMissingTrailingParams);
+ return ::isSignatureValid(FT, TableRef, NumArgs, IsVarArg, OverloadTys, OS);
}
bool Intrinsic::isSignatureValid(Function *F,
diff --git a/llvm/test/TableGen/intrinsic-default-args.td b/llvm/test/TableGen/intrinsic-default-args.td
index 1afffc964360b..e372a870880e7 100644
--- a/llvm/test/TableGen/intrinsic-default-args.td
+++ b/llvm/test/TableGen/intrinsic-default-args.td
@@ -6,7 +6,7 @@
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_NEGATIVE 2>&1 | FileCheck %s --check-prefix=ERR-NEGATIVE
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_RANGE 2>&1 | FileCheck %s --check-prefix=ERR-RANGE
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_NONINT 2>&1 | FileCheck %s --check-prefix=ERR-NONINT
-// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DTEST_OVERLOADED | FileCheck %s --check-prefix=OVERLOADED
+// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_OVERLOADED 2>&1 | FileCheck %s --check-prefix=ERR-OVERLOADED
// RUN: not llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s -DTEST_INTRINSICS_SUPPRESS_DEFS -DERROR_GAP 2>&1 | FileCheck %s --check-prefix=ERR-GAP
include "llvm/IR/Intrinsics.td"
@@ -77,29 +77,12 @@ def int_test_nonint :
[ImmArg<ArgIndex<1>, DefaultValue<1>>]>;
#endif
-#ifdef TEST_OVERLOADED
+// A default on an overloaded intrinsic is rejected (not yet supported).
+#ifdef ERROR_OVERLOADED
+// ERR-OVERLOADED: error: default argument values are not supported for overloaded intrinsics
def int_test_overloaded :
Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, llvm_i32_ty],
[ImmArg<ArgIndex<1>, DefaultValue<5>>]>;
-
-def int_test_overloaded_two_defaults :
- Intrinsic<[llvm_anyint_ty],
- [LLVMMatchType<0>, llvm_i1_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<1>, DefaultValue<0>>,
- ImmArg<ArgIndex<2>, DefaultValue<7>>]>;
-
-def int_test_overloaded_ptr :
- Intrinsic<[llvm_anyptr_ty], [LLVMMatchType<0>, llvm_i32_ty],
- [ImmArg<ArgIndex<1>, DefaultValue<1>>]>;
-
-// Header (1 << 32) | 1 = 4294967297 for single trailing defaults.
-// Header (2 << 32) | 1 = 8589934593 for two trailing defaults.
-// OVERLOADED: static constexpr uint64_t DefaultArgValuesTable[] = {
-// OVERLOADED-NEXT: 0, // offset 0: sentinel for intrinsics without defaults
-// OVERLOADED-DAG: {{.*}}4294967297,{{.*}}5,{{.*}}0,
-// OVERLOADED-DAG: {{.*}}8589934593,{{.*}}0,{{.*}}7,{{.*}}0,
-// OVERLOADED-DAG: {{.*}}4294967297,{{.*}}1,{{.*}}0,
-// OVERLOADED: Intrinsic::getAllDefaultArgValues(ID IID) {
#endif
// Defaults must form a contiguous trailing block.
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 7bcfcc6d0f87f..b4d1264253af7 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -388,8 +388,14 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
for (auto &Attrs : ArgumentAttributes)
llvm::sort(Attrs);
- // Validate: defaults must form a contiguous trailing block ending at
- // the last parameter (mirrors C++ default-argument rules).
+ if (isOverloaded &&
+ llvm::any_of(ParamDefaultValues, [](const std::optional<uint64_t> &DV) {
+ return DV.has_value();
+ }))
+ PrintFatalError(TheDef->getLoc(),
+ "default argument values are not supported for "
+ "overloaded intrinsics");
+
unsigned NumParams = IS.ParamTys.size();
bool SeenDefault = false;
for (unsigned i = 0; i < NumParams; ++i) {
@@ -403,8 +409,6 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
}
}
- // Validate each declared default: the parameter must be an integer type and
- // the value (an unsigned bit pattern) must fit in the declared width.
for (unsigned i = 0; i < ParamDefaultValues.size(); ++i) {
if (!ParamDefaultValues[i].has_value())
continue;
@@ -528,14 +532,10 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
addArgAttribute(ArgNo, ImmArg);
- // If a DefaultValue (not the NoDefault sentinel) was supplied, record it.
- // NoDefault is recognized by its Value field being unset (?).
const Record *DefaultField = R->getValueAsDef("Default");
const RecordVal *ValueField = DefaultField->getValue("Value");
if (ValueField && !isa<UnsetInit>(ValueField->getValue())) {
int64_t Value = DefaultField->getValueAsInt("Value");
- // Defaults are stored as an unsigned bit pattern; a negative literal
- // would silently wrap, so reject it with a clear message.
if (Value < 0)
PrintFatalError(TheDef->getLoc(), "default argument value " +
Twine(Value) + " on parameter " +
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index 23086c422ec1c..8b1129910e2a4 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -933,21 +933,10 @@ void Intrinsic::printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS, const Cons
void IntrinsicEmitter::EmitDefaultArgValuesTable(
const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
- // Build the per-intrinsic default-value sequences:
- // [Header = (NumDefaults << 32) | FirstDefault, val0, val1, ...]
- // Each value is the (non-negative) default for one parameter, stored as a
- // uint64_t bit pattern.
- //
- // Offset 0 of the values table is reserved as the "no defaults" sentinel
- // (a single 0 word, decoding to NumDefaults = 0). Intrinsics without
- // defaults point to offset 0; real sequences are emitted after it.
- // SequenceToOffsetTable deduplicates the real sequences.
-
+ // Sequences are [NumDefaults << 32 | FirstDefault, values...].
using Sequence = SmallVector<uint64_t, 8>;
SequenceToOffsetTable<Sequence> Table;
- // An empty Sequence means "no defaults" (maps to the reserved offset 0);
- // otherwise it holds the intrinsic's value sequence.
SmallVector<Sequence> PerIntrinsic;
PerIntrinsic.reserve(Ints.size());
@@ -957,7 +946,6 @@ void IntrinsicEmitter::EmitDefaultArgValuesTable(
continue;
}
- // Find the first parameter with a default.
unsigned FirstDefault = 0;
for (size_t j = 0U, N = Int.ParamDefaultValues.size(); j < N; ++j) {
if (Int.ParamDefaultValues[j].has_value()) {
@@ -984,17 +972,11 @@ void IntrinsicEmitter::EmitDefaultArgValuesTable(
IfDefEmitter IfDef(OS, "GET_INTRINSIC_DEFAULT_ARG_VALUES");
- // Emit the flat values table. Offset 0 is the reserved "no defaults"
- // sentinel; the deduplicated real sequences follow it.
OS << "static constexpr uint64_t DefaultArgValuesTable[] = {\n";
OS << " 0, // offset 0: sentinel for intrinsics without defaults\n";
Table.emit(OS, [](raw_ostream &OS, uint64_t Val) { OS << " " << Val; });
OS << "};\n\n";
- // Emit the per-intrinsic offset table. Entry #0 is for the invalid
- // Intrinsic::not_intrinsic (IID 0); it and every intrinsic without defaults
- // point to the reserved sentinel at offset 0. Real sequences are shifted by
- // +1 to skip past the sentinel slot.
OS << "static constexpr uint32_t DefaultArgValuesTableOffset[] = {\n";
OS << " 0, // not_intrinsic\n";
for (const Sequence &Seq : PerIntrinsic) {
@@ -1005,7 +987,6 @@ void IntrinsicEmitter::EmitDefaultArgValuesTable(
}
OS << "};\n\n";
- // Emit the lookup function body.
OS << R"(
std::pair<unsigned, ArrayRef<uint64_t>>
Intrinsic::getAllDefaultArgValues(ID IID) {
>From e336322864e71daca23bfe54d6c64c2375d48a11 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Fri, 21 Aug 2026 05:58:18 +0000
Subject: [PATCH 4/6] Address comments
---
llvm/lib/IR/AutoUpgrade.cpp | 13 ++++++++++---
1 file changed, 10 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index b03826fd6a57d..d6a7435c72ea4 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -5397,9 +5397,15 @@ static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
unsigned OldArgCount = CI->arg_size();
unsigned NewArgCount = NewFn->arg_size();
+ if (OldArgCount < FirstDefault)
+ return false;
+
+ // More arguments than the new intrinsic accepts, cannot upgrade.
if (OldArgCount > NewArgCount)
return false;
+ // The call already passes the defaulted arguments explicitly; only the
+ // callee is still the old, shorter declaration, so retarget it.
if (OldArgCount == NewArgCount) {
if (CI->getFunctionType() != NewFn->getFunctionType())
return false;
@@ -5407,9 +5413,8 @@ static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
return true;
}
- if (OldArgCount < FirstDefault)
- return false;
-
+ // OldArgCount < NewArgCount: Fill in each missing trailing default
+ // argument from the table.
SmallVector<Value *, 8> NewArgs(CI->args());
FunctionType *NewFT = NewFn->getFunctionType();
@@ -5418,11 +5423,13 @@ static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
"missing argument outside the default range");
Type *ParamTy = NewFT->getParamType(Idx);
+ // Only integer types are supported (i1, i8, i16, i32, i64).
if (!ParamTy->isIntegerTy())
return false;
NewArgs.push_back(ConstantInt::get(ParamTy, Defaults[Idx - FirstDefault]));
}
+ // Preserve operand bundles by creating the call with them.
SmallVector<OperandBundleDef, 1> OpBundles;
CI->getOperandBundlesAsDefs(OpBundles);
CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
>From d1091938e02bd7808ae16f50ce5d09b8ee733ba5 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Fri, 21 Aug 2026 11:17:19 +0000
Subject: [PATCH 5/6] Address comments
---
llvm/lib/IR/AutoUpgrade.cpp | 20 ++++++++++----------
1 file changed, 10 insertions(+), 10 deletions(-)
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index d6a7435c72ea4..d619b9a78c0c1 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1467,35 +1467,35 @@ static bool convertIntrinsicValidType(StringRef Name,
return false;
}
-static bool getDefaultArgUpgradeInfo(Function *F, Intrinsic::ID IID,
- unsigned &FullArgCount) {
+static unsigned getFullArgCountForDefaultArgUpgrade(Function *F,
+ Intrinsic::ID IID) {
auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
if (Defaults.empty())
- return false;
+ return 0;
if (Intrinsic::isOverloaded(IID))
- return false;
+ return 0;
- FullArgCount = FirstDefault + Defaults.size();
+ unsigned FullArgCount = FirstDefault + Defaults.size();
// Only trailing default arguments can be missing.
if (F->arg_size() < FirstDefault || F->arg_size() >= FullArgCount)
- return false;
+ return 0;
- return true;
+ return FullArgCount;
}
static bool upgradeIntrinsicWithDefaultArgs(Function *F, Function *&NewFn) {
Intrinsic::ID IID = F->getIntrinsicID();
- unsigned FullArgCount;
- if (!getDefaultArgUpgradeInfo(F, IID, FullArgCount))
+ unsigned FullArgCount = getFullArgCountForDefaultArgUpgrade(F, IID);
+ if (!FullArgCount)
return false;
rename(F);
NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
assert(NewFn->arg_size() == FullArgCount &&
- "default argument table does not match intrinsic signature");
+ "total number of default args does not match intrinsic signature");
return true;
}
>From c641db92feba69dc9db18816fe806190f8f69210 Mon Sep 17 00:00:00 2001
From: DharuniRAcharya <dharunira at nvidia.com>
Date: Tue, 25 Aug 2026 05:48:36 +0000
Subject: [PATCH 6/6] Example usecase with collector_b flag as default argument
to tcgen05.mma intrinsics
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 14 ++--
llvm/lib/IR/AutoUpgrade.cpp | 88 --------------------------
2 files changed, 9 insertions(+), 93 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 4357ad367d269..a77801d3b86a9 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2916,13 +2916,16 @@ class DefaultAttrsIntrinsicFlags<list<LLVMType> ret_types,
list<LLVMType> param_types,
list<LLVMType> flags,
list<IntrinsicProperty> intr_properties,
- string name = "">
+ string name = "",
+ DefaultValue last_flag_default = NoDefault>
: DefaultAttrsIntrinsic<
ret_types,
!listconcat(param_types, flags),
!listconcat(intr_properties,
!foreach(i, !range(flags),
- ImmArg<ArgIndex<!add(i, !size(param_types))>>)),
+ ImmArg<ArgIndex<!add(i, !size(param_types))>,
+ !if(!eq(i, !sub(!size(flags), 1)),
+ last_flag_default, NoDefault)>)),
name>;
// TMA Tensor Copy Intrinsics: S2G -> From Shared to Global memory variants
@@ -3441,7 +3444,7 @@ foreach sparse = [0, 1] in {
def mma.record_name:
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
- mma.intr_name>;
+ mma.intr_name, DefaultValue<0>>;
} // ashift
} // scale_d
} // space
@@ -3485,7 +3488,8 @@ foreach sparse = [0, 1] in {
);
def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
- intrinsic_properties, mma.intr_name>;
+ intrinsic_properties, mma.intr_name,
+ DefaultValue<0>>;
} // ashift
} // scale_d
} // cta_group
@@ -3520,7 +3524,7 @@ foreach sparse = [0, 1] in {
ArgInfo<cta_group, [ArgName<"cta_group">]>,
ArgInfo<collector_usage_a, [ArgName<"collector_a">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>,
ArgInfo<collector_usage_b, [ArgName<"collector_b">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]),
- mma.intr_name>;
+ mma.intr_name, DefaultValue<0>>;
}
} // scale_vec_size
} // kind
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index d619b9a78c0c1..fa42c8789743e 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1345,18 +1345,6 @@ static Intrinsic::ID shouldUpgradeNVPTXBF16Intrinsic(StringRef Name) {
return Intrinsic::not_intrinsic;
}
-static Intrinsic::ID shouldUpgradeNVPTXTcgen05MMAIntrinsic(Function *F,
- StringRef Name) {
- if (!Name.consume_front("tcgen05.mma."))
- return Intrinsic::not_intrinsic;
-
- // tcgen05.mma.ws.* variants do not need collector-b appended.
- if (Name.starts_with("ws"))
- return Intrinsic::not_intrinsic;
-
- return F->getIntrinsicID();
-}
-
static bool consumeNVVMPtrAddrSpace(StringRef &Name) {
return Name.consume_front("local") || Name.consume_front("shared") ||
Name.consume_front("global") || Name.consume_front("constant") ||
@@ -1905,13 +1893,6 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
return true;
}
- // Upgrade tcgen05.mma intrinsics missing collector_usage_b.
- IID = shouldUpgradeNVPTXTcgen05MMAIntrinsic(F, Name);
- if (IID != Intrinsic::not_intrinsic) {
- NewFn = Intrinsic::getOrInsertDeclaration(F->getParent(), IID);
- return NewFn != F;
- }
-
// The following nvvm intrinsics correspond exactly to an LLVM idiom, but
// not to an intrinsic alone. We expand them in UpgradeIntrinsicCall.
//
@@ -5899,75 +5880,6 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
NewCall = Builder.CreateCall(NewFn, Args);
break;
}
- case Intrinsic::nvvm_tcgen05_mma_shared:
- case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg1:
- case Intrinsic::nvvm_tcgen05_mma_shared_disable_output_lane_cg2:
- case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_shared_mxf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block16:
- case Intrinsic::nvvm_tcgen05_mma_shared_mxf4nvf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_shared_mxf8f6f4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_shared_scale_d:
- case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg1:
- case Intrinsic::nvvm_tcgen05_mma_shared_scale_d_disable_output_lane_cg2:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg1:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_disable_output_lane_cg2:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block16:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf4nvf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_mxf8f6f4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg1:
- case Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d_disable_output_lane_cg2:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg1_ashift:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_disable_output_lane_cg2_ashift:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block16:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf4nvf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_mxf8f6f4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1:
- case Intrinsic::
- nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg1_ashift:
- case Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2:
- case Intrinsic::
- nvvm_tcgen05_mma_sp_tensor_scale_d_disable_output_lane_cg2_ashift:
- case Intrinsic::nvvm_tcgen05_mma_tensor:
- case Intrinsic::nvvm_tcgen05_mma_tensor_ashift:
- case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1:
- case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg1_ashift:
- case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2:
- case Intrinsic::nvvm_tcgen05_mma_tensor_disable_output_lane_cg2_ashift:
- case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block16:
- case Intrinsic::nvvm_tcgen05_mma_tensor_mxf4nvf4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale:
- case Intrinsic::nvvm_tcgen05_mma_tensor_mxf8f6f4_block_scale_block32:
- case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d:
- case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift:
- case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1:
- case Intrinsic::
- nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg1_ashift:
- case Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2:
- case Intrinsic::
- nvvm_tcgen05_mma_tensor_scale_d_disable_output_lane_cg2_ashift: {
- SmallVector<Value *, 12> Args(CI->args());
- Args.push_back(Builder.getInt32(0)); // collector_usage_b = discard(0)
- NewCall = Builder.CreateCall(NewFn, Args);
- break;
- }
case Intrinsic::riscv_sha256sig0:
case Intrinsic::riscv_sha256sig1:
case Intrinsic::riscv_sha256sum0:
More information about the llvm-commits
mailing list