[llvm] [NVPTX] Refactor protytype emission out of ISel (PR #203663)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 12 20:19:19 PDT 2026
https://github.com/AlexMaclean created https://github.com/llvm/llvm-project/pull/203663
Previously we horribly abused the TargetExternalSymbol operand to store an assembly string. This required hacks and pushed assembly emission into the SDAG ISel phase. Moving this will allow us to unify prototype and function declaration/definition printing.
>From c34d21d323599c54fc1110411bc3dc9dc5d31988 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Fri, 12 Jun 2026 20:14:38 -0700
Subject: [PATCH] [NVPTX] Refactor protytype emission out of ISel
---
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp | 10 --
.../NVPTX/MCTargetDesc/NVPTXInstPrinter.h | 2 -
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 108 ++++++++++++++--
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h | 2 +
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 122 ++----------------
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 4 -
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 12 --
.../Target/NVPTX/NVPTXMachineFunctionInfo.h | 15 +++
llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll | 2 +-
.../CodeGen/NVPTX/convert-call-to-indirect.ll | 12 +-
llvm/test/CodeGen/NVPTX/indirect_byval.ll | 6 +-
.../CodeGen/NVPTX/lower-args-gridconstant.ll | 12 +-
llvm/test/CodeGen/NVPTX/vaargs.ll | 6 +-
13 files changed, 140 insertions(+), 173 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
index 5af9e75bca399..beed605c56bd4 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp
@@ -447,16 +447,6 @@ void NVPTXInstPrinter::printHexu32imm(const MCInst *MI, int OpNum,
O << formatHex(Imm) << "U";
}
-void NVPTXInstPrinter::printProtoIdent(const MCInst *MI, int OpNum,
- const MCSubtargetInfo &,
- raw_ostream &O) {
- const MCOperand &Op = MI->getOperand(OpNum);
- assert(Op.isExpr() && "Call prototype is not an MCExpr?");
- const MCExpr *Expr = Op.getExpr();
- const MCSymbol &Sym = cast<MCSymbolRefExpr>(Expr)->getSymbol();
- O << Sym.getName();
-}
-
void NVPTXInstPrinter::printPrmtMode(const MCInst *MI, int OpNum,
const MCSubtargetInfo &, raw_ostream &O) {
const MCOperand &MO = MI->getOperand(OpNum);
diff --git a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
index 3f5dd99a2be05..4d9115a1e1b9d 100644
--- a/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
+++ b/llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h
@@ -54,8 +54,6 @@ class NVPTXInstPrinter : public MCInstPrinter {
const MCSubtargetInfo &STI, raw_ostream &O);
void printHexu32imm(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
raw_ostream &O);
- void printProtoIdent(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
- raw_ostream &O);
void printPrmtMode(const MCInst *MI, int OpNum, const MCSubtargetInfo &STI,
raw_ostream &O);
void printTmaReductionMode(const MCInst *MI, int OpNum,
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index cd5c27cebc182..7aec09d9ebed1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -32,6 +32,8 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/Sequence.h"
#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringExtras.h"
@@ -191,14 +193,6 @@ void NVPTXAsmPrinter::emitInstruction(const MachineInstr *MI) {
void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
OutMI.setOpcode(MI->getOpcode());
- // Special: Do not mangle symbol operand of CALL_PROTOTYPE
- if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
- const MachineOperand &MO = MI->getOperand(0);
- OutMI.addOperand(GetSymbolRef(
- OutContext.getOrCreateSymbol(Twine(MO.getSymbolName()))));
- return;
- }
-
for (const auto MO : MI->operands())
OutMI.addOperand(lowerOperand(MO));
}
@@ -320,6 +314,99 @@ void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
printReturnValStr(&F, O);
}
+void NVPTXAsmPrinter::emitCallPrototype(const CallBase &CB,
+ unsigned UniqueCallSite,
+ raw_ostream &O) const {
+ const DataLayout &DL = getDataLayout();
+ const NVPTXSubtarget &STI = MF->getSubtarget<NVPTXSubtarget>();
+ const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
+ const auto PtrVT = TLI->getPointerTy(DL);
+ Type *RetTy = CB.getFunctionType()->getReturnType();
+
+ O << "prototype_" << UniqueCallSite << " : .callprototype ";
+
+ if (RetTy->isVoidTy()) {
+ O << "()";
+ } else {
+ O << "(";
+ if (shouldPassAsArray(RetTy)) {
+ const Align RetAlign =
+ getPTXParamAlign(&CB, RetTy, AttributeList::ReturnIndex, DL);
+ O << ".param .align " << RetAlign.value() << " .b8 _["
+ << DL.getTypeAllocSize(RetTy) << "]";
+ } else if (RetTy->isFloatingPointTy() || RetTy->isIntegerTy()) {
+ unsigned size = 0;
+ if (auto *ITy = dyn_cast<IntegerType>(RetTy)) {
+ size = ITy->getBitWidth();
+ } else {
+ assert(RetTy->isFloatingPointTy() &&
+ "Floating point type expected here");
+ size = RetTy->getPrimitiveSizeInBits();
+ }
+ // PTX ABI requires all scalar return values to be at least 32
+ // bits in size. fp16 normally uses .b16 as its storage type in
+ // PTX, so its size must be adjusted here, too.
+ size = promoteScalarArgumentSize(size);
+
+ O << ".param .b" << size << " _";
+ } else if (isa<PointerType>(RetTy)) {
+ O << ".param .b" << PtrVT.getSizeInBits() << " _";
+ } else {
+ llvm_unreachable("Unknown return type");
+ }
+ O << ") ";
+ }
+ O << "_ (";
+
+ auto MakeArg = [&](const unsigned I) {
+ Type *Ty = CB.getArgOperand(I)->getType();
+
+ if (CB.paramHasAttr(I, Attribute::ByVal)) {
+ // Indirect calls need strict ABI alignment so we disable optimizations by
+ // not providing a function to optimize.
+ Type *ETy = CB.getParamByValType(I);
+ // Mirror the byval alignment computed by SelectionDAGBuilder: prefer an
+ // explicit stack/param alignment, otherwise fall back to the byval type
+ // alignment.
+ MaybeAlign InitialAlign = CB.getParamStackAlign(I);
+ if (!InitialAlign)
+ InitialAlign = CB.getParamAlign(I);
+ Align ByValAlign =
+ InitialAlign.value_or(TLI->getByValTypeAlignment(ETy, DL));
+ Align ParamByValAlign =
+ getDeviceByValParamAlign(/*F=*/nullptr, ETy, ByValAlign, DL);
+
+ O << ".param .align " << ParamByValAlign.value() << " .b8 _["
+ << DL.getTypeAllocSize(ETy) << "]";
+ return;
+ }
+
+ if (shouldPassAsArray(Ty)) {
+ Align ParamAlign =
+ getPTXParamAlign(&CB, Ty, I + AttributeList::FirstArgIndex, DL);
+ O << ".param .align " << ParamAlign.value() << " .b8 _["
+ << DL.getTypeAllocSize(Ty) << "]";
+ return;
+ }
+ // scalar type
+ unsigned sz = 0;
+ if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
+ sz = promoteScalarArgumentSize(ITy->getBitWidth());
+ } else if (isa<PointerType>(Ty)) {
+ sz = PtrVT.getSizeInBits();
+ } else {
+ sz = Ty->getPrimitiveSizeInBits();
+ }
+ O << ".param .b" << sz << " _";
+ };
+ interleave(seq(CB.arg_size()), O, MakeArg, ", ");
+
+ O << ")";
+ if (shouldEmitPTXNoReturn(&CB, TM))
+ O << " .noreturn";
+ O << ";\n";
+}
+
// Return true if MBB is the header of a loop marked with
// llvm.loop.unroll.disable or llvm.loop.unroll.count=1.
bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
@@ -417,6 +504,11 @@ void NVPTXAsmPrinter::emitFunctionBodyStart() {
SmallString<128> Str;
raw_svector_ostream O(Str);
emitDemotedVars(&MF->getFunction(), O);
+
+ const auto *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
+ for (const auto &[Id, CB] : MFI->getCallPrototypes())
+ emitCallPrototype(*CB, Id, O);
+
OutStreamer->emitRawText(O.str());
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h
index bc5f20860083a..8a9b811fec200 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h
@@ -188,6 +188,8 @@ class LLVM_LIBRARY_VISIBILITY NVPTXAsmPrinter : public AsmPrinter {
void encodeDebugInfoRegisterNumbers(const MachineFunction &MF);
void printReturnValStr(const Function *, raw_ostream &O);
void printReturnValStr(const MachineFunction &MF, raw_ostream &O);
+ void emitCallPrototype(const CallBase &CB, unsigned UniqueCallSite,
+ raw_ostream &O) const;
bool PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
const char *ExtraCode, raw_ostream &) override;
void printOperand(const MachineInstr *MI, unsigned OpNum, raw_ostream &O);
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7555845847935..f49da9c104662 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -15,6 +15,7 @@
#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "NVPTX.h"
#include "NVPTXISelDAGToDAG.h"
+#include "NVPTXMachineFunctionInfo.h"
#include "NVPTXSelectionDAGInfo.h"
#include "NVPTXSubtarget.h"
#include "NVPTXTargetMachine.h"
@@ -1189,102 +1190,6 @@ SDValue NVPTXTargetLowering::getSqrtEstimate(SDValue Operand, SelectionDAG &DAG,
}
}
-std::string NVPTXTargetLowering::getPrototype(
- const DataLayout &DL, Type *RetTy, const ArgListTy &Args,
- const SmallVectorImpl<ISD::OutputArg> &Outs, const CallBase &CB,
- unsigned UniqueCallSite) const {
- auto PtrVT = getPointerTy(DL);
-
- std::string Prototype;
- raw_string_ostream O(Prototype);
- O << "prototype_" << UniqueCallSite << " : .callprototype ";
-
- if (RetTy->isVoidTy()) {
- O << "()";
- } else {
- O << "(";
- if (shouldPassAsArray(RetTy)) {
- const Align RetAlign =
- getPTXParamAlign(&CB, RetTy, AttributeList::ReturnIndex, DL);
- O << ".param .align " << RetAlign.value() << " .b8 _["
- << DL.getTypeAllocSize(RetTy) << "]";
- } else if (RetTy->isFloatingPointTy() || RetTy->isIntegerTy()) {
- unsigned size = 0;
- if (auto *ITy = dyn_cast<IntegerType>(RetTy)) {
- size = ITy->getBitWidth();
- } else {
- assert(RetTy->isFloatingPointTy() &&
- "Floating point type expected here");
- size = RetTy->getPrimitiveSizeInBits();
- }
- // PTX ABI requires all scalar return values to be at least 32
- // bits in size. fp16 normally uses .b16 as its storage type in
- // PTX, so its size must be adjusted here, too.
- size = promoteScalarArgumentSize(size);
-
- O << ".param .b" << size << " _";
- } else if (isa<PointerType>(RetTy)) {
- O << ".param .b" << PtrVT.getSizeInBits() << " _";
- } else {
- llvm_unreachable("Unknown return type");
- }
- O << ") ";
- }
- O << "_ (";
-
- auto AllOuts = ArrayRef(Outs);
- auto MakeArg = [&](const unsigned I) {
- const auto ArgOuts =
- AllOuts.take_while([I](auto O) { return O.OrigArgIndex == I; });
- AllOuts = AllOuts.drop_front(ArgOuts.size());
-
- Type *Ty = Args[I].Ty;
-
- if (ArgOuts[0].Flags.isByVal()) {
- // Indirect calls need strict ABI alignment so we disable optimizations by
- // not providing a function to optimize.
- Type *ETy = Args[I].IndirectType;
- Align InitialAlign = ArgOuts[0].Flags.getNonZeroByValAlign();
- Align ParamByValAlign =
- getDeviceByValParamAlign(/*F=*/nullptr, ETy, InitialAlign, DL);
-
- O << ".param .align " << ParamByValAlign.value() << " .b8 _["
- << ArgOuts[0].Flags.getByValSize() << "]";
- return;
- }
-
- if (shouldPassAsArray(Ty)) {
- Align ParamAlign =
- getPTXParamAlign(&CB, Ty, I + AttributeList::FirstArgIndex, DL);
- O << ".param .align " << ParamAlign.value() << " .b8 _["
- << DL.getTypeAllocSize(Ty) << "]";
- return;
- }
- // i8 types in IR will be i16 types in SDAG
- assert((getValueType(DL, Ty) == ArgOuts[0].VT ||
- (getValueType(DL, Ty) == MVT::i8 && ArgOuts[0].VT == MVT::i16)) &&
- "type mismatch between callee prototype and arguments");
- // scalar type
- unsigned sz = 0;
- if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
- sz = promoteScalarArgumentSize(ITy->getBitWidth());
- } else if (isa<PointerType>(Ty)) {
- sz = PtrVT.getSizeInBits();
- } else {
- sz = Ty->getPrimitiveSizeInBits();
- }
- O << ".param .b" << sz << " _";
- };
- interleave(seq(Args.size()), O, MakeArg, ", ");
-
- O << ")";
- if (shouldEmitPTXNoReturn(&CB, *nvTM))
- O << " .noreturn";
- O << ";";
-
- return Prototype;
-}
-
static MachinePointerInfo refinePtrAS(SDValue &Ptr, SelectionDAG &DAG,
const DataLayout &DL,
const TargetLowering &TL) {
@@ -1566,22 +1471,15 @@ SDValue NVPTXTargetLowering::LowerCall(TargetLowering::CallLoweringInfo &CLI,
CalleeFunc->addFnAttr("nvptx-libcall-callee", "true");
}
- if (IsIndirectCall) {
- // This is indirect function call case : PTX requires a prototype of the
- // form
- // proto_0 : .callprototype(.param .b32 _) _ (.param .b32 _);
- // to be emitted, and the label has to used as the last arg of call
- // instruction.
- // The prototype is embedded in a string and put as the operand for a
- // CallPrototype SDNode which will print out to the value of the string.
- std::string Proto =
- getPrototype(DL, RetTy, Args, CLI.Outs, *CB, UniqueCallSite);
- const char *ProtoStr = nvTM->getStrPool().save(Proto).data();
- const SDValue PrototypeDeclare = DAG.getNode(
- NVPTXISD::CallPrototype, dl, MVT::Other,
- {StartChain, DAG.getTargetExternalSymbol(ProtoStr, MVT::i32)});
- CallPrereqs.push_back(PrototypeDeclare);
- }
+ // In the indirect function call case, PTX requires a prototype of the form:
+ // proto_0 : .callprototype(.param .b32 _) _ (.param .b32 _);
+ // Where the label is to be used as the last arg of the call instruction.
+ // We record the call site here and emit all prototypes at the
+ // start of the function in the AsmPrinter.
+ if (IsIndirectCall)
+ DAG.getMachineFunction()
+ .getInfo<NVPTXMachineFunctionInfo>()
+ ->addCallPrototype(UniqueCallSite, CB);
const bool IsUnknownIntrinsic =
CalleeF && CalleeF->isIntrinsic() &&
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 76892c229a842..82a356d5fbd33 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -82,10 +82,6 @@ class NVPTXTargetLowering : public TargetLowering {
SDValue LowerSTACKSAVE(SDValue Op, SelectionDAG &DAG) const;
SDValue LowerSTACKRESTORE(SDValue Op, SelectionDAG &DAG) const;
- std::string getPrototype(const DataLayout &DL, Type *, const ArgListTy &,
- const SmallVectorImpl<ISD::OutputArg> &,
- const CallBase &CB, unsigned UniqueCallSite) const;
-
SDValue LowerReturn(SDValue Chain, CallingConv::ID CallConv, bool isVarArg,
const SmallVectorImpl<ISD::OutputArg> &Outs,
const SmallVectorImpl<SDValue> &OutVals, const SDLoc &dl,
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index ff3e5144c6fbe..79f012912cb74 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1876,18 +1876,6 @@ def : Pat<(declare_array_param externalsym:$a, imm:$align, imm:$size),
def : Pat<(declare_scalar_param externalsym:$a, imm:$size),
(DECLARE_PARAM_scalar (to_texternsym $a), imm:$size)>;
-// Call prototype wrapper, this is a dummy instruction that just prints it's
-// operand which is string defining the prototype.
-def SDTCallPrototype : SDTypeProfile<0, 1, [SDTCisInt<0>]>;
-def CallPrototype :
- SDNode<"NVPTXISD::CallPrototype", SDTCallPrototype,
- [SDNPHasChain, SDNPSideEffect]>;
-def ProtoIdent : Operand<i32> { let PrintMethod = "printProtoIdent"; }
-def CALL_PROTOTYPE :
- NVPTXInst<(outs), (ins ProtoIdent:$ident),
- "$ident", [(CallPrototype (i32 texternalsym:$ident))]>;
-
-
foreach t = [I32RT, I64RT] in {
defvar inst_name = "MOV" # t.Size # "_PARAM";
def inst_name : BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src), "mov.b" # t.Size>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
index 8feae341893aa..866dba0942c01 100644
--- a/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
+++ b/llvm/lib/Target/NVPTX/NVPTXMachineFunctionInfo.h
@@ -16,14 +16,21 @@
#include "llvm/ADT/StringRef.h"
#include "llvm/CodeGen/MachineFunction.h"
+#include <map>
namespace llvm {
+class CallBase;
+
class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
private:
/// Stores a mapping from index to symbol name for image handles that are
/// replaced with image references
SmallVector<std::string, 8> ImageHandleList;
+ /// Stores a mapping from a unique call-site id to the call instruction that
+ /// needs an indirect-call prototype emitted.
+ std::map<unsigned, const CallBase *> CallPrototypes;
+
public:
NVPTXMachineFunctionInfo(const Function &F, const TargetSubtargetInfo *STI) {}
@@ -52,6 +59,14 @@ class NVPTXMachineFunctionInfo : public MachineFunctionInfo {
bool checkImageHandleSymbol(StringRef Symbol) const {
return llvm::is_contained(ImageHandleList, Symbol);
}
+
+ void addCallPrototype(unsigned Id, const CallBase *CB) {
+ CallPrototypes.try_emplace(Id, CB);
+ }
+
+ const std::map<unsigned, const CallBase *> &getCallPrototypes() const {
+ return CallPrototypes;
+ }
};
}
diff --git a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
index 483d48a1012c6..2a36c4162ba02 100644
--- a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
@@ -13,11 +13,11 @@ target triple = "nvptx64-nvidia-cuda"
%"class.sycl::_V1::detail::half_impl::half" = type { half }
%complex_half = type { half, half }
+; CHECK: prototype_0 : .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);
; CHECK: .param .align 2 .b8 param2[4];
; CHECK: .param .align 2 .b8 retval0[4];
; CHECK-DAG: st.param.b16 [param2], %rs{{[0-9]+}};
; CHECK-DAG: st.param.b16 [param2+2], %rs{{[0-9]+}};
-; CHECK: prototype_0 : .callprototype (.param .align 2 .b8 _[4]) _ (.param .b32 _, .param .b32 _, .param .align 2 .b8 _[4]);
; CHECK: call (retval0),
define weak_odr void @foo() {
entry:
diff --git a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
index 4c60d1b253690..1510be1b0daa9 100644
--- a/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-call-to-indirect.ll
@@ -10,14 +10,13 @@ define %struct.64 @test_return_type_mismatch(ptr %p) {
; CHECK-LABEL: test_return_type_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<32>;
-; CHECK-EMPTY:
+; CHECK-NEXT: prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_param_0];
; CHECK-NEXT: { // callseq 0, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .align 1 .b8 retval0[8];
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; CHECK-NEXT: prototype_0 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _);
; CHECK-NEXT: mov.b64 %rd2, callee;
; CHECK-NEXT: call (retval0), %rd2, (param0), prototype_0;
; CHECK-NEXT: ld.param.b8 %rd3, [retval0+7];
@@ -67,12 +66,11 @@ define i64 @test_param_type_mismatch(ptr %p) {
; CHECK-LABEL: test_param_type_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
+; CHECK-NEXT: prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: { // callseq 1, 0
; CHECK-NEXT: .param .b64 param0;
; CHECK-NEXT: .param .b64 retval0;
-; CHECK-NEXT: prototype_1 : .callprototype (.param .b64 _) _ (.param .b64 _);
; CHECK-NEXT: st.param.b64 [param0], 7;
; CHECK-NEXT: mov.b64 %rd1, callee;
; CHECK-NEXT: call (retval0), %rd1, (param0), prototype_1;
@@ -88,7 +86,7 @@ define i64 @test_param_count_mismatch(ptr %p) {
; CHECK-LABEL: test_param_count_mismatch(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-EMPTY:
+; CHECK-NEXT: prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_param_count_mismatch_param_0];
; CHECK-NEXT: { // callseq 2, 0
@@ -96,7 +94,6 @@ define i64 @test_param_count_mismatch(ptr %p) {
; CHECK-NEXT: .param .b64 param1;
; CHECK-NEXT: .param .b64 retval0;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; CHECK-NEXT: prototype_2 : .callprototype (.param .b64 _) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: st.param.b64 [param1], 7;
; CHECK-NEXT: mov.b64 %rd2, callee;
; CHECK-NEXT: call (retval0), %rd2, (param0, param1), prototype_2;
@@ -115,7 +112,7 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: .reg .b64 %SP;
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b64 %rd<33>;
-; CHECK-EMPTY:
+; CHECK-NEXT: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: mov.b64 %SPL, __local_depot3;
; CHECK-NEXT: ld.param.b64 %rd1, [test_return_type_mismatch_variadic_param_0];
@@ -126,7 +123,6 @@ define %struct.64 @test_return_type_mismatch_variadic(ptr %p) {
; CHECK-NEXT: .param .align 1 .b8 retval0[8];
; CHECK-NEXT: st.param.b64 [param1], %rd2;
; CHECK-NEXT: st.param.b64 [param0], %rd1;
-; CHECK-NEXT: prototype_3 : .callprototype (.param .align 1 .b8 _[8]) _ (.param .b64 _, .param .b64 _);
; CHECK-NEXT: mov.b64 %rd3, callee_variadic;
; CHECK-NEXT: call (retval0), %rd3, (param0, param1), prototype_3;
; CHECK-NEXT: ld.param.b8 %rd4, [retval0+7];
diff --git a/llvm/test/CodeGen/NVPTX/indirect_byval.ll b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
index e1fecdb76bd4d..c982aaf70d4a2 100644
--- a/llvm/test/CodeGen/NVPTX/indirect_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/indirect_byval.ll
@@ -18,7 +18,7 @@ define internal i32 @foo() {
; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-EMPTY:
+; CHECK-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot0;
; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
@@ -32,7 +32,6 @@ define internal i32 @foo() {
; CHECK-NEXT: add.u64 %rd3, %SPL, 1;
; CHECK-NEXT: ld.local.b8 %rs1, [%rd3];
; CHECK-NEXT: st.param.b8 [param0], %rs1;
-; CHECK-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_0;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
; CHECK-NEXT: } // callseq 0
@@ -55,7 +54,7 @@ define internal i32 @bar() {
; CHECK-NEXT: .reg .b64 %SPL;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-NEXT: .reg .b64 %rd<5>;
-; CHECK-EMPTY:
+; CHECK-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK-NEXT: cvta.local.u64 %SP, %SPL;
@@ -69,7 +68,6 @@ define internal i32 @bar() {
; CHECK-NEXT: add.u64 %rd3, %SPL, 8;
; CHECK-NEXT: ld.local.b64 %rd4, [%rd3];
; CHECK-NEXT: st.param.b64 [param0], %rd4;
-; CHECK-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .align 8 .b8 _[8], .param .b64 _);
; CHECK-NEXT: call (retval0), %rd1, (param0, param1), prototype_1;
; CHECK-NEXT: ld.param.b32 %r1, [retval0];
; CHECK-NEXT: } // callseq 1
diff --git a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
index 4d9cb6ff97413..d2417337c68f8 100644
--- a/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
+++ b/llvm/test/CodeGen/NVPTX/lower-args-gridconstant.ll
@@ -122,7 +122,7 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-LABEL: grid_const_escape(
; PTX: {
; PTX-NEXT: .reg .b64 %rd<4>;
-; PTX-EMPTY:
+; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %rd1, grid_const_escape_param_0;
; PTX-NEXT: cvta.param.u64 %rd2, %rd1;
@@ -130,7 +130,6 @@ define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.gri
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param0], %rd2;
-; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-NEXT: mov.b64 %rd3, escape;
; PTX-NEXT: call (retval0), %rd3, (param0), prototype_0;
; PTX-NEXT: } // callseq 0
@@ -153,7 +152,7 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-NEXT: .reg .b64 %SPL;
; PTX-NEXT: .reg .b32 %r<2>;
; PTX-NEXT: .reg .b64 %rd<8>;
-; PTX-EMPTY:
+; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %SPL, __local_depot4;
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
@@ -173,7 +172,6 @@ define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4
; PTX-NEXT: st.param.b64 [param2], %rd3;
; PTX-NEXT: st.param.b64 [param1], %rd5;
; PTX-NEXT: st.param.b64 [param0], %rd4;
-; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
; PTX-NEXT: mov.b64 %rd7, escape3;
; PTX-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;
; PTX-NEXT: } // callseq 1
@@ -254,7 +252,7 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX: {
; PTX-NEXT: .reg .b32 %r<3>;
; PTX-NEXT: .reg .b64 %rd<6>;
-; PTX-EMPTY:
+; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0;
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1];
@@ -267,7 +265,6 @@ define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) align 4 "nvvm.g
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param0], %rd4;
-; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-NEXT: mov.b64 %rd5, escape;
; PTX-NEXT: call (retval0), %rd5, (param0), prototype_2;
; PTX-NEXT: } // callseq 2
@@ -293,7 +290,7 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX: {
; PTX-NEXT: .reg .b32 %r<4>;
; PTX-NEXT: .reg .b64 %rd<6>;
-; PTX-EMPTY:
+; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-NEXT: // %bb.0:
; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0;
; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1];
@@ -307,7 +304,6 @@ define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) align 4
; PTX-NEXT: .param .b64 param0;
; PTX-NEXT: .param .b32 retval0;
; PTX-NEXT: st.param.b64 [param0], %rd4;
-; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
; PTX-NEXT: mov.b64 %rd5, escape;
; PTX-NEXT: call (retval0), %rd5, (param0), prototype_3;
; PTX-NEXT: } // callseq 3
diff --git a/llvm/test/CodeGen/NVPTX/vaargs.ll b/llvm/test/CodeGen/NVPTX/vaargs.ll
index 0a9c59f9ef040..ea0807a88bd68 100644
--- a/llvm/test/CodeGen/NVPTX/vaargs.ll
+++ b/llvm/test/CodeGen/NVPTX/vaargs.ll
@@ -188,7 +188,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .reg .b32 %SPL;
; CHECK32-NEXT: .reg .b32 %r<9>;
; CHECK32-NEXT: .reg .b64 %rd<3>;
-; CHECK32-EMPTY:
+; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT: // %bb.0: // %entry
; CHECK32-NEXT: mov.b32 %SPL, __local_depot1;
; CHECK32-NEXT: cvta.local.u32 %SP, %SPL;
@@ -210,7 +210,6 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK32-NEXT: .param .b32 param1;
; CHECK32-NEXT: .param .b32 retval0;
; CHECK32-NEXT: st.param.b32 [param1], %r7;
-; CHECK32-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b32 _);
; CHECK32-NEXT: st.param.b32 [param0], 4;
; CHECK32-NEXT: call (retval0), %r5, (param0, param1), prototype_1;
; CHECK32-NEXT: ld.param.b32 %r8, [retval0];
@@ -225,7 +224,7 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .reg .b64 %SPL;
; CHECK64-NEXT: .reg .b32 %r<3>;
; CHECK64-NEXT: .reg .b64 %rd<9>;
-; CHECK64-EMPTY:
+; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT: // %bb.0: // %entry
; CHECK64-NEXT: mov.b64 %SPL, __local_depot1;
; CHECK64-NEXT: cvta.local.u64 %SP, %SPL;
@@ -247,7 +246,6 @@ define i32 @test_foo(i32 %i, i64 %l, double %d, ptr %p) {
; CHECK64-NEXT: .param .b64 param1;
; CHECK64-NEXT: .param .b32 retval0;
; CHECK64-NEXT: st.param.b64 [param1], %rd8;
-; CHECK64-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b32 _, .param .b64 _);
; CHECK64-NEXT: st.param.b32 [param0], 4;
; CHECK64-NEXT: call (retval0), %rd6, (param0, param1), prototype_1;
; CHECK64-NEXT: ld.param.b32 %r2, [retval0];
More information about the llvm-commits
mailing list