Skip to content

Commit fadac9b

Browse files
DharuniRAcharyaDharuniRAcharya
committed
[LLVM-Tablegen] Pretty Printing Arguments in LLVM Intrinsics
This patch adds LLVM infrastructure to support pretty printing arguments of the intrinsics. The motivation is to increase the readability of LLVM intrinsics and facilitate easy modifications and debugging of LLVM IR. This adds a property ArgInfo<ArgIndex, "argName", "functionName"> to the intrinsic arguments that enables printing self-explanatory inline comment for the arguments. The addition of pretty print support can provide a simple, low-overhead feature that enhances usability of LLVM intrinsics without disrupting existing workflows. Link to the RFC: https://discourse.llvm.org/t/rfc-pretty-printing-immediate-arguments-in-llvm-intrinsics/88536 Signed-off-by: Dharuni R Acharya <[email protected]>
1 parent 3f3ca76 commit fadac9b

File tree

11 files changed

+350
-9
lines changed

11 files changed

+350
-9
lines changed

‎llvm/include/llvm/IR/Intrinsics.h‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,8 @@ class LLVMContext;
3030
class Module;
3131
class AttributeList;
3232
class AttributeSet;
33+
class raw_ostream;
34+
class Constant;
3335

3436
/// This namespace contains an enum with a value for every intrinsic/builtin
3537
/// function known by LLVM. The enum values are returned by
@@ -81,6 +83,9 @@ namespace Intrinsic {
8183
/// Returns true if the intrinsic can be overloaded.
8284
LLVM_ABI bool isOverloaded(ID id);
8385

86+
/// Returns true if the intrinsic has pretty printed immediate arguments.
87+
LLVM_ABI bool hasPrettyPrintedArgs(ID id);
88+
8489
/// isTargetIntrinsic - Returns true if IID is an intrinsic specific to a
8590
/// certain target. If it is a generic intrinsic false is returned.
8691
LLVM_ABI bool isTargetIntrinsic(ID IID);
@@ -284,6 +289,10 @@ namespace Intrinsic {
284289
/// N.
285290
LLVM_ABI Intrinsic::ID getDeinterleaveIntrinsicID(unsigned Factor);
286291

292+
/// Print the argument info for the arguments with ArgInfo.
293+
LLVM_ABI void printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS,
294+
const Constant *ImmArgVal);
295+
287296
} // namespace Intrinsic
288297

289298
} // namespace llvm

‎llvm/include/llvm/IR/Intrinsics.td‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -87,6 +87,14 @@ class NoUndef<AttrIndex idx> : IntrinsicProperty {
8787
int ArgNo = idx.Value;
8888
}
8989

90+
// ArgInfo - The specified argument has an argument name and an optional argument printing
91+
// function for diagnostic output.
92+
class ArgInfo<AttrIndex idx, string argname, string funcname = ""> : IntrinsicProperty {
93+
int ArgNo = idx.Value;
94+
string ArgName = argname;
95+
string FunctionName = funcname;
96+
}
97+
9098
// NonNull - The return value or specified argument is not null.
9199
class NonNull<AttrIndex idx> : IntrinsicProperty {
92100
int ArgNo = idx.Value;

‎llvm/include/llvm/IR/IntrinsicsNVVM.td‎

Lines changed: 15 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2871,7 +2871,14 @@ foreach sp = [0, 1] in {
28712871
defvar nargs = !size(args);
28722872
defvar scale_d_imm = ArgIndex<!sub(nargs, 1)>;
28732873
defvar scale_d_imm_range = [ImmArg<scale_d_imm>, Range<scale_d_imm, 0, 16>];
2874-
defvar intrinsic_properties = !listconcat(
2874+
2875+
// Check if this is the specific llvm.nvvm.tcgen05.mma.tensor intrinsic.
2876+
defvar is_target_intrinsic = !and(!eq(sp, 0),
2877+
!eq(space, "tensor"),
2878+
!eq(scale_d, 0),
2879+
!eq(ashift, 0));
2880+
2881+
defvar base_properties = !listconcat(
28752882
mma.common_intr_props,
28762883
!if(!eq(scale_d, 1), scale_d_imm_range, []),
28772884
[Range<ArgIndex<nargs>, 0, !if(!eq(scale_d, 1), 2, 4)>, // kind
@@ -2881,6 +2888,13 @@ foreach sp = [0, 1] in {
28812888
]
28822889
);
28832890

2891+
defvar intrinsic_properties = !if(is_target_intrinsic,
2892+
!listconcat(base_properties,
2893+
[ArgInfo<ArgIndex<nargs>, "kind", "printTcgen05MMAKind">,
2894+
ArgInfo<ArgIndex<!add(nargs, 1)>, "cta_group">,
2895+
ArgInfo<ArgIndex<!add(nargs, 2)>, "collector", "printTcgen05CollectorUsageOp">]),
2896+
base_properties);
2897+
28842898
def mma.record:
28852899
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
28862900
mma.intr>;

‎llvm/include/llvm/IR/NVVMIntrinsicUtils.h‎

Lines changed: 48 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,11 @@
1818
#include <stdint.h>
1919

2020
#include "llvm/ADT/APFloat.h"
21+
#include "llvm/ADT/APInt.h"
22+
#include "llvm/IR/Constants.h"
2123
#include "llvm/IR/Intrinsics.h"
2224
#include "llvm/IR/IntrinsicsNVPTX.h"
25+
#include "llvm/Support/raw_ostream.h"
2326

2427
namespace llvm {
2528
namespace nvvm {
@@ -659,6 +662,51 @@ inline APFloat::roundingMode GetFMARoundingMode(Intrinsic::ID IntrinsicID) {
659662
llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM fma");
660663
}
661664

665+
inline void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal) {
666+
if (const ConstantInt *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
667+
uint64_t Val = CI->getZExtValue();
668+
switch (static_cast<Tcgen05MMAKind>(Val)) {
669+
case Tcgen05MMAKind::F16:
670+
OS << "f16";
671+
return;
672+
case Tcgen05MMAKind::TF32:
673+
OS << "tf32";
674+
return;
675+
case Tcgen05MMAKind::F8F6F4:
676+
OS << "f8f6f4";
677+
return;
678+
case Tcgen05MMAKind::I8:
679+
OS << "i8";
680+
return;
681+
}
682+
}
683+
llvm_unreachable(
684+
"printTcgen05MMAKind called with invalid value for immediate argument");
685+
}
686+
687+
inline void printTcgen05CollectorUsageOp(raw_ostream &OS,
688+
const Constant *ImmArgVal) {
689+
if (const ConstantInt *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
690+
uint64_t Val = CI->getZExtValue();
691+
switch (static_cast<Tcgen05CollectorUsageOp>(Val)) {
692+
case Tcgen05CollectorUsageOp::DISCARD:
693+
OS << "discard";
694+
return;
695+
case Tcgen05CollectorUsageOp::LASTUSE:
696+
OS << "lastuse";
697+
return;
698+
case Tcgen05CollectorUsageOp::FILL:
699+
OS << "fill";
700+
return;
701+
case Tcgen05CollectorUsageOp::USE:
702+
OS << "use";
703+
return;
704+
}
705+
}
706+
llvm_unreachable("printTcgen05CollectorUsageOp called with invalid value for "
707+
"immediate argument");
708+
}
709+
662710
} // namespace nvvm
663711
} // namespace llvm
664712
#endif // LLVM_IR_NVVMINTRINSICUTILS_H

‎llvm/lib/IR/AsmWriter.cpp‎

Lines changed: 33 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,7 @@
5353
#include "llvm/IR/Instruction.h"
5454
#include "llvm/IR/Instructions.h"
5555
#include "llvm/IR/IntrinsicInst.h"
56+
#include "llvm/IR/Intrinsics.h"
5657
#include "llvm/IR/LLVMContext.h"
5758
#include "llvm/IR/Metadata.h"
5859
#include "llvm/IR/Module.h"
@@ -4575,12 +4576,38 @@ void AssemblyWriter::printInstruction(const Instruction &I) {
45754576
Out << ' ';
45764577
writeOperand(Operand, false);
45774578
Out << '(';
4579+
bool HasPrettyPrintedArgs =
4580+
isa<IntrinsicInst>(CI) &&
4581+
Intrinsic::hasPrettyPrintedArgs(CI->getIntrinsicID());
4582+
45784583
ListSeparator LS;
4579-
for (unsigned op = 0, Eop = CI->arg_size(); op < Eop; ++op) {
4580-
Out << LS;
4581-
writeParamOperand(CI->getArgOperand(op), PAL.getParamAttrs(op));
4582-
}
4584+
if (HasPrettyPrintedArgs) {
4585+
Function *CalledFunc = CI->getCalledFunction();
4586+
auto PrintArgComment = [&](unsigned ArgNo) {
4587+
const Constant *ConstArg = dyn_cast<Constant>(CI->getArgOperand(ArgNo));
4588+
if (!ConstArg)
4589+
return;
4590+
std::string ArgComment;
4591+
raw_string_ostream ArgCommentStream(ArgComment);
4592+
Intrinsic::ID IID = CalledFunc->getIntrinsicID();
4593+
Intrinsic::printImmArg(IID, ArgNo, ArgCommentStream, ConstArg);
4594+
if (ArgComment.empty())
4595+
return;
4596+
Out << "/* " << ArgComment << " */ ";
4597+
};
45834598

4599+
for (unsigned ArgNo = 0, NumArgs = CI->arg_size(); ArgNo < NumArgs;
4600+
++ArgNo) {
4601+
Out << LS;
4602+
PrintArgComment(ArgNo);
4603+
writeParamOperand(CI->getArgOperand(ArgNo), PAL.getParamAttrs(ArgNo));
4604+
}
4605+
} else {
4606+
for (unsigned op = 0, Eop = CI->arg_size(); op < Eop; ++op) {
4607+
Out << LS;
4608+
writeParamOperand(CI->getArgOperand(op), PAL.getParamAttrs(op));
4609+
}
4610+
}
45844611
// Emit an ellipsis if this is a musttail call in a vararg function. This
45854612
// is only to aid readability, musttail calls forward varargs by default.
45864613
if (CI->isMustTailCall() && CI->getParent() &&
@@ -5004,12 +5031,10 @@ void AssemblyWriter::printUseLists(const Function *F) {
50045031
//===----------------------------------------------------------------------===//
50055032

50065033
void Function::print(raw_ostream &ROS, AssemblyAnnotationWriter *AAW,
5007-
bool ShouldPreserveUseListOrder,
5008-
bool IsForDebug) const {
5034+
bool ShouldPreserveUseListOrder, bool IsForDebug) const {
50095035
SlotTracker SlotTable(this->getParent());
50105036
formatted_raw_ostream OS(ROS);
5011-
AssemblyWriter W(OS, SlotTable, this->getParent(), AAW,
5012-
IsForDebug,
5037+
AssemblyWriter W(OS, SlotTable, this->getParent(), AAW, IsForDebug,
50135038
ShouldPreserveUseListOrder);
50145039
W.printFunction(this);
50155040
}

‎llvm/lib/IR/Intrinsics.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,7 @@
3232
#include "llvm/IR/IntrinsicsX86.h"
3333
#include "llvm/IR/IntrinsicsXCore.h"
3434
#include "llvm/IR/Module.h"
35+
#include "llvm/IR/NVVMIntrinsicUtils.h"
3536
#include "llvm/IR/Type.h"
3637

3738
using namespace llvm;
@@ -601,6 +602,12 @@ bool Intrinsic::isOverloaded(ID id) {
601602
#undef GET_INTRINSIC_OVERLOAD_TABLE
602603
}
603604

605+
bool Intrinsic::hasPrettyPrintedArgs(ID id){
606+
#define GET_INTRINSIC_PRETTY_PRINT_TABLE
607+
#include "llvm/IR/IntrinsicImpl.inc"
608+
#undef GET_INTRINSIC_PRETTY_PRINT_TABLE
609+
}
610+
604611
/// Table of per-target intrinsic name tables.
605612
#define GET_INTRINSIC_TARGET_DATA
606613
#include "llvm/IR/IntrinsicImpl.inc"
@@ -1129,3 +1136,7 @@ Intrinsic::ID Intrinsic::getDeinterleaveIntrinsicID(unsigned Factor) {
11291136
assert(Factor >= 2 && Factor <= 8 && "Unexpected factor");
11301137
return InterleaveIntrinsics[Factor - 2].Deinterleave;
11311138
}
1139+
1140+
#define GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
1141+
#include "llvm/IR/IntrinsicImpl.inc"
1142+
#undef GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS

‎llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll‎

Lines changed: 41 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,41 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
2+
; NOTE: This sample test demonstrates the --print-formatted-intrinsics feature for NVPTX intrinsics
3+
; RUN: llvm-as < %s | llvm-dis --print-formatted-intrinsics | FileCheck %s
4+
5+
target triple = "nvptx64-nvidia-cuda"
6+
7+
define void @tcgen05_mma_fp16_cta1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d) {
8+
; CHECK-LABEL: define void @tcgen05_mma_fp16_cta1(
9+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f16 */ i32 0, /* cta_group= */ i32 1, /* collector=discard */ i32 0)
10+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 0, i32 1, i32 0)
11+
12+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f16 */ i32 0, /* cta_group= */ i32 1, /* collector=lastuse */ i32 1)
13+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 0, i32 1, i32 1)
14+
15+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f16 */ i32 0, /* cta_group= */ i32 1, /* collector=fill */ i32 2)
16+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 0, i32 1, i32 2)
17+
18+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f16 */ i32 0, /* cta_group= */ i32 1, /* collector=use */ i32 3)
19+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 0, i32 1, i32 3)
20+
21+
ret void
22+
}
23+
24+
define void @tcgen05_mma_f8f6f4_cta2(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d) {
25+
; CHECK-LABEL: define void @tcgen05_mma_f8f6f4_cta2(
26+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 2, /* collector=discard */ i32 0)
27+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 2, i32 2, i32 0)
28+
29+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 2, /* collector=lastuse */ i32 1)
30+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 2, i32 2, i32 1)
31+
32+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 2, /* collector=fill */ i32 2)
33+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 2, i32 2, i32 2)
34+
35+
; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 2, /* collector=use */ i32 3)
36+
call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d, i32 2, i32 2, i32 3)
37+
38+
ret void
39+
}
40+
41+
declare void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6), ptr addrspace(6), i64, i32, i1, i32, i32, i32)

‎llvm/test/TableGen/intrinsic-arginfo.td‎

Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
1+
// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s | FileCheck %s
2+
3+
// Test ArgInfo property for pretty-printing intrinsic arguments.
4+
// This test verifies that TableGen generates the correct pretty-printing code
5+
// for intrinsics that use the ArgInfo property.
6+
7+
include "llvm/IR/Intrinsics.td"
8+
9+
// Simple intrinsic with two arguments that have ArgInfo.
10+
def int_dummy_foo_bar : DefaultAttrsIntrinsic<
11+
[llvm_i32_ty],
12+
[llvm_i32_ty, // data
13+
llvm_i32_ty, // mode
14+
llvm_i32_ty], // stride
15+
[IntrNoMem,
16+
ImmArg<ArgIndex<1>>,
17+
ArgInfo<ArgIndex<1>, "mode", "printDummyMode">,
18+
ArgInfo<ArgIndex<2>, "stride">]>;
19+
20+
// A custom floating point add with rounding and sat mode.
21+
def int_my_fadd_f32 : DefaultAttrsIntrinsic<
22+
[llvm_float_ty],
23+
[llvm_float_ty, // a
24+
llvm_float_ty, // b
25+
llvm_i32_ty, // rounding_mode
26+
llvm_i1_ty], // saturation_mode
27+
[IntrNoMem,
28+
ImmArg<ArgIndex<2>>,
29+
ImmArg<ArgIndex<3>>,
30+
ArgInfo<ArgIndex<2>, "rounding_mode", "printRoundingMode">,
31+
ArgInfo<ArgIndex<3>, "saturation_mode">]>;
32+
33+
// CHECK: #ifdef GET_INTRINSIC_PRETTY_PRINT_TABLE
34+
// CHECK-NEXT: static constexpr uint8_t PPTable[] = {
35+
36+
// CHECK: #endif // GET_INTRINSIC_PRETTY_PRINT_TABLE
37+
38+
// CHECK: #ifdef GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
39+
// CHECK-NEXT: void Intrinsic::printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS, const Constant *ImmArgVal) {
40+
41+
// CHECK: case dummy_foo_bar:
42+
// CHECK-NEXT: switch (ArgIdx) {
43+
44+
// CHECK-NEXT: case 1:
45+
// CHECK-NEXT: OS << "mode=";
46+
// CHECK-NEXT: printDummyMode(OS, ImmArgVal);
47+
// CHECK-NEXT: return;
48+
49+
// CHECK-NEXT: case 2:
50+
// CHECK-NEXT: OS << "stride=";
51+
// CHECK-NEXT: return;
52+
53+
// CHECK-NEXT: }
54+
// CHECK-NEXT: break;
55+
56+
// CHECK: case my_fadd_f32:
57+
// CHECK-NEXT: switch (ArgIdx) {
58+
59+
// CHECK-NEXT: case 2:
60+
// CHECK-NEXT: OS << "rounding_mode=";
61+
// CHECK-NEXT: printRoundingMode(OS, ImmArgVal);
62+
// CHECK-NEXT: return;
63+
64+
// CHECK-NEXT: case 3:
65+
// CHECK-NEXT: OS << "saturation_mode=";
66+
// CHECK-NEXT: return;
67+
68+
// CHECK-NEXT: }
69+
// CHECK-NEXT: break;
70+
71+
// CHECK: #endif // GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS

‎llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -441,6 +441,14 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
441441
int64_t Lower = R->getValueAsInt("Lower");
442442
int64_t Upper = R->getValueAsInt("Upper");
443443
addArgAttribute(ArgNo, Range, Lower, Upper);
444+
} else if (R->isSubClassOf("ArgInfo")) {
445+
unsigned ArgNo = R->getValueAsInt("ArgNo");
446+
if (ArgNo < 1)
447+
PrintFatalError(R->getLoc(),
448+
"ArgInfo requires ArgNo >= 1 (0 is return value)");
449+
StringRef ArgName = R->getValueAsString("ArgName");
450+
StringRef FuncName = R->getValueAsString("FunctionName");
451+
addPrettyPrintFunction(ArgNo - 1, ArgName, FuncName);
444452
} else {
445453
llvm_unreachable("Unknown property!");
446454
}
@@ -468,3 +476,18 @@ void CodeGenIntrinsic::addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V,
468476
ArgumentAttributes.resize(Idx + 1);
469477
ArgumentAttributes[Idx].emplace_back(AK, V, V2);
470478
}
479+
480+
void CodeGenIntrinsic::addPrettyPrintFunction(unsigned ArgIdx,
481+
StringRef ArgName,
482+
StringRef FuncName) {
483+
auto It = llvm::find_if(PrettyPrintFunctions, [ArgIdx](const auto &Info) {
484+
return Info.ArgIdx == ArgIdx;
485+
});
486+
if (It != PrettyPrintFunctions.end()) {
487+
PrintFatalError(TheDef->getLoc(), "ArgInfo for argument " + Twine(ArgIdx) +
488+
" is already defined as '" +
489+
It->FuncName + "'");
490+
return;
491+
}
492+
PrettyPrintFunctions.emplace_back(ArgIdx, ArgName, FuncName);
493+
}

0 commit comments

Comments
 (0)