[llvm] 39e7712 - [LLVM-Tablegen] Pretty Printing Arguments in LLVM Intrinsics (#162629)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Nov 17 09:00:44 PST 2025
Author: Dharuni R Acharya
Date: 2025-11-17T09:00:40-08:00
New Revision: 39e7712ac520ccfc43383b3e9d6ea8cf2958b8e3
URL: https://github.com/llvm/llvm-project/commit/39e7712ac520ccfc43383b3e9d6ea8cf2958b8e3
DIFF: https://github.com/llvm/llvm-project/commit/39e7712ac520ccfc43383b3e9d6ea8cf2958b8e3.diff
LOG: [LLVM-Tablegen] Pretty Printing Arguments in LLVM Intrinsics (#162629)
This patch adds LLVM infrastructure to support pretty printing of the
intrinsic arguments.
The motivation is to improve the readability of LLVM intrinsics and
facilitate easy
modifications and debugging of LLVM IR.
This feature adds a property `ArgInfo<ArgIndex, [ArgName<"argName">,
ImmArgPrinter<"functionName">]>`
to the intrinsic arguments to print self-explanatory inline comments for
the arguments.
The addition of pretty print support can provide a simple, low-overhead
feature that
enhances the usability of LLVM intrinsics without disrupting existing
workflows.
Link to the RFC, where this feature was discussed:
https://discourse.llvm.org/t/rfc-pretty-printing-immediate-arguments-in-llvm-intrinsics/88536
---------
Signed-off-by: Dharuni R Acharya <dharunira at nvidia.com>
Co-authored-by: Rahul Joshi <rjoshi at nvidia.com>
Added:
llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll
llvm/test/TableGen/intrinsic-arginfo.td
Modified:
llvm/include/llvm/IR/Intrinsics.h
llvm/include/llvm/IR/Intrinsics.td
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/include/llvm/IR/NVVMIntrinsicUtils.h
llvm/lib/IR/AsmWriter.cpp
llvm/lib/IR/Intrinsics.cpp
llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
Removed:
################################################################################
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index 9577d0141f168..c91fc254ebe11 100644
--- a/llvm/include/llvm/IR/Intrinsics.h
+++ b/llvm/include/llvm/IR/Intrinsics.h
@@ -30,6 +30,8 @@ class LLVMContext;
class Module;
class AttributeList;
class AttributeSet;
+class raw_ostream;
+class Constant;
/// This namespace contains an enum with a value for every intrinsic/builtin
/// function known by LLVM. The enum values are returned by
@@ -81,6 +83,9 @@ namespace Intrinsic {
/// Returns true if the intrinsic can be overloaded.
LLVM_ABI bool isOverloaded(ID id);
+ /// Returns true if the intrinsic has pretty printed immediate arguments.
+ LLVM_ABI bool hasPrettyPrintedArgs(ID id);
+
/// 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);
@@ -284,6 +289,10 @@ namespace Intrinsic {
/// N.
LLVM_ABI Intrinsic::ID getDeinterleaveIntrinsicID(unsigned Factor);
+ /// Print the argument info for the arguments with ArgInfo.
+ LLVM_ABI void printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS,
+ const Constant *ImmArgVal);
+
} // namespace Intrinsic
} // namespace llvm
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index 07aa2faffa7c5..27f404a1be65c 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -142,6 +142,25 @@ class Range<AttrIndex idx, int lower, int upper> : IntrinsicProperty {
int Upper = upper;
}
+// ArgProperty - Base class for argument properties that can be specified in ArgInfo.
+class ArgProperty;
+
+// ArgName - Specifies the name of an argument for pretty-printing.
+class ArgName<string name> : ArgProperty {
+ string Name = name;
+}
+
+// ImmArgPrinter - Specifies a custom printer function for immediate arguments.
+class ImmArgPrinter<string funcname> : ArgProperty {
+ string FuncName = funcname;
+}
+
+// ArgInfo - The specified argument has properties defined by a list of ArgProperty objects.
+class ArgInfo<ArgIndex idx, list<ArgProperty> arg_properties> : IntrinsicProperty {
+ int ArgNo = idx.Value;
+ list<ArgProperty> Properties = arg_properties;
+}
+
def IntrNoReturn : IntrinsicProperty;
// Applied by default.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 21badc2692037..1b485dc8ccd1e 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2955,7 +2955,14 @@ foreach sp = [0, 1] in {
defvar nargs = !size(args);
defvar scale_d_imm = ArgIndex<!sub(nargs, 1)>;
defvar scale_d_imm_range = [ImmArg<scale_d_imm>, Range<scale_d_imm, 0, 16>];
- defvar intrinsic_properties = !listconcat(
+
+ // Check if this is the specific llvm.nvvm.tcgen05.mma.tensor intrinsic.
+ defvar is_target_intrinsic = !and(!eq(sp, 0),
+ !eq(space, "tensor"),
+ !eq(scale_d, 0),
+ !eq(ashift, 0));
+
+ defvar base_properties = !listconcat(
mma.common_intr_props,
!if(!eq(scale_d, 1), scale_d_imm_range, []),
[Range<ArgIndex<nargs>, 0, !if(!eq(scale_d, 1), 2, 4)>, // kind
@@ -2965,6 +2972,13 @@ foreach sp = [0, 1] in {
]
);
+ defvar intrinsic_properties = !if(is_target_intrinsic,
+ !listconcat(base_properties,
+ [ArgInfo<ArgIndex<nargs>, [ArgName<"kind">, ImmArgPrinter<"printTcgen05MMAKind">]>,
+ ArgInfo<ArgIndex<!add(nargs, 1)>, [ArgName<"cta_group">]>,
+ ArgInfo<ArgIndex<!add(nargs, 2)>, [ArgName<"collector">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]),
+ base_properties);
+
def mma.record_name:
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
mma.intr_name>;
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index d55100e5e709d..d383769043605 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -18,8 +18,11 @@
#include <stdint.h>
#include "llvm/ADT/APFloat.h"
+#include "llvm/ADT/APInt.h"
+#include "llvm/IR/Constants.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
+#include "llvm/Support/raw_ostream.h"
namespace llvm {
namespace nvvm {
@@ -659,6 +662,51 @@ inline APFloat::roundingMode GetFMARoundingMode(Intrinsic::ID IntrinsicID) {
llvm_unreachable("Invalid FP instrinsic rounding mode for NVVM fma");
}
+inline void printTcgen05MMAKind(raw_ostream &OS, const Constant *ImmArgVal) {
+ if (const auto *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
+ uint64_t Val = CI->getZExtValue();
+ switch (static_cast<Tcgen05MMAKind>(Val)) {
+ case Tcgen05MMAKind::F16:
+ OS << "f16";
+ return;
+ case Tcgen05MMAKind::TF32:
+ OS << "tf32";
+ return;
+ case Tcgen05MMAKind::F8F6F4:
+ OS << "f8f6f4";
+ return;
+ case Tcgen05MMAKind::I8:
+ OS << "i8";
+ return;
+ }
+ }
+ llvm_unreachable(
+ "printTcgen05MMAKind called with invalid value for immediate argument");
+}
+
+inline void printTcgen05CollectorUsageOp(raw_ostream &OS,
+ const Constant *ImmArgVal) {
+ if (const auto *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
+ uint64_t Val = CI->getZExtValue();
+ switch (static_cast<Tcgen05CollectorUsageOp>(Val)) {
+ case Tcgen05CollectorUsageOp::DISCARD:
+ OS << "discard";
+ return;
+ case Tcgen05CollectorUsageOp::LASTUSE:
+ OS << "lastuse";
+ return;
+ case Tcgen05CollectorUsageOp::FILL:
+ OS << "fill";
+ return;
+ case Tcgen05CollectorUsageOp::USE:
+ OS << "use";
+ return;
+ }
+ }
+ llvm_unreachable("printTcgen05CollectorUsageOp called with invalid value for "
+ "immediate argument");
+}
+
} // namespace nvvm
} // namespace llvm
#endif // LLVM_IR_NVVMINTRINSICUTILS_H
diff --git a/llvm/lib/IR/AsmWriter.cpp b/llvm/lib/IR/AsmWriter.cpp
index 4d4ffe93a8067..94a1aa3087377 100644
--- a/llvm/lib/IR/AsmWriter.cpp
+++ b/llvm/lib/IR/AsmWriter.cpp
@@ -53,6 +53,7 @@
#include "llvm/IR/Instruction.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/IntrinsicInst.h"
+#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/Metadata.h"
#include "llvm/IR/Module.h"
@@ -4576,12 +4577,38 @@ void AssemblyWriter::printInstruction(const Instruction &I) {
Out << ' ';
writeOperand(Operand, false);
Out << '(';
+ bool HasPrettyPrintedArgs =
+ isa<IntrinsicInst>(CI) &&
+ Intrinsic::hasPrettyPrintedArgs(CI->getIntrinsicID());
+
ListSeparator LS;
- for (unsigned op = 0, Eop = CI->arg_size(); op < Eop; ++op) {
- Out << LS;
- writeParamOperand(CI->getArgOperand(op), PAL.getParamAttrs(op));
+ Function *CalledFunc = CI->getCalledFunction();
+ auto PrintArgComment = [&](unsigned ArgNo) {
+ const auto *ConstArg = dyn_cast<Constant>(CI->getArgOperand(ArgNo));
+ if (!ConstArg)
+ return;
+ std::string ArgComment;
+ raw_string_ostream ArgCommentStream(ArgComment);
+ Intrinsic::ID IID = CalledFunc->getIntrinsicID();
+ Intrinsic::printImmArg(IID, ArgNo, ArgCommentStream, ConstArg);
+ if (ArgComment.empty())
+ return;
+ Out << "/* " << ArgComment << " */ ";
+ };
+ if (HasPrettyPrintedArgs) {
+ for (unsigned ArgNo = 0, NumArgs = CI->arg_size(); ArgNo < NumArgs;
+ ++ArgNo) {
+ Out << LS;
+ PrintArgComment(ArgNo);
+ writeParamOperand(CI->getArgOperand(ArgNo), PAL.getParamAttrs(ArgNo));
+ }
+ } else {
+ for (unsigned ArgNo = 0, NumArgs = CI->arg_size(); ArgNo < NumArgs;
+ ++ArgNo) {
+ Out << LS;
+ writeParamOperand(CI->getArgOperand(ArgNo), PAL.getParamAttrs(ArgNo));
+ }
}
-
// Emit an ellipsis if this is a musttail call in a vararg function. This
// is only to aid readability, musttail calls forward varargs by default.
if (CI->isMustTailCall() && CI->getParent() &&
@@ -5005,12 +5032,10 @@ void AssemblyWriter::printUseLists(const Function *F) {
//===----------------------------------------------------------------------===//
void Function::print(raw_ostream &ROS, AssemblyAnnotationWriter *AAW,
- bool ShouldPreserveUseListOrder,
- bool IsForDebug) const {
+ bool ShouldPreserveUseListOrder, bool IsForDebug) const {
SlotTracker SlotTable(this->getParent());
formatted_raw_ostream OS(ROS);
- AssemblyWriter W(OS, SlotTable, this->getParent(), AAW,
- IsForDebug,
+ AssemblyWriter W(OS, SlotTable, this->getParent(), AAW, IsForDebug,
ShouldPreserveUseListOrder);
W.printFunction(this);
}
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index 526800e217399..859689b9cf168 100644
--- a/llvm/lib/IR/Intrinsics.cpp
+++ b/llvm/lib/IR/Intrinsics.cpp
@@ -32,6 +32,7 @@
#include "llvm/IR/IntrinsicsX86.h"
#include "llvm/IR/IntrinsicsXCore.h"
#include "llvm/IR/Module.h"
+#include "llvm/IR/NVVMIntrinsicUtils.h"
#include "llvm/IR/Type.h"
using namespace llvm;
@@ -601,6 +602,12 @@ bool Intrinsic::isOverloaded(ID id) {
#undef GET_INTRINSIC_OVERLOAD_TABLE
}
+bool Intrinsic::hasPrettyPrintedArgs(ID id){
+#define GET_INTRINSIC_PRETTY_PRINT_TABLE
+#include "llvm/IR/IntrinsicImpl.inc"
+#undef GET_INTRINSIC_PRETTY_PRINT_TABLE
+}
+
/// Table of per-target intrinsic name tables.
#define GET_INTRINSIC_TARGET_DATA
#include "llvm/IR/IntrinsicImpl.inc"
@@ -1142,3 +1149,7 @@ Intrinsic::ID Intrinsic::getDeinterleaveIntrinsicID(unsigned Factor) {
assert(Factor >= 2 && Factor <= 8 && "Unexpected factor");
return InterleaveIntrinsics[Factor - 2].Deinterleave;
}
+
+#define GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
+#include "llvm/IR/IntrinsicImpl.inc"
+#undef GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
diff --git a/llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll b/llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll
new file mode 100644
index 0000000000000..479de53dd90f2
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll
@@ -0,0 +1,50 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; NOTE: This sample test demonstrates the pretty print feature for NVPTX intrinsics
+; RUN: llvm-as < %s | llvm-dis | FileCheck %s
+
+target triple = "nvptx64-nvidia-cuda"
+
+define void @tcgen05_mma_fp16_cta1(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d) {
+ ; CHECK-LABEL: define void @tcgen05_mma_fp16_cta1(
+ ; 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)
+ 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)
+
+ ; 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)
+ 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)
+
+ ; 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)
+ 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)
+
+ ; 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)
+ 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)
+
+ ret void
+}
+
+define void @tcgen05_mma_f8f6f4_cta2(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 %b, i32 %idesc, i1 %enable_inp_d) {
+ ; CHECK-LABEL: define void @tcgen05_mma_f8f6f4_cta2(
+ ; 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)
+ 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)
+
+ ; 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)
+ 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)
+
+ ; 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)
+ 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)
+
+ ; 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)
+ 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)
+
+ ret void
+}
+
+; This test verifies that printImmArg is safe to call on all constant arguments, but only prints comments for arguments that have pretty printing configured.
+define void @test_mixed_constants_edge_case(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor) {
+ ; CHECK-LABEL: define void @test_mixed_constants_edge_case(
+ ; CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 42, i32 100, i1 true, /* kind=i8 */ i32 3, /* cta_group= */ i32 1, /* collector=discard */ i32 0)
+ call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor, i64 42, i32 100, i1 true, i32 3, i32 1, i32 0)
+
+ ret void
+}
+
+declare void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6), ptr addrspace(6), i64, i32, i1, i32, i32, i32)
diff --git a/llvm/test/TableGen/intrinsic-arginfo.td b/llvm/test/TableGen/intrinsic-arginfo.td
new file mode 100644
index 0000000000000..eab1f5e032bc3
--- /dev/null
+++ b/llvm/test/TableGen/intrinsic-arginfo.td
@@ -0,0 +1,71 @@
+// RUN: llvm-tblgen -gen-intrinsic-impl -I %p/../../include %s | FileCheck %s
+
+// Test ArgInfo property for pretty-printing intrinsic arguments.
+// This test verifies that TableGen generates the correct pretty-printing code
+// for intrinsics that use the ArgInfo property.
+
+include "llvm/IR/Intrinsics.td"
+
+// Simple intrinsic with two arguments that have ArgInfo.
+def int_dummy_foo_bar : DefaultAttrsIntrinsic<
+ [llvm_i32_ty],
+ [llvm_i32_ty, // data
+ llvm_i32_ty, // mode
+ llvm_i32_ty], // stride
+ [IntrNoMem,
+ ImmArg<ArgIndex<1>>,
+ ArgInfo<ArgIndex<1>, [ArgName<"mode">, ImmArgPrinter<"printDummyMode">]>,
+ ArgInfo<ArgIndex<2>, [ArgName<"stride">]>]>;
+
+// A custom floating point add with rounding and sat mode.
+def int_my_fadd_f32 : DefaultAttrsIntrinsic<
+ [llvm_float_ty],
+ [llvm_float_ty, // a
+ llvm_float_ty, // b
+ llvm_i32_ty, // rounding_mode
+ llvm_i1_ty], // saturation_mode
+ [IntrNoMem,
+ ImmArg<ArgIndex<2>>,
+ ImmArg<ArgIndex<3>>,
+ ArgInfo<ArgIndex<2>, [ArgName<"rounding_mode">, ImmArgPrinter<"printRoundingMode">]>,
+ ArgInfo<ArgIndex<3>, [ArgName<"saturation_mode">]>]>;
+
+// CHECK: #ifdef GET_INTRINSIC_PRETTY_PRINT_TABLE
+// CHECK-NEXT: static constexpr uint8_t PPTable[] = {
+
+// CHECK: #endif // GET_INTRINSIC_PRETTY_PRINT_TABLE
+
+// CHECK: #ifdef GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
+// CHECK: void Intrinsic::printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS, const Constant *ImmArgVal) {
+
+// CHECK: case dummy_foo_bar:
+// CHECK-NEXT: switch (ArgIdx) {
+
+// CHECK-NEXT: case 1:
+// CHECK-NEXT: OS << "mode=";
+// CHECK-NEXT: printDummyMode(OS, ImmArgVal);
+// CHECK-NEXT: return;
+
+// CHECK-NEXT: case 2:
+// CHECK-NEXT: OS << "stride=";
+// CHECK-NEXT: return;
+
+// CHECK-NEXT: }
+// CHECK-NEXT: break;
+
+// CHECK: case my_fadd_f32:
+// CHECK-NEXT: switch (ArgIdx) {
+
+// CHECK-NEXT: case 2:
+// CHECK-NEXT: OS << "rounding_mode=";
+// CHECK-NEXT: printRoundingMode(OS, ImmArgVal);
+// CHECK-NEXT: return;
+
+// CHECK-NEXT: case 3:
+// CHECK-NEXT: OS << "saturation_mode=";
+// CHECK-NEXT: return;
+
+// CHECK-NEXT: }
+// CHECK-NEXT: break;
+
+// CHECK: #endif // GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index ff894853b9771..228969ab37f85 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -449,6 +449,29 @@ void CodeGenIntrinsic::setProperty(const Record *R) {
int64_t Lower = R->getValueAsInt("Lower");
int64_t Upper = R->getValueAsInt("Upper");
addArgAttribute(ArgNo, Range, Lower, Upper);
+ } else if (R->isSubClassOf("ArgInfo")) {
+ unsigned ArgNo = R->getValueAsInt("ArgNo");
+ if (ArgNo < 1)
+ PrintFatalError(R->getLoc(),
+ "ArgInfo requires ArgNo >= 1 (0 is return value)");
+ const ListInit *Properties = R->getValueAsListInit("Properties");
+ StringRef ArgName;
+ StringRef FuncName;
+
+ for (const Init *PropInit : Properties->getElements()) {
+ if (const auto *PropDef = dyn_cast<DefInit>(PropInit)) {
+ const Record *PropRec = PropDef->getDef();
+
+ if (PropRec->isSubClassOf("ArgName"))
+ ArgName = PropRec->getValueAsString("Name");
+ else if (PropRec->isSubClassOf("ImmArgPrinter"))
+ FuncName = PropRec->getValueAsString("FuncName");
+ else
+ PrintFatalError(PropRec->getLoc(),
+ "Unknown ArgProperty type: " + PropRec->getName());
+ }
+ }
+ addPrettyPrintFunction(ArgNo - 1, ArgName, FuncName);
} else {
llvm_unreachable("Unknown property!");
}
@@ -476,3 +499,16 @@ void CodeGenIntrinsic::addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V,
ArgumentAttributes.resize(Idx + 1);
ArgumentAttributes[Idx].emplace_back(AK, V, V2);
}
+
+void CodeGenIntrinsic::addPrettyPrintFunction(unsigned ArgIdx,
+ StringRef ArgName,
+ StringRef FuncName) {
+ auto It = llvm::find_if(PrettyPrintFunctions, [ArgIdx](const auto &Info) {
+ return Info.ArgIdx == ArgIdx;
+ });
+ if (It != PrettyPrintFunctions.end())
+ PrintFatalError(TheDef->getLoc(), "ArgInfo for argument " + Twine(ArgIdx) +
+ " is already defined as '" +
+ It->FuncName + "'");
+ PrettyPrintFunctions.emplace_back(ArgIdx, ArgName, FuncName);
+}
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index 15e803c4feba1..6ac6f734326d8 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -152,6 +152,22 @@ struct CodeGenIntrinsic {
void addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V = 0,
uint64_t V2 = 0);
+ /// Structure to store pretty print and argument information.
+ struct PrettyPrintArgInfo {
+ unsigned ArgIdx;
+ StringRef ArgName;
+ StringRef FuncName;
+
+ PrettyPrintArgInfo(unsigned Idx, StringRef Name, StringRef Func)
+ : ArgIdx(Idx), ArgName(Name), FuncName(Func) {}
+ };
+
+ /// Vector that stores ArgInfo (ArgIndex, ArgName, FunctionName).
+ SmallVector<PrettyPrintArgInfo> PrettyPrintFunctions;
+
+ void addPrettyPrintFunction(unsigned ArgIdx, StringRef ArgName,
+ StringRef FuncName);
+
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 452d2b08f25c3..3ac23185ef91c 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -60,8 +60,16 @@ class IntrinsicEmitter {
raw_ostream &OS);
void EmitIntrinsicToOverloadTable(const CodeGenIntrinsicTable &Ints,
raw_ostream &OS);
+ void EmitIntrinsicToPrettyPrintTable(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS);
+ void EmitIntrinsicBitTable(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS, StringRef Guard,
+ StringRef TableName, StringRef Comment,
+ function_ref<bool(const CodeGenIntrinsic &Int)> GetProperty);
void EmitGenerator(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
void EmitAttributes(const CodeGenIntrinsicTable &Ints, raw_ostream &OS);
+ void EmitPrettyPrintArguments(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS);
void EmitIntrinsicToBuiltinMap(const CodeGenIntrinsicTable &Ints,
bool IsClang, raw_ostream &OS);
};
@@ -109,6 +117,12 @@ void IntrinsicEmitter::run(raw_ostream &OS, bool Enums) {
// Emit the intrinsic parameter attributes.
EmitAttributes(Ints, OS);
+ // Emit the intrinsic ID -> pretty print table.
+ EmitIntrinsicToPrettyPrintTable(Ints, OS);
+
+ // Emit Pretty Print attribute.
+ EmitPrettyPrintArguments(Ints, OS);
+
// Emit code to translate Clang builtins into LLVM intrinsics.
EmitIntrinsicToBuiltinMap(Ints, true, OS);
@@ -240,6 +254,29 @@ static constexpr IntrinsicTargetInfo TargetInfos[] = {
)";
}
+/// Helper function to emit a bit table for intrinsic properties.
+/// This is used for both overload and pretty print bit tables.
+void IntrinsicEmitter::EmitIntrinsicBitTable(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS, StringRef Guard,
+ StringRef TableName, StringRef Comment,
+ function_ref<bool(const CodeGenIntrinsic &Int)> GetProperty) {
+ OS << formatv("// {}\n", Comment);
+ OS << formatv("#ifdef {}\n", Guard);
+ OS << formatv("static constexpr uint8_t {}[] = {{\n", TableName);
+ OS << " 0\n ";
+ for (auto [I, Int] : enumerate(Ints)) {
+ // Add one to the index so we emit a null bit for the invalid #0 intrinsic.
+ size_t Idx = I + 1;
+ if (Idx % 8 == 0)
+ OS << ",\n 0";
+ if (GetProperty(Int))
+ OS << " | (1<<" << Idx % 8 << ')';
+ }
+ OS << "\n};\n\n";
+ OS << formatv("return ({}[id/8] & (1 << (id%8))) != 0;\n", TableName);
+ OS << formatv("#endif // {}\n\n", Guard);
+}
+
void IntrinsicEmitter::EmitIntrinsicToNameTable(
const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
// Built up a table of the intrinsic names.
@@ -276,24 +313,10 @@ static constexpr unsigned IntrinsicNameOffsetTable[] = {
void IntrinsicEmitter::EmitIntrinsicToOverloadTable(
const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
- OS << R"(// Intrinsic ID to overload bitset.
-#ifdef GET_INTRINSIC_OVERLOAD_TABLE
-static constexpr uint8_t OTable[] = {
- 0
- )";
- for (auto [I, Int] : enumerate(Ints)) {
- // Add one to the index so we emit a null bit for the invalid #0 intrinsic.
- size_t Idx = I + 1;
-
- if (Idx % 8 == 0)
- OS << ",\n 0";
- if (Int.isOverloaded)
- OS << " | (1<<" << Idx % 8 << ')';
- }
- OS << "\n};\n\n";
- // OTable contains a true bit at the position if the intrinsic is overloaded.
- OS << "return (OTable[id/8] & (1 << (id%8))) != 0;\n";
- OS << "#endif\n\n";
+ EmitIntrinsicBitTable(
+ Ints, OS, "GET_INTRINSIC_OVERLOAD_TABLE", "OTable",
+ "Intrinsic ID to overload bitset.",
+ [](const CodeGenIntrinsic &Int) { return Int.isOverloaded; });
}
using TypeSigTy = SmallVector<unsigned char>;
@@ -809,6 +832,52 @@ AttributeSet Intrinsic::getFnAttributes(LLVMContext &C, ID id) {{
NoFunctionAttrsID);
}
+void IntrinsicEmitter::EmitIntrinsicToPrettyPrintTable(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
+ EmitIntrinsicBitTable(Ints, OS, "GET_INTRINSIC_PRETTY_PRINT_TABLE", "PPTable",
+ "Intrinsic ID to pretty print bitset.",
+ [](const CodeGenIntrinsic &Int) {
+ return !Int.PrettyPrintFunctions.empty();
+ });
+}
+
+void IntrinsicEmitter::EmitPrettyPrintArguments(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
+ OS << R"(
+#ifdef GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
+void Intrinsic::printImmArg(ID IID, unsigned ArgIdx, raw_ostream &OS, const Constant *ImmArgVal) {
+ using namespace Intrinsic;
+ switch (IID) {
+)";
+
+ for (const CodeGenIntrinsic &Int : Ints) {
+ if (Int.PrettyPrintFunctions.empty())
+ continue;
+
+ OS << " case " << Int.EnumName << ":\n";
+ OS << " switch (ArgIdx) {\n";
+ for (const auto [ArgIdx, ArgName, FuncName] : Int.PrettyPrintFunctions) {
+ OS << " case " << ArgIdx << ":\n";
+ OS << " OS << \"" << ArgName << "=\";\n";
+ if (!FuncName.empty()) {
+ OS << " ";
+ if (!Int.TargetPrefix.empty())
+ OS << Int.TargetPrefix << "::";
+ OS << FuncName << "(OS, ImmArgVal);\n";
+ }
+ OS << " return;\n";
+ }
+ OS << " }\n";
+ OS << " break;\n";
+ }
+ OS << R"( default:
+ break;
+ }
+}
+#endif // GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
+)";
+}
+
void IntrinsicEmitter::EmitIntrinsicToBuiltinMap(
const CodeGenIntrinsicTable &Ints, bool IsClang, raw_ostream &OS) {
StringRef CompilerName = IsClang ? "Clang" : "MS";
More information about the llvm-commits
mailing list