[llvm] [LLVM-Tablegen] Pretty Printing Immediate Arguments in LLVM Intrinsics (PR #162629)
Dharuni R Acharya via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 9 22:12:03 PDT 2025
https://github.com/DharuniRAcharya updated https://github.com/llvm/llvm-project/pull/162629
>From 9119a56efa7bd012e33a1e90a58e28b7a2480552 Mon Sep 17 00:00:00 2001
From: Dharuni R Acharya <dharunira at nvidia.com>
Date: Thu, 9 Oct 2025 10:07:02 +0000
Subject: [PATCH] [LLVM-Tablegen] Pretty Printing Immediate Arguments in LLVM
Intrinsics
This patch adds LLVM infrastructure to support pretty printing of immediate 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 PrettyPrintImmArg<ArgIndex, "functionName"> to the intrinsic ImmArgs that enables printing self-explanatory inline comment for the immediate 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 <dharunira at nvidia.com>
---
llvm/include/llvm/IR/Function.h | 4 +-
llvm/include/llvm/IR/Intrinsics.h | 8 +++
llvm/include/llvm/IR/Intrinsics.td | 7 ++
llvm/include/llvm/IR/IntrinsicsNVVM.td | 16 ++++-
llvm/include/llvm/IR/Module.h | 4 +-
llvm/include/llvm/IR/NVVMIntrinsicUtils.h | 67 ++++++++++++++++++
llvm/lib/IR/AsmWriter.cpp | 64 +++++++++++++----
llvm/lib/IR/Intrinsics.cpp | 11 +++
.../NVPTX/tcgen05-mma-tensor-formatted.ll | 41 +++++++++++
llvm/tools/llvm-dis/llvm-dis.cpp | 8 ++-
.../TableGen/Basic/CodeGenIntrinsics.cpp | 22 ++++++
llvm/utils/TableGen/Basic/CodeGenIntrinsics.h | 5 ++
.../utils/TableGen/Basic/IntrinsicEmitter.cpp | 70 +++++++++++++++++++
13 files changed, 306 insertions(+), 21 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll
diff --git a/llvm/include/llvm/IR/Function.h b/llvm/include/llvm/IR/Function.h
index d3497716ca844..4b41dc137d9de 100644
--- a/llvm/include/llvm/IR/Function.h
+++ b/llvm/include/llvm/IR/Function.h
@@ -929,8 +929,8 @@ class LLVM_ABI Function : public GlobalObject, public ilist_node<Function> {
/// Print the function to an output stream with an optional
/// AssemblyAnnotationWriter.
void print(raw_ostream &OS, AssemblyAnnotationWriter *AAW = nullptr,
- bool ShouldPreserveUseListOrder = false,
- bool IsForDebug = false) const;
+ bool ShouldPreserveUseListOrder = false, bool IsForDebug = false,
+ bool PrettyPrintIntrinsicArgs = false) const;
/// viewCFG - This function is meant for use from the debugger. You can just
/// say 'call F->viewCFG()' and a ghostview window should pop up from the
diff --git a/llvm/include/llvm/IR/Intrinsics.h b/llvm/include/llvm/IR/Intrinsics.h
index 9577d0141f168..ce1360498341e 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,9 @@ namespace Intrinsic {
/// N.
LLVM_ABI Intrinsic::ID getDeinterleaveIntrinsicID(unsigned Factor);
+ 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 96da698538314..f56a7603f4963 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -87,6 +87,13 @@ class NoUndef<AttrIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
+// PrettyPrintImmArg - The specified immediate argument has a custom pretty-print
+// function for diagnostic output.
+class PrettyPrintImmArg<AttrIndex idx, string funcname> : IntrinsicProperty {
+ int ArgNo = idx.Value;
+ string FunctionName = funcname;
+}
+
// NonNull - The return value or specified argument is not null.
class NonNull<AttrIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 3af1750ffcf3f..2b7db120e75f2 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2871,7 +2871,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
@@ -2881,6 +2888,13 @@ foreach sp = [0, 1] in {
]
);
+ defvar intrinsic_properties = !if(is_target_intrinsic,
+ !listconcat(base_properties,
+ [PrettyPrintImmArg<ArgIndex<nargs>, "printTcgen05MMAKind">,
+ PrettyPrintImmArg<ArgIndex<!add(nargs, 1)>, "printCTAGroupKind">,
+ PrettyPrintImmArg<ArgIndex<!add(nargs, 2)>, "printTcgen05CollectorUsageOp">]),
+ base_properties);
+
def mma.record:
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
mma.intr>;
diff --git a/llvm/include/llvm/IR/Module.h b/llvm/include/llvm/IR/Module.h
index a99937a90cbb7..e74035dc68dc0 100644
--- a/llvm/include/llvm/IR/Module.h
+++ b/llvm/include/llvm/IR/Module.h
@@ -865,8 +865,8 @@ class LLVM_ABI Module {
/// uselistorder directives so that use-lists can be recreated when reading
/// the assembly.
void print(raw_ostream &OS, AssemblyAnnotationWriter *AAW,
- bool ShouldPreserveUseListOrder = false,
- bool IsForDebug = false) const;
+ bool ShouldPreserveUseListOrder = false, bool IsForDebug = false,
+ bool PrettyPrintIntrinsicArgs = false) const;
/// Dump the module to stderr (for debugging).
void dump() const;
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index d55100e5e709d..1f835aa96321a 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,70 @@ 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 ConstantInt *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
+ uint64_t Val = CI->getZExtValue();
+ switch (static_cast<Tcgen05MMAKind>(Val)) {
+ case Tcgen05MMAKind::F16:
+ OS << "kind::f16";
+ return;
+ case Tcgen05MMAKind::TF32:
+ OS << "kind::tf32";
+ return;
+ case Tcgen05MMAKind::F8F6F4:
+ OS << "kind::f8f6f4";
+ return;
+ case Tcgen05MMAKind::I8:
+ OS << "kind::i8";
+ return;
+ }
+ }
+ llvm_unreachable(
+ "printTcgen05MMAKind called with invalid value for immediate argument");
+}
+
+inline void printTcgen05CollectorUsageOp(raw_ostream &OS,
+ const Constant *ImmArgVal) {
+ if (const ConstantInt *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
+ uint64_t Val = CI->getZExtValue();
+ switch (static_cast<Tcgen05CollectorUsageOp>(Val)) {
+ case Tcgen05CollectorUsageOp::DISCARD:
+ OS << "collector::a::discard";
+ return;
+ case Tcgen05CollectorUsageOp::LASTUSE:
+ OS << "collector::a::lastuse";
+ return;
+ case Tcgen05CollectorUsageOp::FILL:
+ OS << "collector::a::fill";
+ return;
+ case Tcgen05CollectorUsageOp::USE:
+ OS << "collector::a::use";
+ return;
+ }
+ }
+ llvm_unreachable("printTcgen05CollectorUsageOp called with invalid value for "
+ "immediate argument");
+}
+
+inline void printCTAGroupKind(raw_ostream &OS, const Constant *ImmArgVal) {
+ if (const ConstantInt *CI = dyn_cast<ConstantInt>(ImmArgVal)) {
+ uint64_t Val = CI->getZExtValue();
+ switch (static_cast<CTAGroupKind>(Val)) {
+ case CTAGroupKind::CG_NONE:
+ OS << "cta_group::0";
+ return;
+ case CTAGroupKind::CG_1:
+ OS << "cta_group::1";
+ return;
+ case CTAGroupKind::CG_2:
+ OS << "cta_group::2";
+ return;
+ }
+ }
+ llvm_unreachable(
+ "printCTAGroupKind 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 ae086bcd3902d..049362c6d6bd0 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"
@@ -2831,6 +2832,7 @@ class AssemblyWriter {
SetVector<const Comdat *> Comdats;
bool IsForDebug;
bool ShouldPreserveUseListOrder;
+ bool PrettyPrintIntrinsicArgs;
UseListOrderMap UseListOrders;
SmallVector<StringRef, 8> MDNames;
/// Synchronization scope names registered with LLVMContext.
@@ -2841,7 +2843,8 @@ class AssemblyWriter {
/// Construct an AssemblyWriter with an external SlotTracker
AssemblyWriter(formatted_raw_ostream &o, SlotTracker &Mac, const Module *M,
AssemblyAnnotationWriter *AAW, bool IsForDebug,
- bool ShouldPreserveUseListOrder = false);
+ bool ShouldPreserveUseListOrder = false,
+ bool PrettyPrintIntrinsicArgs = false);
AssemblyWriter(formatted_raw_ostream &o, SlotTracker &Mac,
const ModuleSummaryIndex *Index, bool IsForDebug);
@@ -2930,10 +2933,12 @@ class AssemblyWriter {
AssemblyWriter::AssemblyWriter(formatted_raw_ostream &o, SlotTracker &Mac,
const Module *M, AssemblyAnnotationWriter *AAW,
- bool IsForDebug, bool ShouldPreserveUseListOrder)
+ bool IsForDebug, bool ShouldPreserveUseListOrder,
+ bool PrettyPrintIntrinsicArgs)
: Out(o), TheModule(M), Machine(Mac), TypePrinter(M), AnnotationWriter(AAW),
IsForDebug(IsForDebug),
- ShouldPreserveUseListOrder(ShouldPreserveUseListOrder) {
+ ShouldPreserveUseListOrder(ShouldPreserveUseListOrder),
+ PrettyPrintIntrinsicArgs(PrettyPrintIntrinsicArgs) {
if (!TheModule)
return;
for (const GlobalObject &GO : TheModule->global_objects())
@@ -2944,7 +2949,8 @@ AssemblyWriter::AssemblyWriter(formatted_raw_ostream &o, SlotTracker &Mac,
AssemblyWriter::AssemblyWriter(formatted_raw_ostream &o, SlotTracker &Mac,
const ModuleSummaryIndex *Index, bool IsForDebug)
: Out(o), TheIndex(Index), Machine(Mac), TypePrinter(/*Module=*/nullptr),
- IsForDebug(IsForDebug), ShouldPreserveUseListOrder(false) {}
+ IsForDebug(IsForDebug), ShouldPreserveUseListOrder(false),
+ PrettyPrintIntrinsicArgs(false) {}
void AssemblyWriter::writeOperand(const Value *Operand, bool PrintType) {
if (!Operand) {
@@ -4561,12 +4567,40 @@ void AssemblyWriter::printInstruction(const Instruction &I) {
Out << ' ';
writeOperand(Operand, false);
Out << '(';
+ bool HasPrettyPrintedArgs =
+ PrettyPrintIntrinsicArgs && 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));
- }
+ if (HasPrettyPrintedArgs) {
+ Function *CalledFunc = CI->getCalledFunction();
+ auto PrintArgComment = [&](unsigned ArgNo) {
+ if (!CalledFunc->hasParamAttribute(ArgNo, Attribute::ImmArg))
+ return;
+ const Constant *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 << " */ ";
+ };
+ 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 op = 0, Eop = CI->arg_size(); op < Eop; ++op) {
+ Out << LS;
+ writeParamOperand(CI->getArgOperand(op), PAL.getParamAttrs(op));
+ }
+ }
// 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() &&
@@ -4990,13 +5024,12 @@ void AssemblyWriter::printUseLists(const Function *F) {
//===----------------------------------------------------------------------===//
void Function::print(raw_ostream &ROS, AssemblyAnnotationWriter *AAW,
- bool ShouldPreserveUseListOrder,
- bool IsForDebug) const {
+ bool ShouldPreserveUseListOrder, bool IsForDebug,
+ bool PrettyPrintIntrinsicArgs) const {
SlotTracker SlotTable(this->getParent());
formatted_raw_ostream OS(ROS);
- AssemblyWriter W(OS, SlotTable, this->getParent(), AAW,
- IsForDebug,
- ShouldPreserveUseListOrder);
+ AssemblyWriter W(OS, SlotTable, this->getParent(), AAW, IsForDebug,
+ ShouldPreserveUseListOrder, PrettyPrintIntrinsicArgs);
W.printFunction(this);
}
@@ -5012,11 +5045,12 @@ void BasicBlock::print(raw_ostream &ROS, AssemblyAnnotationWriter *AAW,
}
void Module::print(raw_ostream &ROS, AssemblyAnnotationWriter *AAW,
- bool ShouldPreserveUseListOrder, bool IsForDebug) const {
+ bool ShouldPreserveUseListOrder, bool IsForDebug,
+ bool PrettyPrintIntrinsicArgs) const {
SlotTracker SlotTable(this);
formatted_raw_ostream OS(ROS);
AssemblyWriter W(OS, SlotTable, this, AAW, IsForDebug,
- ShouldPreserveUseListOrder);
+ ShouldPreserveUseListOrder, PrettyPrintIntrinsicArgs);
W.printModule(this);
}
diff --git a/llvm/lib/IR/Intrinsics.cpp b/llvm/lib/IR/Intrinsics.cpp
index 6797a100ff732..167c95c57cb44 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"
@@ -1129,3 +1136,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..84e8585d1667b
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/tcgen05-mma-tensor-formatted.ll
@@ -0,0 +1,41 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; NOTE: This sample test demonstrates the --print-formatted-intrinsics feature for NVPTX intrinsics
+; RUN: llvm-as %s -o - | llvm-dis --print-formatted-intrinsics | 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::1 */ i32 1, /* collector::a::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::1 */ i32 1, /* collector::a::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::1 */ i32 1, /* collector::a::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::1 */ i32 1, /* collector::a::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::2 */ i32 2, /* collector::a::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::2 */ i32 2, /* collector::a::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::2 */ i32 2, /* collector::a::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::2 */ i32 2, /* collector::a::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
+}
+
+declare void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6), ptr addrspace(6), i64, i32, i1, i32, i32, i32)
\ No newline at end of file
diff --git a/llvm/tools/llvm-dis/llvm-dis.cpp b/llvm/tools/llvm-dis/llvm-dis.cpp
index 2b43d27f292a0..0642b5d1e5b0c 100644
--- a/llvm/tools/llvm-dis/llvm-dis.cpp
+++ b/llvm/tools/llvm-dis/llvm-dis.cpp
@@ -96,6 +96,11 @@ static cl::opt<bool> PrintThinLTOIndexOnly(
cl::desc("Only read thinlto index and print the index as LLVM assembly."),
cl::init(false), cl::Hidden, cl::cat(DisCategory));
+static cl::opt<bool> PrettyPrintIntrinsicArgs(
+ "print-formatted-intrinsics",
+ cl::desc("Enable pretty print format for intrinsic arguments"),
+ cl::init(false), cl::cat(DisCategory));
+
namespace {
static void printDebugLoc(const DebugLoc &DL, formatted_raw_ostream &OS) {
@@ -255,7 +260,8 @@ int main(int argc, char **argv) {
if (!DontPrint) {
if (M) {
M->removeDebugIntrinsicDeclarations();
- M->print(Out->os(), Annotator.get(), PreserveAssemblyUseListOrder);
+ M->print(Out->os(), Annotator.get(), PreserveAssemblyUseListOrder,
+ false, PrettyPrintIntrinsicArgs);
}
if (Index)
Index->print(Out->os());
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
index be7537c83da3a..2e3f5a61a29ac 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.cpp
@@ -441,6 +441,14 @@ 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("PrettyPrintImmArg")) {
+ unsigned ArgNo = R->getValueAsInt("ArgNo");
+ if (ArgNo < 1)
+ PrintFatalError(
+ R->getLoc(),
+ "PrettyPrintImmArg requires ArgNo >= 1 (0 is return value)");
+ StringRef FuncName = R->getValueAsString("FunctionName");
+ addPrettyPrintFunction(ArgNo - 1, FuncName);
} else {
llvm_unreachable("Unknown property!");
}
@@ -468,3 +476,17 @@ 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 FuncName) {
+ auto It = llvm::find_if(PrettyPrintFunctions, [ArgIdx](const auto &Pair) {
+ return Pair.first == ArgIdx;
+ });
+ if (It != PrettyPrintFunctions.end()) {
+ PrintFatalError(TheDef->getLoc(),
+ "PrettyPrintImmArg for argument " + Twine(ArgIdx) +
+ " is already defined as '" + It->second + "'");
+ return;
+ }
+ PrettyPrintFunctions.emplace_back(ArgIdx, FuncName);
+}
diff --git a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
index 2e86149514f46..1ef1843e952bd 100644
--- a/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/Basic/CodeGenIntrinsics.h
@@ -149,6 +149,11 @@ struct CodeGenIntrinsic {
void addArgAttribute(unsigned Idx, ArgAttrKind AK, uint64_t V = 0,
uint64_t V2 = 0);
+ /// Vector that stores ArgIndex, functionName.
+ SmallVector<std::pair<unsigned, StringRef>> PrettyPrintFunctions;
+
+ void addPrettyPrintFunction(unsigned ArgIdx, 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 75dffb18fca5a..bdfc6bf09fe68 100644
--- a/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/Basic/IntrinsicEmitter.cpp
@@ -60,8 +60,12 @@ class IntrinsicEmitter {
raw_ostream &OS);
void EmitIntrinsicToOverloadTable(const CodeGenIntrinsicTable &Ints,
raw_ostream &OS);
+ void EmitIntrinsicToPrettyPrintTable(const CodeGenIntrinsicTable &Ints,
+ raw_ostream &OS);
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 +113,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);
@@ -805,6 +815,66 @@ AttributeSet Intrinsic::getFnAttributes(LLVMContext &C, ID id) {{
NoFunctionAttrsID);
}
+void IntrinsicEmitter::EmitIntrinsicToPrettyPrintTable(
+ const CodeGenIntrinsicTable &Ints, raw_ostream &OS) {
+ OS << R"(// Intrinsic ID to pretty print bitset.
+)"
+ R"(#ifdef GET_INTRINSIC_PRETTY_PRINT_TABLE
+static constexpr uint8_t PPTable[] = {
+ 0
+ )";
+ int CountPerLine = 0;
+ for (auto [I, Int] : enumerate(Ints)) {
+ size_t Idx = I + 1;
+
+ if (Idx % 8 == 0) {
+ OS << ", 0";
+ CountPerLine++;
+ if (CountPerLine == 8) {
+ OS << "\n ";
+ CountPerLine = 0;
+ }
+ }
+ if (!Int.PrettyPrintFunctions.empty())
+ OS << " | (1<<" << Idx % 8 << ')';
+ }
+ OS << "\n};\n\n";
+ OS << "return (PPTable[id/8] & (1 << (id%8))) != 0;\n";
+ OS << "#endif // GET_INTRINSIC_PRETTY_PRINT_TABLE\n\n";
+}
+
+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 auto &Int : Ints) {
+ if (Int.PrettyPrintFunctions.empty())
+ continue;
+
+ OS << " case " << Int.EnumName << ":\n";
+ OS << " switch (ArgIdx) {\n";
+ for (const auto &[ArgIdx, FuncName] : Int.PrettyPrintFunctions) {
+ OS << " case " << ArgIdx << ": ";
+ if (!Int.TargetPrefix.empty())
+ OS << Int.TargetPrefix << "::";
+ OS << FuncName << "(OS, ImmArgVal); break;\n";
+ }
+ OS << " }\n";
+ OS << " break;\n";
+ }
+ OS << R"(
+ }
+}
+#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