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 : 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 : ArgProperty { + string Name = name; +} + +// ImmArgPrinter - Specifies a custom printer function for immediate arguments. +class ImmArgPrinter : ArgProperty { + string FuncName = funcname; +} + +// ArgInfo - The specified argument has properties defined by a list of ArgProperty objects. +class ArgInfo arg_properties> : IntrinsicProperty { + int ArgNo = idx.Value; + list 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; defvar scale_d_imm_range = [ImmArg, Range]; - 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, 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, [ArgName<"kind">, ImmArgPrinter<"printTcgen05MMAKind">]>, + ArgInfo, [ArgName<"cta_group">]>, + ArgInfo, [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 #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(ImmArgVal)) { + uint64_t Val = CI->getZExtValue(); + switch (static_cast(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(ImmArgVal)) { + uint64_t Val = CI->getZExtValue(); + switch (static_cast(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(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(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>, + ArgInfo, [ArgName<"mode">, ImmArgPrinter<"printDummyMode">]>, + ArgInfo, [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>, + ImmArg>, + ArgInfo, [ArgName<"rounding_mode">, ImmArgPrinter<"printRoundingMode">]>, + ArgInfo, [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(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 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 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 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; @@ -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";