[llvm] [LLVM][Tablegen] Add Default arguments support for Intrinsics in TableGen (PR #198557)
Varad Rahul Kamthe via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 4 06:34:35 PDT 2026
https://github.com/varadk27 updated https://github.com/llvm/llvm-project/pull/198557
>From bb412773427ede140100ead0845ac0f596753255 Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Tue, 19 May 2026 13:58:16 +0000
Subject: [PATCH 1/3] [LLVM][Intrinsics] Add default arguments support for
Intrinsics in TableGen
---
llvm/include/llvm/IR/Intrinsics.h | 9 ++
llvm/include/llvm/IR/Intrinsics.td | 10 ++
llvm/include/llvm/IR/IntrinsicsNVVM.td | 49 ++++++++++
llvm/lib/IR/AutoUpgrade.cpp | 92 +++++++++++++++++++
llvm/lib/IR/IRBuilder.cpp | 35 ++++++-
llvm/lib/IR/Intrinsics.cpp | 8 ++
.../NVPTX/default-arg-cp-async-bulk-dummy.ll | 14 +++
.../default-arg-explicit-value-preserved.ll | 13 +++
.../NVPTX/default-arg-move-i32-dummy.ll | 14 +++
.../NVPTX/default-arg-multiple-trailing.ll | 14 +++
.../NVPTX/default-arg-nanosleep-dummy.ll | 14 +++
.../CodeGen/NVPTX/default-arg-partial-fill.ll | 14 +++
.../CodeGen/NVPTX/default-arg-single-val.ll | 15 +++
.../NVPTX/default-arg-tcgen05-fence-dummy.ll | 14 +++
llvm/unittests/IR/IRBuilderTest.cpp | 64 +++++++++++++
.../TableGen/Basic/CodeGenIntrinsics.cpp | 49 ++++++++++
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h | 5 +
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 57 ++++++++++++
18 files changed, 488 insertions(+), 2 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-cp-async-bulk-dummy.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-explicit-value-preserved.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-move-i32-dummy.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-multiple-trailing.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-nanosleep-dummy.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-partial-fill.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
create mode 100644 llvm/test/CodeGen/NVPTX/default-arg-tcgen05-fence-dummy.ll
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index af97285fcee70..40fc0544b6600 100644
--- a/llvm/include/llvm/IR/Intrinsics.h
+++ b/llvm/include/llvm/IR/Intrinsics.h
@@ -91,6 +91,15 @@ LLVM_ABI bool isTriviallyScalarizable(ID id);
/// Returns true if the intrinsic has pretty printed immediate arguments.
LLVM_ABI bool hasPrettyPrintedArgs(ID id);
+/// Returns true if the intrinsic has any arguments with default values.
+/// Used by AutoUpgrade to decide whether to call getDefaultArgValue().
+LLVM_ABI bool hasDefaultArgs(ID id);
+
+/// Returns the default value for argument ArgIdx of intrinsic IID.
+/// Returns std::nullopt if that argument has no default.
+/// Supported types: all integer widths (i1/i8/i16/i32/i64).
+LLVM_ABI std::optional<int64_t> getDefaultArgValue(ID IID, unsigned ArgIdx);
+
/// isTargetIntrinsic - Returns true if IID is an intrinsic specific to a
/// certain target. If it is a generic intrinsic false is returned.
LLVM_ABI bool isTargetIntrinsic(ID IID);
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 993ddd7e33701..2cb2ba79c3c48 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -180,6 +180,16 @@ class ArgInfo<ArgIndex idx, list<ArgProperty> arg_properties> : IntrinsicPropert
list<ArgProperty> Properties = arg_properties;
}
+// DefaultIntArg - Attach an integer default value to an intrinsic argument.
+// The argument must be at the END of the parameter list.
+// Supported param types: llvm_i1_ty (use 0/1 or true/false), llvm_i8_ty, llvm_i16_ty, llvm_i32_ty,
+// llvm_i64_ty.
+// Write any integer literal as the value.
+class DefaultIntArg<ArgIndex idx, int val> : IntrinsicProperty {
+ int ArgNo = idx.Value; // 1-based internally; parser subtracts 1
+ int DefaultVal = val;
+}
+
def IntrNoReturn : IntrinsicProperty;
// Applied by default.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index cc245d4e17f5c..61ad479308f7c 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1301,6 +1301,37 @@ let TargetPrefix = "nvvm" in {
DefaultAttrsIntrinsic<[], [llvm_i32_ty],
[IntrConvergent, IntrNoMem, IntrHasSideEffects]>;
+ //
+ // Test intrinsic for default-arg feature: three required i32 args plus a
+ // fourth i32 with DefaultIntArg = 255.
+ //
+ def int_nvvm_test_add3 : NVVMBuiltin,
+ NVVMPureIntrinsic<[llvm_i32_ty],
+ [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
+ [ImmArg<ArgIndex<3>>,
+ DefaultIntArg<ArgIndex<3>, 255>]>;
+
+ //
+ // Test intrinsic for default-arg feature: four required i32 args plus two
+ // trailing args (i64, i32) with DefaultIntArg = 50 and 14.
+ //
+ def int_nvvm_test_add4 : NVVMBuiltin,
+ NVVMPureIntrinsic<[llvm_i32_ty],
+ [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
+ llvm_i64_ty, llvm_i32_ty],
+ [ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>,
+ DefaultIntArg<ArgIndex<4>, 50>,
+ DefaultIntArg<ArgIndex<5>, 14>]>;
+
+ //
+ // Dummy based on int_nvvm_nanosleep. Adds i1 flag with default = false.
+ //
+ def int_nvvm_nanosleep_dummy : NVVMBuiltin,
+ DefaultAttrsIntrinsic<[], [llvm_i32_ty, llvm_i1_ty],
+ [IntrConvergent, IntrNoMem, IntrHasSideEffects,
+ ImmArg<ArgIndex<1>>,
+ DefaultIntArg<ArgIndex<1>, false>]>;
+
//
// Performance Monitor Events (pm events) intrinsics
//
@@ -2064,6 +2095,12 @@ def int_nvvm_cp_async_bulk_commit_group : Intrinsic<[]>;
def int_nvvm_cp_async_bulk_wait_group :
Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;
+// Dummy based on cp_async_bulk_wait_group. Adds i1 flag with default = false.
+def int_nvvm_cp_async_bulk_wait_group_dummy :
+ Intrinsic<[], [llvm_i32_ty, llvm_i1_ty],
+ [ImmArg<ArgIndex<1>>,
+ DefaultIntArg<ArgIndex<1>, false>]>;
+
def int_nvvm_cp_async_bulk_wait_group_read :
Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;
@@ -2198,6 +2235,12 @@ let IntrProperties = [IntrNoMem] in {
def int_nvvm_move_ptr : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty]>;
}
+// Dummy based on int_nvvm_move_i32. Adds i64 second arg with default = 99.
+def int_nvvm_move_i32_dummy :
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i64_ty],
+ [ImmArg<ArgIndex<1>>,
+ DefaultIntArg<ArgIndex<1>, 99>]>;
+
// For getting the handle from a texture or surface variable
def int_nvvm_texsurf_handle
: NVVMPureIntrinsic<[llvm_i64_ty], [llvm_metadata_ty, llvm_anyptr_ty]>;
@@ -3180,6 +3223,12 @@ def int_nvvm_tcgen05_fence_before_thread_sync : Intrinsic<[], [],
def int_nvvm_tcgen05_fence_after_thread_sync : Intrinsic<[], [],
[IntrNoMem, IntrHasSideEffects]>;
+// Dummy based on tcgen05_fence. Two i32 args + i1 flag with default = false.
+def int_nvvm_tcgen05_fence_dummy :
+ Intrinsic<[], [llvm_i32_ty, llvm_i32_ty, llvm_i1_ty],
+ [IntrConvergent, ImmArg<ArgIndex<2>>,
+ DefaultIntArg<ArgIndex<2>, false>]>;
+
// Tcgen05 cp intrinsics
foreach cta_group = ["cg1", "cg2"] in {
foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 5fed947f9314d..52c97f16d4095 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1278,6 +1278,36 @@ static bool convertIntrinsicValidType(StringRef Name,
return false;
}
+static bool upgradeDeclWithDefaultArgs(Function *F, Function *&NewFn) {
+ // Look up the intrinsic ID by full name, e.g. "llvm.nvvm.add3.i32"
+ Intrinsic::ID IID = Intrinsic::lookupIntrinsicID(F->getName());
+ if (IID == Intrinsic::not_intrinsic)
+ return false;
+
+ // Fast path: this intrinsic has no default args annotated in .td
+ if (!Intrinsic::hasDefaultArgs(IID))
+ return false;
+
+ 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())
+ return false;
+
+ // Verify every missing trailing arg has a default in the table
+ for (unsigned Idx = F->arg_size(); Idx < FullDecl->arg_size(); ++Idx) {
+ if (!Intrinsic::getDefaultArgValue(IID, Idx))
+ return false;
+ }
+
+ NewFn = FullDecl;
+ return true;
+}
+
static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
bool CanUpgradeDebugIntrinsicsToRecords) {
assert(F && "Illegal to upgrade a non-existent Function.");
@@ -1896,6 +1926,9 @@ 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 (upgradeDeclWithDefaultArgs(F, NewFn))
+ return true;
+
return false;
}
@@ -4978,6 +5011,60 @@ static Value *upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI,
return nullptr;
}
+static bool upgradeCallWithDefaultArgs(CallBase *CI, Function *NewFn,
+ IRBuilder<> &Builder) {
+ Intrinsic::ID IID = NewFn->getIntrinsicID();
+
+ // Fast path: this intrinsic has no default args in the table.
+ if (!Intrinsic::hasDefaultArgs(IID))
+ return false;
+
+ 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)
+ return false;
+
+ // Start with the existing arguments from the old call.
+ SmallVector<Value *, 8> NewArgs(CI->args());
+
+ // Fill in each missing trailing argument from the table.
+ FunctionType *NewFT = NewFn->getFunctionType();
+ for (unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
+ std::optional<int64_t> DefaultVal = Intrinsic::getDefaultArgValue(IID, Idx);
+
+ // If any missing arg has no entry in the table, give up.
+ if (!DefaultVal)
+ return false;
+
+ 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, static_cast<uint64_t>(*DefaultVal)));
+ }
+
+ // Preserve operand bundles by creating the call with them.
+ SmallVector<OperandBundleDef, 1> OpBundles;
+ CI->getOperandBundlesAsDefs(OpBundles);
+ CallInst *NewCall = Builder.CreateCall(NewFn, NewArgs, OpBundles);
+
+ NewCall->takeName(CI);
+ NewCall->setCallingConv(CI->getCallingConv());
+ NewCall->copyMetadata(*CI);
+ if (auto *OldCI = dyn_cast<CallInst>(CI))
+ NewCall->setTailCallKind(OldCI->getTailCallKind());
+
+ CI->replaceAllUsesWith(NewCall);
+ CI->eraseFromParent();
+ return true;
+}
+
/// Upgrade a call to an old intrinsic. All argument and return casting must be
/// provided to seamlessly integrate with existing context.
void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
@@ -5083,6 +5170,11 @@ 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 DefaultIntArg
+ // in its .td definition, without needing a dedicated case.
+ if (upgradeCallWithDefaultArgs(CI, NewFn, Builder))
+ return;
DefaultCase();
return;
}
diff --git a/llvm/lib/IR/IRBuilder.cpp b/llvm/lib/IR/IRBuilder.cpp
index 7cdcb60a78aa3..2f978277717de 100644
--- a/llvm/lib/IR/IRBuilder.cpp
+++ b/llvm/lib/IR/IRBuilder.cpp
@@ -909,6 +909,33 @@ CallInst *IRBuilderBase::CreateGCGetPointerOffset(Value *DerivedPtr,
{DerivedPtr}, {}, Name);
}
+// Fill missing trailing arguments from the default-arg table.
+// Returns true if Args was extended; false if no fill was applied (e.g.
+// the intrinsic has no declared defaults, all args were already supplied,
+// or a missing arg has no default in the table).
+static bool fillDefaultArgs(Function *Fn, SmallVectorImpl<Value *> &Args) {
+ Intrinsic::ID IID = Fn->getIntrinsicID();
+ if (IID == Intrinsic::not_intrinsic || !Intrinsic::hasDefaultArgs(IID))
+ return false;
+
+ unsigned Expected = Fn->arg_size();
+ unsigned Supplied = Args.size();
+ if (Supplied >= Expected)
+ return false;
+
+ FunctionType *FT = Fn->getFunctionType();
+ for (unsigned Idx = Supplied; Idx < Expected; ++Idx) {
+ std::optional<int64_t> Default = Intrinsic::getDefaultArgValue(IID, Idx);
+ if (!Default)
+ return false;
+ Type *ParamTy = FT->getParamType(Idx);
+ if (!ParamTy->isIntegerTy())
+ return false;
+ Args.push_back(ConstantInt::get(ParamTy, static_cast<uint64_t>(*Default)));
+ }
+ return true;
+}
+
CallInst *IRBuilderBase::CreateUnaryIntrinsic(Intrinsic::ID ID, Value *V,
FMFSource FMFSource,
const Twine &Name) {
@@ -935,7 +962,9 @@ CallInst *IRBuilderBase::CreateIntrinsic(Intrinsic::ID ID,
ArrayRef<OperandBundleDef> OpBundles) {
Module *M = BB->getModule();
Function *Fn = Intrinsic::getOrInsertDeclaration(M, ID, OverloadTypes);
- return createCallHelper(Fn, Args, Name, FMFSource, OpBundles);
+ SmallVector<Value *, 8> FilledArgs(Args);
+ fillDefaultArgs(Fn, FilledArgs);
+ return createCallHelper(Fn, FilledArgs, Name, FMFSource, OpBundles);
}
CallInst *IRBuilderBase::CreateIntrinsic(Type *RetTy, Intrinsic::ID ID,
@@ -945,7 +974,9 @@ CallInst *IRBuilderBase::CreateIntrinsic(Type *RetTy, Intrinsic::ID ID,
Module *M = BB->getModule();
SmallVector<Type *> ArgTys = llvm::map_to_vector(Args, &Value::getType);
Function *Fn = Intrinsic::getOrInsertDeclaration(M, ID, RetTy, ArgTys);
- return createCallHelper(Fn, Args, Name, FMFSource);
+ SmallVector<Value *, 8> FilledArgs(Args);
+ fillDefaultArgs(Fn, FilledArgs);
+ return createCallHelper(Fn, FilledArgs, Name, FMFSource);
}
CallInst *IRBuilderBase::CreateConstrainedFPBinOp(
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index 008d823315156..bec585e362992 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -639,6 +639,11 @@ bool Intrinsic::hasPrettyPrintedArgs(ID id){
#include "llvm/IR/IntrinsicImpl.inc"
}
+bool Intrinsic::hasDefaultArgs(ID id){
+#define GET_INTRINSIC_DEFAULT_ARG_TABLE
+#include "llvm/IR/IntrinsicImpl.inc"
+}
+
/// Table of per-target intrinsic name tables.
#define GET_INTRINSIC_TARGET_DATA
#include "llvm/IR/IntrinsicImpl.inc"
@@ -1337,3 +1342,6 @@ Intrinsic::ID Intrinsic::getDeinterleaveIntrinsicID(unsigned Factor) {
#define GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
#include "llvm/IR/IntrinsicImpl.inc"
+
+#define GET_INTRINSIC_DEFAULT_ARG_VALUES
+#include "llvm/IR/IntrinsicImpl.inc"
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-cp-async-bulk-dummy.ll b/llvm/test/CodeGen/NVPTX/default-arg-cp-async-bulk-dummy.ll
new file mode 100644
index 0000000000000..7df6494b3a40a
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-cp-async-bulk-dummy.ll
@@ -0,0 +1,14 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; Plain Intrinsic<[]> — void-returning, cp_async_bulk style.
+; Old decl: 1 arg (i32). New: 2 args. Arg[1] = i1, default = false.
+
+declare void @llvm.nvvm.cp.async.bulk.wait.group.dummy(i32)
+
+define void @test_cp_async_bulk_wait_group_dummy(i32 %n) {
+entry:
+ call void @llvm.nvvm.cp.async.bulk.wait.group.dummy(i32 %n)
+ ret void
+}
+; CHECK-LABEL: @test_cp_async_bulk_wait_group_dummy
+; CHECK: call void @llvm.nvvm.cp.async.bulk.wait.group.dummy(i32 %n, i1 false)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-explicit-value-preserved.ll b/llvm/test/CodeGen/NVPTX/default-arg-explicit-value-preserved.ll
new file mode 100644
index 0000000000000..f8510e7b893b3
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-explicit-value-preserved.ll
@@ -0,0 +1,13 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; The 4-arg call passes an explicit value. It must be preserved.
+
+declare i32 @llvm.nvvm.test.add3(i32, i32, i32, i32)
+
+define i32 @test_explicit_value_preserved(i32 %a, i32 %b, i32 %c) {
+entry:
+ %r = call i32 @llvm.nvvm.test.add3(i32 %a, i32 %b, i32 %c, i32 42)
+ ret i32 %r
+}
+; CHECK-LABEL: @test_explicit_value_preserved
+; CHECK: call i32 @llvm.nvvm.test.add3(i32 %a, i32 %b, i32 %c, i32 42)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-move-i32-dummy.ll b/llvm/test/CodeGen/NVPTX/default-arg-move-i32-dummy.ll
new file mode 100644
index 0000000000000..cfc7e0b26a25a
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-move-i32-dummy.ll
@@ -0,0 +1,14 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; DefaultAttrsIntrinsic with return type (no NVVMBuiltin).
+; Old decl: 1 arg (i32). New: 2 args. Arg[1] = i64, default = 99.
+
+declare i32 @llvm.nvvm.move.i32.dummy(i32)
+
+define i32 @test_move_i32_dummy(i32 %a) {
+entry:
+ %r = call i32 @llvm.nvvm.move.i32.dummy(i32 %a)
+ ret i32 %r
+}
+; CHECK-LABEL: @test_move_i32_dummy
+; CHECK: call i32 @llvm.nvvm.move.i32.dummy(i32 %a, i64 99)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-multiple-trailing.ll b/llvm/test/CodeGen/NVPTX/default-arg-multiple-trailing.ll
new file mode 100644
index 0000000000000..c048c79c5d92f
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-multiple-trailing.ll
@@ -0,0 +1,14 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; test_add4 has 6 params; last two have defaults (i64 50, i32 14).
+; Old declaration has only 4 args; both missing defaults are inserted.
+
+declare i32 @llvm.nvvm.test.add4(i32, i32, i32, i32)
+
+define i32 @test_multiple_trailing_defaults(i32 %a, i32 %b, i32 %c, i32 %d) {
+entry:
+ %r = call i32 @llvm.nvvm.test.add4(i32 %a, i32 %b, i32 %c, i32 %d)
+ ret i32 %r
+}
+; CHECK-LABEL: @test_multiple_trailing_defaults
+; CHECK: call i32 @llvm.nvvm.test.add4(i32 %a, i32 %b, i32 %c, i32 %d, i64 50, i32 14)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-nanosleep-dummy.ll b/llvm/test/CodeGen/NVPTX/default-arg-nanosleep-dummy.ll
new file mode 100644
index 0000000000000..95b6f46dfa9fd
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-nanosleep-dummy.ll
@@ -0,0 +1,14 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; NVVMBuiltin + DefaultAttrsIntrinsic (void-returning).
+; Old decl: 1 arg (i32). New: 2 args. Arg[1] = i1, default = false.
+
+declare void @llvm.nvvm.nanosleep.dummy(i32)
+
+define void @test_nanosleep_dummy(i32 %ns) {
+entry:
+ call void @llvm.nvvm.nanosleep.dummy(i32 %ns)
+ ret void
+}
+; CHECK-LABEL: @test_nanosleep_dummy
+; CHECK: call void @llvm.nvvm.nanosleep.dummy(i32 %ns, i1 false)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-partial-fill.ll b/llvm/test/CodeGen/NVPTX/default-arg-partial-fill.ll
new file mode 100644
index 0000000000000..db296f5a1aeda
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-partial-fill.ll
@@ -0,0 +1,14 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; Old decl has 5 args (index 5 missing). Arg[4] passed explicitly (i64 99),
+; must NOT be overridden. Arg[5] filled with default i32 14.
+
+declare i32 @llvm.nvvm.test.add4(i32, i32, i32, i32, i64)
+
+define i32 @test_partial_fill(i32 %a, i32 %b, i32 %c, i32 %d) {
+entry:
+ %r = call i32 @llvm.nvvm.test.add4(i32 %a, i32 %b, i32 %c, i32 %d, i64 99)
+ ret i32 %r
+}
+; CHECK-LABEL: @test_partial_fill
+; CHECK: call i32 @llvm.nvvm.test.add4(i32 %a, i32 %b, i32 %c, i32 %d, i64 99, i32 14)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll b/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
new file mode 100644
index 0000000000000..62074e02cbfaa
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
@@ -0,0 +1,15 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; Verifies that AutoUpgrade correctly fills in a missing trailing argument
+; with the value declared via DefaultIntArg in the intrinsic's TableGen
+; definition.
+
+declare i32 @llvm.nvvm.test.add3(i32, i32, i32)
+
+define i32 @test_default_single_val(i32 %a, i32 %b, i32 %c) {
+entry:
+ %r = call i32 @llvm.nvvm.test.add3(i32 %a, i32 %b, i32 %c)
+ ret i32 %r
+}
+; CHECK-LABEL: @test_default_single_val
+; CHECK: call i32 @llvm.nvvm.test.add3(i32 %a, i32 %b, i32 %c, i32 255)
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-tcgen05-fence-dummy.ll b/llvm/test/CodeGen/NVPTX/default-arg-tcgen05-fence-dummy.ll
new file mode 100644
index 0000000000000..81f0ad507bd42
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/default-arg-tcgen05-fence-dummy.ll
@@ -0,0 +1,14 @@
+; RUN: opt -S < %s | FileCheck %s
+;
+; Plain Intrinsic<[]>, tcgen05 style.
+; Old decl: 2 args (i32, i32). New: 3 args. Arg[2] = i1, default = false.
+
+declare void @llvm.nvvm.tcgen05.fence.dummy(i32, i32)
+
+define void @test_tcgen05_fence_dummy(i32 %flags, i32 %token) {
+entry:
+ call void @llvm.nvvm.tcgen05.fence.dummy(i32 %flags, i32 %token)
+ ret void
+}
+; CHECK-LABEL: @test_tcgen05_fence_dummy
+; CHECK: call void @llvm.nvvm.tcgen05.fence.dummy(i32 %flags, i32 %token, i1 false)
diff --git a/llvm/unittests/IR/IRBuilderTest.cpp b/llvm/unittests/IR/IRBuilderTest.cpp
index b4f1c97f03aca..a4a27d661551d 100644
--- a/llvm/unittests/IR/IRBuilderTest.cpp
+++ b/llvm/unittests/IR/IRBuilderTest.cpp
@@ -15,6 +15,7 @@
#include "llvm/IR/Function.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/IntrinsicsAArch64.h"
+#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
@@ -160,6 +161,69 @@ TEST_F(IRBuilderTest, Intrinsics) {
EXPECT_EQ(II->getIntrinsicID(), Intrinsic::set_rounding);
}
+TEST_F(IRBuilderTest, IntrinsicWithDefaultArgs) {
+ IRBuilder<> Builder(BB);
+
+ // Use test_add3, which has 4 params with a default on the trailing i32.
+ Value *A = Builder.getInt32(1);
+ Value *B = Builder.getInt32(2);
+ Value *C = Builder.getInt32(3);
+
+ // Case 1: caller supplies only 3 of 4 args; default i32 255 is filled.
+ CallInst *Call = Builder.CreateIntrinsic(Intrinsic::nvvm_test_add3,
+ /*Types=*/{}, {A, B, C});
+ EXPECT_EQ(Call->arg_size(), 4u);
+ auto *FilledArg = dyn_cast<ConstantInt>(Call->getArgOperand(3));
+ ASSERT_NE(FilledArg, nullptr);
+ EXPECT_TRUE(FilledArg->getType()->isIntegerTy(32));
+ EXPECT_EQ(FilledArg->getZExtValue(), 255u);
+
+ // Case 2: caller supplies all 4 args explicitly; default NOT applied.
+ Value *Explicit = Builder.getInt32(42);
+ CallInst *Call2 = Builder.CreateIntrinsic(Intrinsic::nvvm_test_add3,
+ /*Types=*/{}, {A, B, C, Explicit});
+ EXPECT_EQ(Call2->arg_size(), 4u);
+ EXPECT_EQ(Call2->getArgOperand(3), Explicit);
+}
+
+TEST_F(IRBuilderTest, IntrinsicWithMultipleDefaultArgs) {
+ IRBuilder<> Builder(BB);
+
+ // int_nvvm_test_add4: (i32, i32, i32, i32, i64, i32) with defaults
+ // ArgIndex<4> = 50 (i64)
+ // ArgIndex<5> = 14 (i32)
+ Value *A = Builder.getInt32(1);
+ Value *B = Builder.getInt32(2);
+ Value *C = Builder.getInt32(3);
+ Value *D = Builder.getInt32(4);
+
+ // Case 1: caller supplies 4 of 6 args; BOTH defaults are filled.
+ CallInst *Call = Builder.CreateIntrinsic(Intrinsic::nvvm_test_add4,
+ /*Types=*/{}, {A, B, C, D});
+ EXPECT_EQ(Call->arg_size(), 6u);
+
+ auto *Arg4 = dyn_cast<ConstantInt>(Call->getArgOperand(4));
+ ASSERT_NE(Arg4, nullptr);
+ EXPECT_TRUE(Arg4->getType()->isIntegerTy(64));
+ EXPECT_EQ(Arg4->getZExtValue(), 50u);
+
+ auto *Arg5 = dyn_cast<ConstantInt>(Call->getArgOperand(5));
+ ASSERT_NE(Arg5, nullptr);
+ EXPECT_TRUE(Arg5->getType()->isIntegerTy(32));
+ EXPECT_EQ(Arg5->getZExtValue(), 14u);
+
+ // Case 2: caller supplies 5 of 6 args; only the trailing i32 is filled,
+ // the explicitly-supplied i64 is preserved.
+ Value *ExplicitI64 = Builder.getInt64(99);
+ CallInst *Call2 = Builder.CreateIntrinsic(
+ Intrinsic::nvvm_test_add4, /*Types=*/{}, {A, B, C, D, ExplicitI64});
+ EXPECT_EQ(Call2->arg_size(), 6u);
+ EXPECT_EQ(Call2->getArgOperand(4), ExplicitI64);
+ auto *FilledArg5 = dyn_cast<ConstantInt>(Call2->getArgOperand(5));
+ ASSERT_NE(FilledArg5, nullptr);
+ EXPECT_EQ(FilledArg5->getZExtValue(), 14u);
+}
+
TEST_F(IRBuilderTest, IntrinsicMangling) {
IRBuilder<> Builder(BB);
Type *VoidTy = Builder.getVoidTy();
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 05ce7f22468c3..30c10b3d33567 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -385,6 +385,35 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
// Sort the argument attributes for later benefit.
for (auto &Attrs : ArgumentAttributes)
llvm::sort(Attrs);
+
+ // Validate: every DefaultIntArg must be on an argument that also has
+ // ImmArg. This narrows the scope of the default-argument feature to
+ // immediate-only parameters, which is the primary intended use case.
+ for (size_t i = 0; i < ParamDefaultValues.size(); ++i) {
+ if (!ParamDefaultValues[i].has_value())
+ continue;
+ if (!isParamImmArg(i))
+ PrintFatalError(TheDef->getLoc(),
+ "DefaultIntArg on argument " + Twine(i) +
+ " requires that argument to also have ImmArg");
+ }
+
+ // Validate: defaults must form a contiguous trailing block ending at
+ // the last parameter (mirrors C++ default-argument rules).
+ unsigned NumParams = IS.ParamTys.size();
+ bool SeenDefault = false;
+ for (unsigned i = 0; i < NumParams; ++i) {
+ bool HasDefault =
+ (i < ParamDefaultValues.size() && ParamDefaultValues[i].has_value());
+ if (HasDefault) {
+ SeenDefault = true;
+ } else if (SeenDefault) {
+ PrintFatalError(TheDef->getLoc(),
+ "DefaultIntArg gap at argument " + Twine(i) +
+ ". Defaults must form a contiguous trailing block "
+ "ending at the last parameter.");
+ }
+ }
}
void CodeGenIntrinsic::setDefaultProperties(
@@ -525,6 +554,14 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
}
}
addPrettyPrintFunction(ArgNo - 1, ArgName, FuncName);
+ } else if (R->isSubClassOf("DefaultIntArg")) {
+ unsigned ArgNo = R->getValueAsInt("ArgNo");
+ if (ArgNo < 1)
+ PrintFatalError(
+ R->getLoc(),
+ "DefaultIntArg requires ArgNo >= 1 (0 is the return value)");
+ int64_t IntVal = R->getValueAsInt("DefaultVal");
+ addDefaultArgValue(ArgNo - 1, IntVal);
} else {
llvm_unreachable("Unknown property!");
}
@@ -581,3 +618,15 @@ void CodeGenIntrinsic::addPrettyPrintFunction(unsigned ArgIdx,
It->FuncName + "'");
PrettyPrintFunctions.emplace_back(ArgIdx, ArgName, FuncName);
}
+
+void CodeGenIntrinsic::addDefaultArgValue(unsigned ArgIdx, int64_t Value) {
+ if (ArgIdx >= ParamDefaultValues.size())
+ ParamDefaultValues.resize(ArgIdx + 1, std::nullopt);
+
+ if (ParamDefaultValues[ArgIdx].has_value())
+ PrintFatalError(TheDef->getLoc(), "Default value for argument " +
+ Twine(ArgIdx) +
+ " is already defined");
+
+ ParamDefaultValues[ArgIdx] = Value;
+}
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index 7d2d21382b319..71da6f3c3490b 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -168,9 +168,14 @@ struct CodeGenIntrinsic {
/// Vector that stores ArgInfo (ArgIndex, ArgName, FunctionName).
SmallVector<PrettyPrintArgInfo> PrettyPrintFunctions;
+ /// Default values for parameters. Index = param index.
+ SmallVector<std::optional<int64_t>> ParamDefaultValues;
+
void addPrettyPrintFunction(unsigned ArgIdx, StringRef ArgName,
StringRef FuncName);
+ void addDefaultArgValue(unsigned ArgIdx, int64_t Value);
+
bool hasProperty(enum SDNP Prop) const { return Properties & (1 << Prop); }
/// Goes through all IntrProperties that have IsDefault value set and sets
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index 02b2cd9850997..cfb146036725a 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -73,6 +73,9 @@ class IntrinsicEmitter {
void EmitAttributes(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
void EmitPrettyPrintArguments(const CodeGenIntrinsicTable &Ints,
raw_ostream &OS);
+ void EmitIntrinsicToDefaultArgTable(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS);
+ void EmitDefaultArgValues(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
void EmitIntrinsicToBuiltinMap(const CodeGenIntrinsicTable &Ints,
bool IsClang, raw_ostream &OS);
};
@@ -129,6 +132,12 @@ void IntrinsicEmitter::run(raw_ostream &OS, bool Enums) {
// Emit Pretty Print attribute.
EmitPrettyPrintArguments(Ints, OS);
+ // Emit the Intrinsic ID -> default argument bit table.
+ EmitIntrinsicToDefaultArgTable(Ints, OS);
+
+ // Emit the default argument value lookup function.
+ EmitDefaultArgValues(Ints, OS);
+
// Emit code to translate Clang builtins into LLVM intrinsics.
EmitIntrinsicToBuiltinMap(Ints, true, OS);
@@ -877,6 +886,54 @@ void Intrinsic::printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS, const Cons
})";
}
+void IntrinsicEmitter::EmitIntrinsicToDefaultArgTable(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
+ EmitIntrinsicBitTable(
+ Ints, OS, "GET_INTRINSIC_DEFAULT_ARG_TABLE", "DefaultArgTable",
+ "Intrinsic ID to default argument bitset.",
+ [](const CodeGenIntrinsic &Int) {
+ return llvm::any_of(
+ Int.ParamDefaultValues,
+ [](const std::optional<int64_t> &V) { return V.has_value(); });
+ });
+}
+
+void IntrinsicEmitter::EmitDefaultArgValues(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS) {
+ OS << R"(
+#ifdef GET_INTRINSIC_DEFAULT_ARG_VALUES
+std::optional<int64_t> Intrinsic::getDefaultArgValue(
+ ID IID, unsigned ArgIdx) {
+ switch (IID) {
+)";
+
+ for (const CodeGenIntrinsic &Int : Ints) {
+ bool HasAny = llvm::any_of(
+ Int.ParamDefaultValues,
+ [](const std::optional<int64_t> &V) { return V.has_value(); });
+ if (!HasAny)
+ continue;
+
+ OS << " case " << Int.EnumName << ": {\n";
+ OS << " switch (ArgIdx) {\n";
+ for (auto [Idx, Val] : llvm::enumerate(Int.ParamDefaultValues)) {
+ if (!Val.has_value())
+ continue;
+ OS << " case " << Idx << ": return " << *Val << "LL;\n";
+ }
+ OS << " default: return std::nullopt;\n";
+ OS << " }\n";
+ OS << " }\n";
+ }
+
+ OS << R"( default:
+ return std::nullopt;
+ }
+}
+#endif // GET_INTRINSIC_DEFAULT_ARG_VALUES
+)";
+}
+
void IntrinsicEmitter::EmitIntrinsicToBuiltinMap(
const CodeGenIntrinsicTable &Ints, bool IsClang, raw_ostream &OS) {
StringRef CompilerName = IsClang ? "Clang" : "MS";
>From 629643516b50c9b448dd0c8de73cddf3239a8c0c Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Tue, 26 May 2026 10:06:50 +0000
Subject: [PATCH 2/3] Merge DefaultVal into ImmArg, drop IRBuilder
---
llvm/include/llvm/IR/Intrinsics.td | 29 +++++----
llvm/include/llvm/IR/IntrinsicsNVVM.td | 31 ++++-----
llvm/lib/IR/AutoUpgrade.cpp | 14 ++--
llvm/lib/IR/IRBuilder.cpp | 35 +---------
.../CodeGen/NVPTX/default-arg-single-val.ll | 2 +-
llvm/unittests/IR/IRBuilderTest.cpp | 64 -------------------
.../TableGen/Basic/CodeGenIntrinsics.cpp | 34 ++++------
7 files changed, 54 insertions(+), 155 deletions(-)
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 2cb2ba79c3c48..2fc629d080f30 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -130,9 +130,24 @@ class Returned<ArgIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
-// ImmArg - The specified argument must be an immediate.
-class ImmArg<ArgIndex idx> : IntrinsicProperty {
+// DefaultVal - 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.
+class DefaultVal<int val> {
+ bit HasDefault = 1;
+ int Value = val;
+}
+
+// Sentinel — used as the default for ImmArg's optional DefaultVal parameter.
+def NoDefault : DefaultVal<0> {
+ let HasDefault = 0;
+}
+
+// ImmArg - The specified argument must be an immediate. The optional
+// second template parameter declares a default value for AutoUpgrade.
+class ImmArg<ArgIndex idx, DefaultVal val = NoDefault> : IntrinsicProperty {
int ArgNo = idx.Value;
+ DefaultVal Default = val;
}
// ReadOnly - The specified argument pointer is not written to through the
@@ -180,16 +195,6 @@ class ArgInfo<ArgIndex idx, list<ArgProperty> arg_properties> : IntrinsicPropert
list<ArgProperty> Properties = arg_properties;
}
-// DefaultIntArg - Attach an integer default value to an intrinsic argument.
-// The argument must be at the END of the parameter list.
-// Supported param types: llvm_i1_ty (use 0/1 or true/false), llvm_i8_ty, llvm_i16_ty, llvm_i32_ty,
-// llvm_i64_ty.
-// Write any integer literal as the value.
-class DefaultIntArg<ArgIndex idx, int val> : IntrinsicProperty {
- int ArgNo = idx.Value; // 1-based internally; parser subtracts 1
- int DefaultVal = val;
-}
-
def IntrNoReturn : IntrinsicProperty;
// Applied by default.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 61ad479308f7c..17dac72b27977 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1303,34 +1303,31 @@ let TargetPrefix = "nvvm" in {
//
// Test intrinsic for default-arg feature: three required i32 args plus a
- // fourth i32 with DefaultIntArg = 255.
+ // fourth i32 with DefaultVal = 255.
//
- def int_nvvm_test_add3 : NVVMBuiltin,
+ def int_nvvm_test_add3 :
NVVMPureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<3>>,
- DefaultIntArg<ArgIndex<3>, 255>]>;
+ [ImmArg<ArgIndex<3>, DefaultVal<255>>]>;
//
// Test intrinsic for default-arg feature: four required i32 args plus two
- // trailing args (i64, i32) with DefaultIntArg = 50 and 14.
+ // trailing args (i64, i32) with DefaultVal = 50 and 14.
//
- def int_nvvm_test_add4 : NVVMBuiltin,
+ def int_nvvm_test_add4 :
NVVMPureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
llvm_i64_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<4>>, ImmArg<ArgIndex<5>>,
- DefaultIntArg<ArgIndex<4>, 50>,
- DefaultIntArg<ArgIndex<5>, 14>]>;
+ [ImmArg<ArgIndex<4>, DefaultVal<50>>,
+ ImmArg<ArgIndex<5>, DefaultVal<14>>]>;
//
// Dummy based on int_nvvm_nanosleep. Adds i1 flag with default = false.
//
- def int_nvvm_nanosleep_dummy : NVVMBuiltin,
+ def int_nvvm_nanosleep_dummy :
DefaultAttrsIntrinsic<[], [llvm_i32_ty, llvm_i1_ty],
[IntrConvergent, IntrNoMem, IntrHasSideEffects,
- ImmArg<ArgIndex<1>>,
- DefaultIntArg<ArgIndex<1>, false>]>;
+ ImmArg<ArgIndex<1>, DefaultVal<false>>]>;
//
// Performance Monitor Events (pm events) intrinsics
@@ -2098,8 +2095,7 @@ def int_nvvm_cp_async_bulk_wait_group :
// Dummy based on cp_async_bulk_wait_group. Adds i1 flag with default = false.
def int_nvvm_cp_async_bulk_wait_group_dummy :
Intrinsic<[], [llvm_i32_ty, llvm_i1_ty],
- [ImmArg<ArgIndex<1>>,
- DefaultIntArg<ArgIndex<1>, false>]>;
+ [ImmArg<ArgIndex<1>, DefaultVal<false>>]>;
def int_nvvm_cp_async_bulk_wait_group_read :
Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;
@@ -2238,8 +2234,7 @@ let IntrProperties = [IntrNoMem] in {
// Dummy based on int_nvvm_move_i32. Adds i64 second arg with default = 99.
def int_nvvm_move_i32_dummy :
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i64_ty],
- [ImmArg<ArgIndex<1>>,
- DefaultIntArg<ArgIndex<1>, 99>]>;
+ [ImmArg<ArgIndex<1>, DefaultVal<99>>]>;
// For getting the handle from a texture or surface variable
def int_nvvm_texsurf_handle
@@ -3226,8 +3221,8 @@ def int_nvvm_tcgen05_fence_after_thread_sync : Intrinsic<[], [],
// Dummy based on tcgen05_fence. Two i32 args + i1 flag with default = false.
def int_nvvm_tcgen05_fence_dummy :
Intrinsic<[], [llvm_i32_ty, llvm_i32_ty, llvm_i1_ty],
- [IntrConvergent, ImmArg<ArgIndex<2>>,
- DefaultIntArg<ArgIndex<2>, false>]>;
+ [IntrConvergent,
+ ImmArg<ArgIndex<2>, DefaultVal<false>>]>;
// Tcgen05 cp intrinsics
foreach cta_group = ["cg1", "cg2"] in {
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index 52c97f16d4095..d0411a20580a9 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1278,7 +1278,7 @@ static bool convertIntrinsicValidType(StringRef Name,
return false;
}
-static bool upgradeDeclWithDefaultArgs(Function *F, Function *&NewFn) {
+static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn) {
// Look up the intrinsic ID by full name, e.g. "llvm.nvvm.add3.i32"
Intrinsic::ID IID = Intrinsic::lookupIntrinsicID(F->getName());
if (IID == Intrinsic::not_intrinsic)
@@ -1288,6 +1288,8 @@ static bool upgradeDeclWithDefaultArgs(Function *F, Function *&NewFn) {
if (!Intrinsic::hasDefaultArgs(IID))
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;
@@ -1926,7 +1928,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 (upgradeDeclWithDefaultArgs(F, NewFn))
+ if (upgradeIntrinsicDeclWithDefaultArgs(F, NewFn))
return true;
return false;
@@ -5011,8 +5013,8 @@ static Value *upgradeConvertIntrinsicCall(StringRef Name, CallBase *CI,
return nullptr;
}
-static bool upgradeCallWithDefaultArgs(CallBase *CI, Function *NewFn,
- IRBuilder<> &Builder) {
+static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
+ IRBuilder<> &Builder) {
Intrinsic::ID IID = NewFn->getIntrinsicID();
// Fast path: this intrinsic has no default args in the table.
@@ -5171,9 +5173,9 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
switch (NewFn->getIntrinsicID()) {
default: {
// Last resort: try the data-driven default-arg upgrade.
- // Handles any intrinsic annotated with DefaultIntArg
+ // Handles any intrinsic annotated with ImmArg<..., DefaultVal<...>>
// in its .td definition, without needing a dedicated case.
- if (upgradeCallWithDefaultArgs(CI, NewFn, Builder))
+ if (upgradeIntrinsicCallWithDefaultArgs(CI, NewFn, Builder))
return;
DefaultCase();
return;
diff --git a/llvm/lib/IR/IRBuilder.cpp b/llvm/lib/IR/IRBuilder.cpp
index 2f978277717de..7cdcb60a78aa3 100644
--- a/llvm/lib/IR/IRBuilder.cpp
+++ b/llvm/lib/IR/IRBuilder.cpp
@@ -909,33 +909,6 @@ CallInst *IRBuilderBase::CreateGCGetPointerOffset(Value *DerivedPtr,
{DerivedPtr}, {}, Name);
}
-// Fill missing trailing arguments from the default-arg table.
-// Returns true if Args was extended; false if no fill was applied (e.g.
-// the intrinsic has no declared defaults, all args were already supplied,
-// or a missing arg has no default in the table).
-static bool fillDefaultArgs(Function *Fn, SmallVectorImpl<Value *> &Args) {
- Intrinsic::ID IID = Fn->getIntrinsicID();
- if (IID == Intrinsic::not_intrinsic || !Intrinsic::hasDefaultArgs(IID))
- return false;
-
- unsigned Expected = Fn->arg_size();
- unsigned Supplied = Args.size();
- if (Supplied >= Expected)
- return false;
-
- FunctionType *FT = Fn->getFunctionType();
- for (unsigned Idx = Supplied; Idx < Expected; ++Idx) {
- std::optional<int64_t> Default = Intrinsic::getDefaultArgValue(IID, Idx);
- if (!Default)
- return false;
- Type *ParamTy = FT->getParamType(Idx);
- if (!ParamTy->isIntegerTy())
- return false;
- Args.push_back(ConstantInt::get(ParamTy, static_cast<uint64_t>(*Default)));
- }
- return true;
-}
-
CallInst *IRBuilderBase::CreateUnaryIntrinsic(Intrinsic::ID ID, Value *V,
FMFSource FMFSource,
const Twine &Name) {
@@ -962,9 +935,7 @@ CallInst *IRBuilderBase::CreateIntrinsic(Intrinsic::ID ID,
ArrayRef<OperandBundleDef> OpBundles) {
Module *M = BB->getModule();
Function *Fn = Intrinsic::getOrInsertDeclaration(M, ID, OverloadTypes);
- SmallVector<Value *, 8> FilledArgs(Args);
- fillDefaultArgs(Fn, FilledArgs);
- return createCallHelper(Fn, FilledArgs, Name, FMFSource, OpBundles);
+ return createCallHelper(Fn, Args, Name, FMFSource, OpBundles);
}
CallInst *IRBuilderBase::CreateIntrinsic(Type *RetTy, Intrinsic::ID ID,
@@ -974,9 +945,7 @@ CallInst *IRBuilderBase::CreateIntrinsic(Type *RetTy, Intrinsic::ID ID,
Module *M = BB->getModule();
SmallVector<Type *> ArgTys = llvm::map_to_vector(Args, &Value::getType);
Function *Fn = Intrinsic::getOrInsertDeclaration(M, ID, RetTy, ArgTys);
- SmallVector<Value *, 8> FilledArgs(Args);
- fillDefaultArgs(Fn, FilledArgs);
- return createCallHelper(Fn, FilledArgs, Name, FMFSource);
+ return createCallHelper(Fn, Args, Name, FMFSource);
}
CallInst *IRBuilderBase::CreateConstrainedFPBinOp(
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll b/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
index 62074e02cbfaa..e9cad39c262aa 100644
--- a/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
+++ b/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
@@ -1,7 +1,7 @@
; RUN: opt -S < %s | FileCheck %s
;
; Verifies that AutoUpgrade correctly fills in a missing trailing argument
-; with the value declared via DefaultIntArg in the intrinsic's TableGen
+; with the value declared via DefaultVal in the intrinsic's TableGen
; definition.
declare i32 @llvm.nvvm.test.add3(i32, i32, i32)
diff --git a/llvm/unittests/IR/IRBuilderTest.cpp b/llvm/unittests/IR/IRBuilderTest.cpp
index a4a27d661551d..b4f1c97f03aca 100644
--- a/llvm/unittests/IR/IRBuilderTest.cpp
+++ b/llvm/unittests/IR/IRBuilderTest.cpp
@@ -15,7 +15,6 @@
#include "llvm/IR/Function.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/IntrinsicsAArch64.h"
-#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Module.h"
@@ -161,69 +160,6 @@ TEST_F(IRBuilderTest, Intrinsics) {
EXPECT_EQ(II->getIntrinsicID(), Intrinsic::set_rounding);
}
-TEST_F(IRBuilderTest, IntrinsicWithDefaultArgs) {
- IRBuilder<> Builder(BB);
-
- // Use test_add3, which has 4 params with a default on the trailing i32.
- Value *A = Builder.getInt32(1);
- Value *B = Builder.getInt32(2);
- Value *C = Builder.getInt32(3);
-
- // Case 1: caller supplies only 3 of 4 args; default i32 255 is filled.
- CallInst *Call = Builder.CreateIntrinsic(Intrinsic::nvvm_test_add3,
- /*Types=*/{}, {A, B, C});
- EXPECT_EQ(Call->arg_size(), 4u);
- auto *FilledArg = dyn_cast<ConstantInt>(Call->getArgOperand(3));
- ASSERT_NE(FilledArg, nullptr);
- EXPECT_TRUE(FilledArg->getType()->isIntegerTy(32));
- EXPECT_EQ(FilledArg->getZExtValue(), 255u);
-
- // Case 2: caller supplies all 4 args explicitly; default NOT applied.
- Value *Explicit = Builder.getInt32(42);
- CallInst *Call2 = Builder.CreateIntrinsic(Intrinsic::nvvm_test_add3,
- /*Types=*/{}, {A, B, C, Explicit});
- EXPECT_EQ(Call2->arg_size(), 4u);
- EXPECT_EQ(Call2->getArgOperand(3), Explicit);
-}
-
-TEST_F(IRBuilderTest, IntrinsicWithMultipleDefaultArgs) {
- IRBuilder<> Builder(BB);
-
- // int_nvvm_test_add4: (i32, i32, i32, i32, i64, i32) with defaults
- // ArgIndex<4> = 50 (i64)
- // ArgIndex<5> = 14 (i32)
- Value *A = Builder.getInt32(1);
- Value *B = Builder.getInt32(2);
- Value *C = Builder.getInt32(3);
- Value *D = Builder.getInt32(4);
-
- // Case 1: caller supplies 4 of 6 args; BOTH defaults are filled.
- CallInst *Call = Builder.CreateIntrinsic(Intrinsic::nvvm_test_add4,
- /*Types=*/{}, {A, B, C, D});
- EXPECT_EQ(Call->arg_size(), 6u);
-
- auto *Arg4 = dyn_cast<ConstantInt>(Call->getArgOperand(4));
- ASSERT_NE(Arg4, nullptr);
- EXPECT_TRUE(Arg4->getType()->isIntegerTy(64));
- EXPECT_EQ(Arg4->getZExtValue(), 50u);
-
- auto *Arg5 = dyn_cast<ConstantInt>(Call->getArgOperand(5));
- ASSERT_NE(Arg5, nullptr);
- EXPECT_TRUE(Arg5->getType()->isIntegerTy(32));
- EXPECT_EQ(Arg5->getZExtValue(), 14u);
-
- // Case 2: caller supplies 5 of 6 args; only the trailing i32 is filled,
- // the explicitly-supplied i64 is preserved.
- Value *ExplicitI64 = Builder.getInt64(99);
- CallInst *Call2 = Builder.CreateIntrinsic(
- Intrinsic::nvvm_test_add4, /*Types=*/{}, {A, B, C, D, ExplicitI64});
- EXPECT_EQ(Call2->arg_size(), 6u);
- EXPECT_EQ(Call2->getArgOperand(4), ExplicitI64);
- auto *FilledArg5 = dyn_cast<ConstantInt>(Call2->getArgOperand(5));
- ASSERT_NE(FilledArg5, nullptr);
- EXPECT_EQ(FilledArg5->getZExtValue(), 14u);
-}
-
TEST_F(IRBuilderTest, IntrinsicMangling) {
IRBuilder<> Builder(BB);
Type *VoidTy = Builder.getVoidTy();
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 30c10b3d33567..2fccd7615d7be 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -386,18 +386,6 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
for (auto &Attrs : ArgumentAttributes)
llvm::sort(Attrs);
- // Validate: every DefaultIntArg must be on an argument that also has
- // ImmArg. This narrows the scope of the default-argument feature to
- // immediate-only parameters, which is the primary intended use case.
- for (size_t i = 0; i < ParamDefaultValues.size(); ++i) {
- if (!ParamDefaultValues[i].has_value())
- continue;
- if (!isParamImmArg(i))
- PrintFatalError(TheDef->getLoc(),
- "DefaultIntArg on argument " + Twine(i) +
- " requires that argument to also have ImmArg");
- }
-
// Validate: defaults must form a contiguous trailing block ending at
// the last parameter (mirrors C++ default-argument rules).
unsigned NumParams = IS.ParamTys.size();
@@ -409,7 +397,7 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
SeenDefault = true;
} else if (SeenDefault) {
PrintFatalError(TheDef->getLoc(),
- "DefaultIntArg gap at argument " + Twine(i) +
+ "DefaultVal missing at argument " + Twine(i) +
". Defaults must form a contiguous trailing block "
"ending at the last parameter.");
}
@@ -518,6 +506,18 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
} else if (R->isSubClassOf("ImmArg")) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
addArgAttribute(ArgNo, ImmArg);
+
+ // If a DefaultVal (not the NoDefault sentinel) was supplied, record it.
+ const RecordVal *DefField = R->getValue("Default");
+ if (DefField) {
+ if (const auto *DI = dyn_cast<DefInit>(DefField->getValue())) {
+ const Record *DefRec = DI->getDef();
+ if (DefRec->getValueAsBit("HasDefault")) {
+ int64_t Value = DefRec->getValueAsInt("Value");
+ addDefaultArgValue(ArgNo - 1, Value);
+ }
+ }
+ }
} else if (R->isSubClassOf("Align")) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
uint64_t Align = R->getValueAsInt("Align");
@@ -554,14 +554,6 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
}
}
addPrettyPrintFunction(ArgNo - 1, ArgName, FuncName);
- } else if (R->isSubClassOf("DefaultIntArg")) {
- unsigned ArgNo = R->getValueAsInt("ArgNo");
- if (ArgNo < 1)
- PrintFatalError(
- R->getLoc(),
- "DefaultIntArg requires ArgNo >= 1 (0 is the return value)");
- int64_t IntVal = R->getValueAsInt("DefaultVal");
- addDefaultArgValue(ArgNo - 1, IntVal);
} else {
llvm_unreachable("Unknown property!");
}
>From 5cfcf94ce20c1a56414ea09e064a3f28ac73b402 Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <vkamthe at nvidia.com>
Date: Thu, 4 Jun 2026 13:33:17 +0000
Subject: [PATCH 3/3] [Intrinsics] Table-driven default-arg lookup; add
getAllDefaultArgValues
Address review comments:
- Replace per-intrinsic switch with SequenceToOffsetTable + offset array.
- New API returns (FirstDefault, ArrayRef<int64_t>) in one call.
- Rename DefaultVal -> DefaultValue; use ? for NoDefault sentinel.
- Add TableGen-time type + range validation.
---
llvm/include/llvm/IR/Intrinsics.h | 16 +--
llvm/include/llvm/IR/Intrinsics.td | 15 +-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 18 +--
llvm/lib/IR/AutoUpgrade.cpp | 27 ++--
llvm/lib/IR/Intrinsics.cpp | 8 +-
.../CodeGen/NVPTX/default-arg-single-val.ll | 2 +-
.../TableGen/Basic/CodeGenIntrinsics.cpp | 48 +++++--
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 129 ++++++++++++------
8 files changed, 161 insertions(+), 102 deletions(-)
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index 40fc0544b6600..83d92d1690268 100644
--- a/llvm/include/llvm/IR/Intrinsics.h
+++ b/llvm/include/llvm/IR/Intrinsics.h
@@ -91,14 +91,14 @@ LLVM_ABI bool isTriviallyScalarizable(ID id);
/// Returns true if the intrinsic has pretty printed immediate arguments.
LLVM_ABI bool hasPrettyPrintedArgs(ID id);
-/// Returns true if the intrinsic has any arguments with default values.
-/// Used by AutoUpgrade to decide whether to call getDefaultArgValue().
-LLVM_ABI bool hasDefaultArgs(ID id);
-
-/// Returns the default value for argument ArgIdx of intrinsic IID.
-/// Returns std::nullopt if that argument has no default.
-/// Supported types: all integer widths (i1/i8/i16/i32/i64).
-LLVM_ABI std::optional<int64_t> getDefaultArgValue(ID IID, unsigned ArgIdx);
+/// Returns the first default argument index and an ArrayRef of all
+/// default values for the trailing parameters of intrinsic IID.
+/// Returns {0, empty} if the intrinsic has no default arguments.
+///
+/// The defaults are stored contiguously starting at FirstDefault and
+/// extending to the last parameter (mirrors C++ default-argument
+/// rules).
+LLVM_ABI std::pair<unsigned, ArrayRef<int64_t>> getAllDefaultArgValues(ID IID);
/// isTargetIntrinsic - Returns true if IID is an intrinsic specific to a
/// certain target. If it is a generic intrinsic false is returned.
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 2fc629d080f30..98f52efba33cd 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -130,24 +130,23 @@ class Returned<ArgIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
-// DefaultVal - A value carrier paired with ImmArg to declare a default for
+// 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.
-class DefaultVal<int val> {
- bit HasDefault = 1;
+class DefaultValue<int val> {
int Value = val;
}
-// Sentinel — used as the default for ImmArg's optional DefaultVal parameter.
-def NoDefault : DefaultVal<0> {
- let HasDefault = 0;
+// Sentinel — used as the default for ImmArg's optional DefaultValue parameter.
+def NoDefault : DefaultValue<0> {
+ let Value = ?;
}
// ImmArg - The specified argument must be an immediate. The optional
// second template parameter declares a default value for AutoUpgrade.
-class ImmArg<ArgIndex idx, DefaultVal val = NoDefault> : IntrinsicProperty {
+class ImmArg<ArgIndex idx, DefaultValue val = NoDefault> : IntrinsicProperty {
int ArgNo = idx.Value;
- DefaultVal Default = val;
+ DefaultValue Default = val;
}
// ReadOnly - The specified argument pointer is not written to through the
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 17dac72b27977..ba118c4c5efb3 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -1303,23 +1303,23 @@ let TargetPrefix = "nvvm" in {
//
// Test intrinsic for default-arg feature: three required i32 args plus a
- // fourth i32 with DefaultVal = 255.
+ // fourth i32 with DefaultValue = 255.
//
def int_nvvm_test_add3 :
NVVMPureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<3>, DefaultVal<255>>]>;
+ [ImmArg<ArgIndex<3>, DefaultValue<255>>]>;
//
// Test intrinsic for default-arg feature: four required i32 args plus two
- // trailing args (i64, i32) with DefaultVal = 50 and 14.
+ // trailing args (i64, i32) with DefaultValue = 50 and 14.
//
def int_nvvm_test_add4 :
NVVMPureIntrinsic<[llvm_i32_ty],
[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
llvm_i64_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<4>, DefaultVal<50>>,
- ImmArg<ArgIndex<5>, DefaultVal<14>>]>;
+ [ImmArg<ArgIndex<4>, DefaultValue<50>>,
+ ImmArg<ArgIndex<5>, DefaultValue<14>>]>;
//
// Dummy based on int_nvvm_nanosleep. Adds i1 flag with default = false.
@@ -1327,7 +1327,7 @@ let TargetPrefix = "nvvm" in {
def int_nvvm_nanosleep_dummy :
DefaultAttrsIntrinsic<[], [llvm_i32_ty, llvm_i1_ty],
[IntrConvergent, IntrNoMem, IntrHasSideEffects,
- ImmArg<ArgIndex<1>, DefaultVal<false>>]>;
+ ImmArg<ArgIndex<1>, DefaultValue<false>>]>;
//
// Performance Monitor Events (pm events) intrinsics
@@ -2095,7 +2095,7 @@ def int_nvvm_cp_async_bulk_wait_group :
// Dummy based on cp_async_bulk_wait_group. Adds i1 flag with default = false.
def int_nvvm_cp_async_bulk_wait_group_dummy :
Intrinsic<[], [llvm_i32_ty, llvm_i1_ty],
- [ImmArg<ArgIndex<1>, DefaultVal<false>>]>;
+ [ImmArg<ArgIndex<1>, DefaultValue<false>>]>;
def int_nvvm_cp_async_bulk_wait_group_read :
Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;
@@ -2234,7 +2234,7 @@ let IntrProperties = [IntrNoMem] in {
// Dummy based on int_nvvm_move_i32. Adds i64 second arg with default = 99.
def int_nvvm_move_i32_dummy :
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i64_ty],
- [ImmArg<ArgIndex<1>, DefaultVal<99>>]>;
+ [ImmArg<ArgIndex<1>, DefaultValue<99>>]>;
// For getting the handle from a texture or surface variable
def int_nvvm_texsurf_handle
@@ -3222,7 +3222,7 @@ def int_nvvm_tcgen05_fence_after_thread_sync : Intrinsic<[], [],
def int_nvvm_tcgen05_fence_dummy :
Intrinsic<[], [llvm_i32_ty, llvm_i32_ty, llvm_i1_ty],
[IntrConvergent,
- ImmArg<ArgIndex<2>, DefaultVal<false>>]>;
+ ImmArg<ArgIndex<2>, DefaultValue<false>>]>;
// Tcgen05 cp intrinsics
foreach cta_group = ["cg1", "cg2"] in {
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index d0411a20580a9..d14eb01f4b4b7 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -1279,13 +1279,12 @@ static bool convertIntrinsicValidType(StringRef Name,
}
static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn) {
- // Look up the intrinsic ID by full name, e.g. "llvm.nvvm.add3.i32"
Intrinsic::ID IID = Intrinsic::lookupIntrinsicID(F->getName());
if (IID == Intrinsic::not_intrinsic)
return false;
- // Fast path: this intrinsic has no default args annotated in .td
- if (!Intrinsic::hasDefaultArgs(IID))
+ auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
+ if (Defaults.empty())
return false;
// Overloaded intrinsics are out of scope for the default-arg feature
@@ -1300,9 +1299,10 @@ static bool upgradeIntrinsicDeclWithDefaultArgs(Function *F, Function *&NewFn) {
if (F->arg_size() >= FullDecl->arg_size())
return false;
- // Verify every missing trailing arg has a default in the table
+ // Every missing trailing arg must fall within the default range
+ // [FirstDefault, FirstDefault + Defaults.size()).
for (unsigned Idx = F->arg_size(); Idx < FullDecl->arg_size(); ++Idx) {
- if (!Intrinsic::getDefaultArgValue(IID, Idx))
+ if (Idx < FirstDefault || Idx >= FirstDefault + Defaults.size())
return false;
}
@@ -5017,8 +5017,8 @@ static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
IRBuilder<> &Builder) {
Intrinsic::ID IID = NewFn->getIntrinsicID();
- // Fast path: this intrinsic has no default args in the table.
- if (!Intrinsic::hasDefaultArgs(IID))
+ auto [FirstDefault, Defaults] = Intrinsic::getAllDefaultArgValues(IID);
+ if (Defaults.empty())
return false;
unsigned OldArgCount = CI->arg_size();
@@ -5035,20 +5035,15 @@ static bool upgradeIntrinsicCallWithDefaultArgs(CallBase *CI, Function *NewFn,
// Fill in each missing trailing argument from the table.
FunctionType *NewFT = NewFn->getFunctionType();
for (unsigned Idx = OldArgCount; Idx < NewArgCount; ++Idx) {
- std::optional<int64_t> DefaultVal = Intrinsic::getDefaultArgValue(IID, Idx);
-
- // If any missing arg has no entry in the table, give up.
- if (!DefaultVal)
+ if (Idx < FirstDefault || Idx >= FirstDefault + Defaults.size())
return false;
-
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, static_cast<uint64_t>(*DefaultVal)));
+ NewArgs.push_back(ConstantInt::get(
+ ParamTy, static_cast<uint64_t>(Defaults[Idx - FirstDefault])));
}
// Preserve operand bundles by creating the call with them.
@@ -5173,7 +5168,7 @@ void llvm::UpgradeIntrinsicCall(CallBase *CI, Function *NewFn) {
switch (NewFn->getIntrinsicID()) {
default: {
// Last resort: try the data-driven default-arg upgrade.
- // Handles any intrinsic annotated with ImmArg<..., DefaultVal<...>>
+ // Handles any intrinsic annotated with ImmArg<..., DefaultValue<...>>
// in its .td definition, without needing a dedicated case.
if (upgradeIntrinsicCallWithDefaultArgs(CI, NewFn, Builder))
return;
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index bec585e362992..767187eec4e8c 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -639,11 +639,6 @@ bool Intrinsic::hasPrettyPrintedArgs(ID id){
#include "llvm/IR/IntrinsicImpl.inc"
}
-bool Intrinsic::hasDefaultArgs(ID id){
-#define GET_INTRINSIC_DEFAULT_ARG_TABLE
-#include "llvm/IR/IntrinsicImpl.inc"
-}
-
/// Table of per-target intrinsic name tables.
#define GET_INTRINSIC_TARGET_DATA
#include "llvm/IR/IntrinsicImpl.inc"
@@ -1343,5 +1338,8 @@ Intrinsic::ID Intrinsic::getDeinterleaveIntrinsicID(unsigned Factor) {
#define GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
#include "llvm/IR/IntrinsicImpl.inc"
+// Emit the default-argument values table and lookup function
+// (Intrinsic::getAllDefaultArgValues).
#define GET_INTRINSIC_DEFAULT_ARG_VALUES
#include "llvm/IR/IntrinsicImpl.inc"
+#undef GET_INTRINSIC_DEFAULT_ARG_VALUES
diff --git a/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll b/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
index e9cad39c262aa..52e8d3a3dc73f 100644
--- a/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
+++ b/llvm/test/CodeGen/NVPTX/default-arg-single-val.ll
@@ -1,7 +1,7 @@
; RUN: opt -S < %s | FileCheck %s
;
; Verifies that AutoUpgrade correctly fills in a missing trailing argument
-; with the value declared via DefaultVal in the intrinsic's TableGen
+; with the value declared via DefaultValue in the intrinsic's TableGen
; definition.
declare i32 @llvm.nvvm.test.add3(i32, i32, i32)
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index 2fccd7615d7be..a62b495ed488e 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -397,11 +397,40 @@ CodeGenIntrinsic::CodeGenIntrinsic(const Record *R,
SeenDefault = true;
} else if (SeenDefault) {
PrintFatalError(TheDef->getLoc(),
- "DefaultVal missing at argument " + Twine(i) +
+ "DefaultValue missing for argument " + Twine(i) +
". Defaults must form a contiguous trailing block "
"ending at the last parameter.");
}
}
+
+ // Validate each declared default:
+ // (a) the parameter is an integer type
+ // (b) the default value fits in the declared integer width
+ for (unsigned i = 0; i < ParamDefaultValues.size(); ++i) {
+ if (!ParamDefaultValues[i].has_value())
+ continue;
+ const Record *ParamTy = IS.ParamTys[i];
+ if (!ParamTy->isSubClassOf("LLVMType")) {
+ PrintFatalError(TheDef->getLoc(),
+ "DefaultValue at argument " + Twine(i) +
+ " requires an integer parameter type");
+ }
+ const Record *VT = ParamTy->getValueAsDef("VT");
+ if (!VT->getValueAsBit("isInteger")) {
+ PrintFatalError(TheDef->getLoc(),
+ "DefaultValue at argument " + Twine(i) +
+ " requires an integer parameter type");
+ }
+ unsigned Width = VT->getValueAsInt("Size");
+ int64_t Value = *ParamDefaultValues[i];
+ int64_t MaxUnsigned = (Width >= 64) ? INT64_MAX : ((1LL << Width) - 1);
+ int64_t MinSigned = (Width >= 64) ? INT64_MIN : -(1LL << (Width - 1));
+ if (Value < MinSigned || Value > MaxUnsigned) {
+ PrintFatalError(TheDef->getLoc(),
+ "DefaultValue " + Twine(Value) + " out of range for i" +
+ Twine(Width) + " parameter at argument " + Twine(i));
+ }
+ }
}
void CodeGenIntrinsic::setDefaultProperties(
@@ -507,16 +536,13 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
addArgAttribute(ArgNo, ImmArg);
- // If a DefaultVal (not the NoDefault sentinel) was supplied, record it.
- const RecordVal *DefField = R->getValue("Default");
- if (DefField) {
- if (const auto *DI = dyn_cast<DefInit>(DefField->getValue())) {
- const Record *DefRec = DI->getDef();
- if (DefRec->getValueAsBit("HasDefault")) {
- int64_t Value = DefRec->getValueAsInt("Value");
- addDefaultArgValue(ArgNo - 1, Value);
- }
- }
+ // 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");
+ addDefaultArgValue(ArgNo - 1, Value);
}
} else if (R->isSubClassOf("Align")) {
unsigned ArgNo = R->getValueAsInt("ArgNo");
diff --git a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
index cfb146036725a..a7d9dec285962 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -73,9 +73,8 @@ class IntrinsicEmitter {
void EmitAttributes(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
void EmitPrettyPrintArguments(const CodeGenIntrinsicTable &Ints,
raw_ostream &OS);
- void EmitIntrinsicToDefaultArgTable(const CodeGenIntrinsicTable &Ints,
- raw_ostream &OS);
- void EmitDefaultArgValues(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
+ void EmitDefaultArgValuesTable(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS);
void EmitIntrinsicToBuiltinMap(const CodeGenIntrinsicTable &Ints,
bool IsClang, raw_ostream &OS);
};
@@ -132,11 +131,8 @@ void IntrinsicEmitter::run(raw_ostream &OS, bool Enums) {
// Emit Pretty Print attribute.
EmitPrettyPrintArguments(Ints, OS);
- // Emit the Intrinsic ID -> default argument bit table.
- EmitIntrinsicToDefaultArgTable(Ints, OS);
-
- // Emit the default argument value lookup function.
- EmitDefaultArgValues(Ints, OS);
+ // Emit the default-argument values table and lookup function.
+ EmitDefaultArgValuesTable(Ints, OS);
// Emit code to translate Clang builtins into LLVM intrinsics.
EmitIntrinsicToBuiltinMap(Ints, true, OS);
@@ -886,52 +882,97 @@ void Intrinsic::printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS, const Cons
})";
}
-void IntrinsicEmitter::EmitIntrinsicToDefaultArgTable(
+void IntrinsicEmitter::EmitDefaultArgValuesTable(
const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
- EmitIntrinsicBitTable(
- Ints, OS, "GET_INTRINSIC_DEFAULT_ARG_TABLE", "DefaultArgTable",
- "Intrinsic ID to default argument bitset.",
- [](const CodeGenIntrinsic &Int) {
- return llvm::any_of(
- Int.ParamDefaultValues,
- [](const std::optional<int64_t> &V) { return V.has_value(); });
- });
-}
+ // Build the per-intrinsic default-value sequences:
+ // [Header = (NumDefaults << 32) | FirstDefault, val0, val1, ...]
+ // Each value is the int64_t default for one parameter, stored as a
+ // uint64_t bit pattern (sign is preserved via reinterpret).
+ //
+ // Intrinsics with no defaults share a single sentinel sequence {0},
+ // which decodes to Header = 0 → NumDefaults = 0 → empty defaults.
+ // SequenceToOffsetTable deduplicates identical sequences.
-void IntrinsicEmitter::EmitDefaultArgValues(const CodeGenIntrinsicTable &Ints,
- raw_ostream &OS) {
- OS << R"(
-#ifdef GET_INTRINSIC_DEFAULT_ARG_VALUES
-std::optional<int64_t> Intrinsic::getDefaultArgValue(
- ID IID, unsigned ArgIdx) {
- switch (IID) {
-)";
+ using Sequence = std::vector<uint64_t>;
- for (const CodeGenIntrinsic &Int : Ints) {
- bool HasAny = llvm::any_of(
- Int.ParamDefaultValues,
- [](const std::optional<int64_t> &V) { return V.has_value(); });
- if (!HasAny)
+ SequenceToOffsetTable<Sequence> Table;
+ std::vector<Sequence> PerIntrinsic(Ints.size());
+
+ // Sentinel "no defaults" sequence. Registered so every intrinsic without
+ // defaults has a valid offset to point at. The actual offset is whatever
+ // SequenceToOffsetTable assigns after suffix-dedup (not necessarily 0).
+ Sequence Sentinel = {0ULL};
+ Table.add(Sentinel);
+
+ for (size_t i = 0; i < Ints.size(); ++i) {
+ const CodeGenIntrinsic &Int = Ints[i];
+ if (Int.ParamDefaultValues.empty()) {
+ PerIntrinsic[i] = Sentinel;
continue;
+ }
- OS << " case " << Int.EnumName << ": {\n";
- OS << " switch (ArgIdx) {\n";
- for (auto [Idx, Val] : llvm::enumerate(Int.ParamDefaultValues)) {
- if (!Val.has_value())
- continue;
- OS << " case " << Idx << ": return " << *Val << "LL;\n";
+ // Find the first parameter with a default.
+ unsigned FirstDefault = 0;
+ for (size_t j = 0; j < Int.ParamDefaultValues.size(); ++j) {
+ if (Int.ParamDefaultValues[j].has_value()) {
+ FirstDefault = j;
+ break;
+ }
}
- OS << " default: return std::nullopt;\n";
- OS << " }\n";
- OS << " }\n";
+ unsigned NumDefaults =
+ static_cast<unsigned>(Int.ParamDefaultValues.size()) - FirstDefault;
+
+ Sequence Seq;
+ Seq.reserve(NumDefaults + 1);
+ uint64_t Header = (static_cast<uint64_t>(NumDefaults) << 32) | FirstDefault;
+ Seq.push_back(Header);
+ for (size_t j = FirstDefault; j < Int.ParamDefaultValues.size(); ++j) {
+ assert(Int.ParamDefaultValues[j].has_value() &&
+ "Default block must be contiguous");
+ Seq.push_back(static_cast<uint64_t>(*Int.ParamDefaultValues[j]));
+ }
+ Table.add(Seq);
+ PerIntrinsic[i] = std::move(Seq);
}
- OS << R"( default:
- return std::nullopt;
- }
+ Table.layout();
+
+ OS << "#ifdef GET_INTRINSIC_DEFAULT_ARG_VALUES\n";
+
+ // Emit the deduplicated flat values table.
+ OS << "static const uint64_t DefaultArgValuesTable[] = {\n";
+ Table.emit(OS, [](raw_ostream &OS, uint64_t Val) {
+ OS << " 0x" << Twine::utohexstr(Val) << "ULL";
+ });
+ OS << "};\n\n";
+
+ // Emit the per-intrinsic offset table. Entry #0 is for the invalid
+ // Intrinsic::not_intrinsic (IID 0) and must point to the sentinel
+ // (which decodes to "no defaults"). Real intrinsics follow.
+ OS << "static const uint32_t DefaultArgValuesTableOffset[] = {\n";
+ OS << " " << Table.get(Sentinel) << ", // not_intrinsic\n";
+ for (size_t i = 0; i < Ints.size(); ++i)
+ OS << " " << Table.get(PerIntrinsic[i]) << ",\n";
+ OS << "};\n\n";
+
+ // Emit the lookup function body.
+ OS << R"(
+std::pair<unsigned, ArrayRef<int64_t>>
+Intrinsic::getAllDefaultArgValues(ID IID) {
+ uint32_t Offset = DefaultArgValuesTableOffset[IID];
+ uint64_t Header = DefaultArgValuesTable[Offset];
+ uint32_t FirstDefault = Header & 0xFFFFFFFFu;
+ uint32_t NumDefaults = (Header >> 32) & 0xFFFFFFFFu;
+ if (NumDefaults == 0)
+ return {0, {}};
+ return {FirstDefault,
+ ArrayRef<int64_t>(reinterpret_cast<const int64_t *>(
+ &DefaultArgValuesTable[Offset + 1]),
+ NumDefaults)};
}
-#endif // GET_INTRINSIC_DEFAULT_ARG_VALUES
)";
+
+ OS << "#endif // GET_INTRINSIC_DEFAULT_ARG_VALUES\n";
}
void IntrinsicEmitter::EmitIntrinsicToBuiltinMap(
More information about the llvm-commits
mailing list