Skip to content

Commit 39e7712

Browse files
DharuniRAcharyajurahul
DharuniRAcharya
and
jurahul
authored
[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 <[email protected]> Co-authored-by: Rahul Joshi <[email protected]>
1 parent c66f1fd commit 39e7712

File tree

11 files changed

+395
-27
lines changed

11 files changed

+395
-27
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: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -142,6 +142,25 @@ class Range<AttrIndex idx, int lower, int upper> : IntrinsicProperty {
142142
int Upper = upper;
143143
}
144144

145+
// ArgProperty - Base class for argument properties that can be specified in ArgInfo.
146+
class ArgProperty;
147+
148+
// ArgName - Specifies the name of an argument for pretty-printing.
149+
class ArgName<string name> : ArgProperty {
150+
string Name = name;
151+
}
152+
153+
// ImmArgPrinter - Specifies a custom printer function for immediate arguments.
154+
class ImmArgPrinter<string funcname> : ArgProperty {
155+
string FuncName = funcname;
156+
}
157+
158+
// ArgInfo - The specified argument has properties defined by a list of ArgProperty objects.
159+
class ArgInfo<ArgIndex idx, list<ArgProperty> arg_properties> : IntrinsicProperty {
160+
int ArgNo = idx.Value;
161+
list<ArgProperty> Properties = arg_properties;
162+
}
163+
145164
def IntrNoReturn : IntrinsicProperty;
146165

147166
// Applied by default.

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

Lines changed: 15 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2955,7 +2955,14 @@ foreach sp = [0, 1] in {
29552955
defvar nargs = !size(args);
29562956
defvar scale_d_imm = ArgIndex<!sub(nargs, 1)>;
29572957
defvar scale_d_imm_range = [ImmArg<scale_d_imm>, Range<scale_d_imm, 0, 16>];
2958-
defvar intrinsic_properties = !listconcat(
2958+
2959+
// Check if this is the specific llvm.nvvm.tcgen05.mma.tensor intrinsic.
2960+
defvar is_target_intrinsic = !and(!eq(sp, 0),
2961+
!eq(space, "tensor"),
2962+
!eq(scale_d, 0),
2963+
!eq(ashift, 0));
2964+
2965+
defvar base_properties = !listconcat(
29592966
mma.common_intr_props,
29602967
!if(!eq(scale_d, 1), scale_d_imm_range, []),
29612968
[Range<ArgIndex<nargs>, 0, !if(!eq(scale_d, 1), 2, 4)>, // kind
@@ -2965,6 +2972,13 @@ foreach sp = [0, 1] in {
29652972
]
29662973
);
29672974

2975+
defvar intrinsic_properties = !if(is_target_intrinsic,
2976+
!listconcat(base_properties,
2977+
[ArgInfo<ArgIndex<nargs>, [ArgName<"kind">, ImmArgPrinter<"printTcgen05MMAKind">]>,
2978+
ArgInfo<ArgIndex<!add(nargs, 1)>, [ArgName<"cta_group">]>,
2979+
ArgInfo<ArgIndex<!add(nargs, 2)>, [ArgName<"collector">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]),
2980+
base_properties);
2981+
29682982
def mma.record_name:
29692983
DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
29702984
mma.intr_name>;

‎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 auto *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 auto *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"
@@ -4576,12 +4577,38 @@ void AssemblyWriter::printInstruction(const Instruction &I) {
45764577
Out << ' ';
45774578
writeOperand(Operand, false);
45784579
Out << '(';
4580+
bool HasPrettyPrintedArgs =
4581+
isa<IntrinsicInst>(CI) &&
4582+
Intrinsic::hasPrettyPrintedArgs(CI->getIntrinsicID());
4583+
45794584
ListSeparator LS;
4580-
for (unsigned op = 0, Eop = CI->arg_size(); op < Eop; ++op) {
4581-
Out << LS;
4582-
writeParamOperand(CI->getArgOperand(op), PAL.getParamAttrs(op));
4585+
Function *CalledFunc = CI->getCalledFunction();
4586+
auto PrintArgComment = [&](unsigned ArgNo) {
4587+
const auto *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+
};
4598+
if (HasPrettyPrintedArgs) {
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 ArgNo = 0, NumArgs = CI->arg_size(); ArgNo < NumArgs;
4607+
++ArgNo) {
4608+
Out << LS;
4609+
writeParamOperand(CI->getArgOperand(ArgNo), PAL.getParamAttrs(ArgNo));
4610+
}
45834611
}
4584-
45854612
// Emit an ellipsis if this is a musttail call in a vararg function. This
45864613
// is only to aid readability, musttail calls forward varargs by default.
45874614
if (CI->isMustTailCall() && CI->getParent() &&
@@ -5005,12 +5032,10 @@ void AssemblyWriter::printUseLists(const Function *F) {
50055032
//===----------------------------------------------------------------------===//
50065033

50075034
void Function::print(raw_ostream &ROS, AssemblyAnnotationWriter *AAW,
5008-
bool ShouldPreserveUseListOrder,
5009-
bool IsForDebug) const {
5035+
bool ShouldPreserveUseListOrder, bool IsForDebug) const {
50105036
SlotTracker SlotTable(this->getParent());
50115037
formatted_raw_ostream OS(ROS);
5012-
AssemblyWriter W(OS, SlotTable, this->getParent(), AAW,
5013-
IsForDebug,
5038+
AssemblyWriter W(OS, SlotTable, this->getParent(), AAW, IsForDebug,
50145039
ShouldPreserveUseListOrder);
50155040
W.printFunction(this);
50165041
}

‎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"
@@ -1142,3 +1149,7 @@ Intrinsic::ID Intrinsic::getDeinterleaveIntrinsicID(unsigned Factor) {
11421149
assert(Factor >= 2 && Factor <= 8 && "Unexpected factor");
11431150
return InterleaveIntrinsics[Factor - 2].Deinterleave;
11441151
}
1152+
1153+
#define GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS
1154+
#include "llvm/IR/IntrinsicImpl.inc"
1155+
#undef GET_INTRINSIC_PRETTY_PRINT_ARGUMENTS

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

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,50 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
2+
; NOTE: This sample test demonstrates the pretty print feature for NVPTX intrinsics
3+
; RUN: llvm-as < %s | llvm-dis | 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+
; This test verifies that printImmArg is safe to call on all constant arguments, but only prints comments for arguments that have pretty printing configured.
42+
define void @test_mixed_constants_edge_case(ptr addrspace(6) %dtmem, ptr addrspace(6) %atensor) {
43+
; CHECK-LABEL: define void @test_mixed_constants_edge_case(
44+
; 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)
45+
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)
46+
47+
ret void
48+
}
49+
50+
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>, [ArgName<"mode">, ImmArgPrinter<"printDummyMode">]>,
18+
ArgInfo<ArgIndex<2>, [ArgName<"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>, [ArgName<"rounding_mode">, ImmArgPrinter<"printRoundingMode">]>,
31+
ArgInfo<ArgIndex<3>, [ArgName<"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: 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

0 commit comments

Comments
 (0)