[llvm] [NVPTX] Consolidate AsmPrinter type and global variable printing (PR #222240)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 28 12:52:36 PDT 2026
https://github.com/AlexMaclean updated https://github.com/llvm/llvm-project/pull/222240
>From 6dc857d3e4f595bd2913f94325cce38bd12c300a Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Tue, 8 Sep 2026 22:41:21 -0700
Subject: [PATCH 1/2] [NVPTX] Consolidate AsmPrinter type and global variable
printing
---
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp | 681 ++++++------------
llvm/lib/Target/NVPTX/NVPTXUtilities.h | 16 +-
llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll | 2 +-
llvm/test/CodeGen/NVPTX/callchain.ll | 2 +-
llvm/test/CodeGen/NVPTX/empty-type.ll | 2 +-
llvm/test/CodeGen/NVPTX/fp-global.ll | 55 ++
llvm/test/CodeGen/NVPTX/i1-param.ll | 2 +-
llvm/test/CodeGen/NVPTX/managed.ll | 2 +-
llvm/test/CodeGen/NVPTX/noreturn.ll | 2 +-
llvm/test/CodeGen/NVPTX/odd-integer-width.ll | 37 +
llvm/test/CodeGen/NVPTX/param-align.ll | 10 +-
llvm/test/CodeGen/NVPTX/vector-global.ll | 7 +
llvm/test/DebugInfo/NVPTX/debug-info.ll | 4 +-
13 files changed, 340 insertions(+), 482 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/fp-global.ll
create mode 100644 llvm/test/CodeGen/NVPTX/odd-integer-width.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index e2b55b3be4dc0..9444c61ad1e8a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -251,8 +251,6 @@ class NVPTXAsmPrinter : public AsmPrinter {
void emitFunctionParamList(const Function *, raw_ostream &O);
void setAndEmitFunctionVirtualRegisters(const MachineFunction &MF);
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;
void emitJumpTable(const MachineJumpTableEntry &MJT, unsigned MJTI) const;
@@ -313,14 +311,16 @@ class NVPTXAsmPrinter : public AsmPrinter {
// List of variables demoted to a function scope.
std::map<const Function *, std::vector<const GlobalVariable *>> localDecls;
- void emitPTXGlobalVariable(const GlobalVariable *GVar, raw_ostream &O,
- const NVPTXSubtarget &STI);
+ /// Print the state space, alignment, type, name, and — when
+ /// \p EmitInitializer is set — the initializer of \p GVar. Passing false
+ /// prints a declaration whose type still matches the definition, as an
+ /// `.extern` forward declaration requires.
void emitPTXGlobalVariableDefinition(const GlobalVariable *GVar,
raw_ostream &O,
const NVPTXSubtarget &STI,
bool EmitInitializer);
void emitPTXAddressSpace(unsigned int AddressSpace, raw_ostream &O) const;
- std::string getPTXFundamentalTypeStr(Type *Ty, bool = true) const;
+ std::string getPTXFundamentalTypeStr(Type *Ty) const;
void printScalarConstant(const Constant *CPV, raw_ostream &O);
void printFPConstant(const ConstantFP *Fp, raw_ostream &O) const;
void bufferLEByte(const Constant *CPV, int Bytes, AggBuffer *aggBuffer);
@@ -374,21 +374,6 @@ class NVPTXAsmPrinter : public AsmPrinter {
} // end anonymous namespace
-static StringRef getTextureName(const Value &V) {
- assert(V.hasName() && "Found texture variable with no name");
- return V.getName();
-}
-
-static StringRef getSurfaceName(const Value &V) {
- assert(V.hasName() && "Found surface variable with no name");
- return V.getName();
-}
-
-static StringRef getSamplerName(const Value &V) {
- assert(V.hasName() && "Found sampler variable with no name");
- return V.getName();
-}
-
/// Emits initial debug location directive.
static void emitInitialRawDwarfLocDirective(const MachineFunction &MF,
DwarfDebug *DD,
@@ -690,118 +675,64 @@ MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
return MCOperand::createExpr(Expr);
}
-void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
- const DataLayout &DL = getDataLayout();
- const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
- const auto *TLI = cast<NVPTXTargetLowering>(STI.getTargetLowering());
+template <typename OwnerT>
+static void printParam(const OwnerT *Owner, Type *Ty, unsigned AttrIdx,
+ bool IsByVal, bool IsKernel, StringRef Name,
+ const DataLayout &DL, raw_ostream &O) {
+ O << ".param ";
+
+ if (IsByVal || shouldPassAsArray(Ty)) {
+ const Align ParamAlign =
+ IsByVal && !IsKernel ? getDeviceByValParamAlign(Owner, Ty, AttrIdx, DL)
+ : getPTXParamAlign(Owner, Ty, AttrIdx, DL);
+ O << ".align " << ParamAlign.value() << " .b8 " << Name << "["
+ << DL.getTypeAllocSize(Ty) << "]";
+ return;
+ }
+
+ assert((Ty->isFloatingPointTy() || Ty->isIntOrPtrTy()) &&
+ "Unknown parameter type");
+ const unsigned Size = DL.getTypeSizeInBits(Ty).getFixedValue();
+ O << ".b"
+ << (IsKernel ? promoteScalarKernelArgumentSize(Size)
+ : promoteScalarArgumentSize(Size))
+ << " " << Name;
+}
+
+template <typename OwnerT>
+static void printReturnValClause(const OwnerT *Owner, StringRef Name,
+ const DataLayout &DL, raw_ostream &O) {
+ Type *RetTy = Owner->getFunctionType()->getReturnType();
- Type *Ty = F->getReturnType();
// A void or zero-sized return type (e.g. an empty struct) produces no return
// parameter.
- if (Ty->isVoidTy() || Ty->isEmptyTy())
+ if (RetTy->isVoidTy() || RetTy->isEmptyTy())
return;
- O << " (";
- auto PrintScalarRetVal = [&](unsigned Size) {
- O << ".param .b" << promoteScalarArgumentSize(Size) << " func_retval0";
- };
- if (shouldPassAsArray(Ty)) {
- const unsigned TotalSize = DL.getTypeAllocSize(Ty);
- const Align RetAlignment =
- getPTXParamAlign(F, Ty, AttributeList::ReturnIndex, DL);
- O << ".param .align " << RetAlignment.value() << " .b8 func_retval0["
- << TotalSize << "]";
- } else if (Ty->isFloatingPointTy()) {
- PrintScalarRetVal(Ty->getPrimitiveSizeInBits());
- } else if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
- PrintScalarRetVal(ITy->getBitWidth());
- } else if (isa<PointerType>(Ty)) {
- PrintScalarRetVal(TLI->getPointerTy(DL).getSizeInBits());
- } else
- llvm_unreachable("Unknown return type");
+ // Only device functions return a value, so no kernel promotion applies.
+ O << "(";
+ printParam(Owner, RetTy, AttributeList::ReturnIndex, /*IsByVal=*/false,
+ /*IsKernel=*/false, Name, DL, O);
O << ") ";
}
-void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
- raw_ostream &O) {
- const Function &F = MF.getFunction();
- 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() || RetTy->isEmptyTy()) {
- 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 << ") ";
- }
+ printReturnValClause(&CB, "_", DL, O);
O << "_ (";
auto MakeArg = [&](const unsigned I) {
- Type *Ty = CB.getArgOperand(I)->getType();
+ const bool IsByVal = CB.isByValArgument(I);
+ Type *Ty =
+ IsByVal ? CB.getParamByValType(I) : CB.getArgOperand(I)->getType();
- if (CB.isByValArgument(I)) {
- Type *ETy = CB.getParamByValType(I);
- Align ParamByValAlign = getDeviceByValParamAlign(
- &CB, ETy, I + AttributeList::FirstArgIndex, 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 << " _";
+ printParam(&CB, Ty, I + AttributeList::FirstArgIndex, IsByVal,
+ /*IsKernel=*/false, "_", DL, O);
};
const FunctionType *FTy = CB.getFunctionType();
@@ -897,7 +828,7 @@ void NVPTXAsmPrinter::emitFunctionEntryLabel() {
O << ".entry ";
else {
O << ".func ";
- printReturnValStr(*MF, O);
+ printReturnValClause(F, "func_retval0", getDataLayout(), O);
}
CurrentFnSym->print(O, MAI);
@@ -1066,11 +997,12 @@ void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
void NVPTXAsmPrinter::emitDeclarationWithName(const Function *F, MCSymbol *S,
raw_ostream &O) {
emitLinkageDirective(F, O);
- if (isKernelFunction(*F))
+ if (isKernelFunction(*F)) {
O << ".entry ";
- else
+ } else {
O << ".func ";
- printReturnValStr(F, O);
+ printReturnValClause(F, "func_retval0", getDataLayout(), O);
+ }
S->print(O, MAI);
O << "\n";
emitFunctionParamList(F, O);
@@ -1434,12 +1366,16 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
const PTXOpaqueType OpaqueType = getPTXOpaqueType(*GVar);
if (OpaqueType == PTXOpaqueType::Texture) {
- O << ".global .texref " << getTextureName(*GVar) << ";\n";
+ O << ".global .texref ";
+ getSymbol(GVar)->print(O, MAI);
+ O << ";\n";
return;
}
if (OpaqueType == PTXOpaqueType::Surface) {
- O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";
+ O << ".global .surfref ";
+ getSymbol(GVar)->print(O, MAI);
+ O << ";\n";
return;
}
@@ -1447,13 +1383,14 @@ void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
// (extern) declarations, no definition or initializer
// Currently the only known declaration is for an automatic __local
// (.shared) promoted to global.
- emitPTXGlobalVariable(GVar, O, STI);
+ emitPTXGlobalVariableDefinition(GVar, O, STI, /*EmitInitializer=*/false);
O << ";\n";
return;
}
if (OpaqueType == PTXOpaqueType::Sampler) {
- O << ".global .samplerref " << getSamplerName(*GVar);
+ O << ".global .samplerref ";
+ getSymbol(GVar)->print(O, MAI);
const Constant *Initializer = nullptr;
if (GVar->hasInitializer())
@@ -1542,7 +1479,6 @@ void NVPTXAsmPrinter::emitPTXGlobalVariableDefinition(
Type *ETy = GVar->getValueType();
- O << ".";
emitPTXAddressSpace(GVar->getAddressSpace(), O);
if (isManaged(*GVar)) {
@@ -1555,115 +1491,87 @@ void NVPTXAsmPrinter::emitPTXGlobalVariableDefinition(
O << " .align "
<< GVar->getAlign().value_or(DL.getPrefTypeAlign(ETy)).value();
+ const Constant *Initializer = nullptr;
+ if (GVar->hasInitializer()) {
+ const Constant *Init = GVar->getInitializer();
+ if (!Init->isNullValue() && !isa<UndefValue>(Init)) {
+ if (GVar->getAddressSpace() != ADDRESS_SPACE_GLOBAL &&
+ GVar->getAddressSpace() != ADDRESS_SPACE_CONST)
+ report_fatal_error("initial value of '" + GVar->getName() +
+ "' is not allowed in addrspace(" +
+ Twine(GVar->getAddressSpace()) + ")");
+ Initializer = Init;
+ }
+ }
+
if (ETy->isPointerTy() || ((ETy->isIntegerTy() || ETy->isFloatingPointTy()) &&
ETy->getScalarSizeInBits() <= 64)) {
- O << " .";
- // Special case: ABI requires that we use .u8 for predicates
- if (ETy->isIntegerTy(1))
- O << "u8";
- else
- O << getPTXFundamentalTypeStr(ETy, false);
- O << " ";
+ O << " ." << getPTXFundamentalTypeStr(ETy) << " ";
getSymbol(GVar)->print(O, MAI);
- // Ptx allows variable initilization only for constant and global state
- // spaces.
- if (EmitInitializer && GVar->hasInitializer()) {
- if ((GVar->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
- (GVar->getAddressSpace() == ADDRESS_SPACE_CONST)) {
- const Constant *Initializer = GVar->getInitializer();
- // 'undef' is treated as there is no value specified.
- if (!Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
- O << " = ";
- printScalarConstant(Initializer, O);
- }
- } else {
- // The frontend adds zero-initializer to device and constant variables
- // that don't have an initial value, and UndefValue to shared
- // variables, so skip warning for this case.
- if (!GVar->getInitializer()->isNullValue() &&
- !isa<UndefValue>(GVar->getInitializer())) {
- report_fatal_error("initial value of '" + GVar->getName() +
- "' is not allowed in addrspace(" +
- Twine(GVar->getAddressSpace()) + ")");
- }
- }
+ if (EmitInitializer && Initializer) {
+ O << " = ";
+ printScalarConstant(Initializer, O);
}
- } else {
- // Although PTX has direct support for struct type and array type and
- // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
- // targets that support these high level field accesses. Structs, arrays
- // and vectors are lowered into arrays of bytes.
- switch (ETy->getTypeID()) {
- case Type::IntegerTyID: // Integers larger than 64 bits
- case Type::FP128TyID:
- case Type::StructTyID:
- case Type::ArrayTyID:
- case Type::FixedVectorTyID: {
- const uint64_t ElementSize = DL.getTypeStoreSize(ETy);
- // Ptx allows variable initilization only for constant and
- // global state spaces.
- if (((GVar->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
- (GVar->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
- GVar->hasInitializer()) {
- const Constant *Initializer = GVar->getInitializer();
- if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) {
- AggBuffer aggBuffer(ElementSize, *this);
- bufferAggregateConstant(Initializer, &aggBuffer);
- if (aggBuffer.numSymbols()) {
- const unsigned int ptrSize = MAI.getCodePointerSize();
- if (ElementSize % ptrSize ||
- !aggBuffer.allSymbolsAligned(ptrSize)) {
- // Print in bytes and use the mask() operator for pointers.
- if (!STI.hasMaskOperator())
- report_fatal_error(
- "initialized packed aggregate with pointers '" +
- GVar->getName() +
- "' requires at least PTX ISA version 7.1");
- O << " .u8 ";
- getSymbol(GVar)->print(O, MAI);
- O << "[" << ElementSize << "]";
- if (EmitInitializer) {
- O << " = {";
- aggBuffer.printBytes(O);
- O << "}";
- }
- } else {
- O << " .u" << ptrSize * 8 << " ";
- getSymbol(GVar)->print(O, MAI);
- O << "[" << ElementSize / ptrSize << "]";
- if (EmitInitializer) {
- O << " = {";
- aggBuffer.printWords(O);
- O << "}";
- }
- }
- } else {
- O << " .b8 ";
- getSymbol(GVar)->print(O, MAI);
- O << "[" << ElementSize << "]";
- if (EmitInitializer) {
- O << " = {";
- aggBuffer.printBytes(O);
- O << "}";
- }
- }
- } else {
- O << " .b8 ";
- getSymbol(GVar)->print(O, MAI);
- if (ElementSize)
- O << "[" << ElementSize << "]";
- }
- } else {
- O << " .b8 ";
- getSymbol(GVar)->print(O, MAI);
- if (ElementSize)
- O << "[" << ElementSize << "]";
+ return;
+ }
+
+ // Although PTX has direct support for struct type and array type and LLVM IR
+ // is very similar to PTX, the LLVM CodeGen does not support for targets that
+ // support these high level field accesses. Structs, arrays and vectors are
+ // lowered into arrays of bytes.
+ assert((ETy->isIntegerTy() || ETy->isFP128Ty() || ETy->isAggregateType() ||
+ isa<FixedVectorType>(ETy)) &&
+ "type not supported yet");
+
+ const uint64_t ElementSize = DL.getTypeStoreSize(ETy);
+
+ if (!Initializer) {
+ O << " .b8 ";
+ getSymbol(GVar)->print(O, MAI);
+ if (ElementSize)
+ O << "[" << ElementSize << "]";
+ else if (!EmitInitializer)
+ O << "[]";
+ return;
+ }
+
+ AggBuffer aggBuffer(ElementSize, *this);
+ bufferAggregateConstant(Initializer, &aggBuffer);
+ if (aggBuffer.numSymbols()) {
+ const unsigned int ptrSize = MAI.getCodePointerSize();
+ if (ElementSize % ptrSize || !aggBuffer.allSymbolsAligned(ptrSize)) {
+ // Print in bytes and use the mask() operator for pointers.
+ if (!STI.hasMaskOperator())
+ report_fatal_error("initialized packed aggregate with pointers '" +
+ GVar->getName() +
+ "' requires at least PTX ISA version 7.1");
+ O << " .u8 ";
+ getSymbol(GVar)->print(O, MAI);
+ O << "[" << ElementSize << "]";
+ if (EmitInitializer) {
+ O << " = {";
+ aggBuffer.printBytes(O);
+ O << "}";
+ }
+ } else {
+ O << " .u" << ptrSize * 8 << " ";
+ getSymbol(GVar)->print(O, MAI);
+ O << "[" << ElementSize / ptrSize << "]";
+ if (EmitInitializer) {
+ O << " = {";
+ aggBuffer.printWords(O);
+ O << "}";
}
- break;
}
- default:
- llvm_unreachable("type not supported yet");
+ } else {
+ O << " .b8 ";
+ getSymbol(GVar)->print(O, MAI);
+ O << "[" << ElementSize << "]";
+ if (EmitInitializer) {
+ O << " = {";
+ aggBuffer.printBytes(O);
+ O << "}";
}
}
}
@@ -1769,129 +1677,67 @@ void NVPTXAsmPrinter::emitDemotedVars(const Function *F, raw_ostream &O) {
}
}
-void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
- raw_ostream &O) const {
+/// The PTX state space directive for \p AddressSpace, or an empty string if it
+/// does not name one, as is the case for the generic address space.
+static StringRef getPTXAddressSpaceName(unsigned AddressSpace) {
switch (AddressSpace) {
case ADDRESS_SPACE_LOCAL:
- O << "local";
- break;
+ return ".local";
case ADDRESS_SPACE_GLOBAL:
- O << "global";
- break;
+ return ".global";
case ADDRESS_SPACE_CONST:
- O << "const";
- break;
+ return ".const";
case ADDRESS_SPACE_SHARED:
- O << "shared";
- break;
+ return ".shared";
default:
+ return {};
+ }
+}
+
+/// The PTX opaque type directive for an image or sampler handle, or an empty
+/// string for PTXOpaqueType::None.
+static StringRef getPTXOpaqueTypeName(PTXOpaqueType OpaqueType) {
+ switch (OpaqueType) {
+ case PTXOpaqueType::Sampler:
+ return ".samplerref";
+ case PTXOpaqueType::Texture:
+ return ".texref";
+ case PTXOpaqueType::Surface:
+ return ".surfref";
+ case PTXOpaqueType::None:
+ return {};
+ }
+ llvm_unreachable("unexpected PTXOpaqueType");
+}
+
+void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
+ raw_ostream &O) const {
+ const StringRef Name = getPTXAddressSpaceName(AddressSpace);
+ if (Name.empty())
report_fatal_error("Bad address space found while emitting PTX: " +
llvm::Twine(AddressSpace));
- break;
- }
+ O << Name;
}
-std::string
-NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
+std::string NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty) const {
switch (Ty->getTypeID()) {
- case Type::IntegerTyID: {
- unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
- if (NumBits == 1)
- return "pred";
- if (NumBits <= 64) {
- std::string name = "u";
- return name + utostr(NumBits);
- }
- llvm_unreachable("Integer too large");
- break;
+ case Type::IntegerTyID:
+ case Type::PointerTyID: {
+ const uint64_t NumBits = getDataLayout().getTypeStoreSizeInBits(Ty);
+ assert(NumBits <= 64 && "type too large");
+ return "u" + utostr(promoteScalarKernelArgumentSize(NumBits));
}
case Type::BFloatTyID:
case Type::HalfTyID:
- // fp16 and bf16 are stored as .b16 for compatibility with pre-sm_53
- // PTX assembly.
- return "b16";
case Type::FloatTyID:
- return "f32";
case Type::DoubleTyID:
- return "f64";
- case Type::PointerTyID: {
- unsigned PtrSize = TM.getPointerSizeInBits(Ty->getPointerAddressSpace());
- assert((PtrSize == 64 || PtrSize == 32) && "Unexpected pointer size");
-
- if (PtrSize == 64)
- if (useB4PTR)
- return "b64";
- else
- return "u64";
- else if (useB4PTR)
- return "b32";
- else
- return "u32";
- }
+ return "b" + utostr(Ty->getScalarSizeInBits());
default:
break;
}
llvm_unreachable("unexpected type");
}
-void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
- raw_ostream &O,
- const NVPTXSubtarget &STI) {
- const DataLayout &DL = getDataLayout();
-
- // GlobalVariables are always constant pointers themselves.
- Type *ETy = GVar->getValueType();
-
- O << ".";
- emitPTXAddressSpace(GVar->getType()->getAddressSpace(), O);
- if (isManaged(*GVar)) {
- if (!STI.hasFeature(NVPTX::PTX40) || !STI.hasFeature(NVPTX::SM30))
- report_fatal_error(
- ".attribute(.managed) requires PTX version >= 4.0 and sm_30");
-
- O << " .attribute(.managed)";
- }
- O << " .align "
- << GVar->getAlign().value_or(DL.getPrefTypeAlign(ETy)).value();
-
- // Special case for i128/fp128
- if (ETy->getScalarSizeInBits() == 128) {
- O << " .b8 ";
- getSymbol(GVar)->print(O, MAI);
- O << "[16]";
- return;
- }
-
- if (ETy->isFloatingPointTy() || ETy->isIntOrPtrTy()) {
- O << " ." << getPTXFundamentalTypeStr(ETy) << " ";
- getSymbol(GVar)->print(O, MAI);
- return;
- }
-
- int64_t ElementSize = 0;
-
- // Although PTX has direct support for struct type and array type and LLVM IR
- // is very similar to PTX, the LLVM CodeGen does not support for targets that
- // support these high level field accesses. Structs and arrays are lowered
- // into arrays of bytes.
- switch (ETy->getTypeID()) {
- case Type::StructTyID:
- case Type::ArrayTyID:
- case Type::FixedVectorTyID:
- ElementSize = DL.getTypeStoreSize(ETy);
- O << " .b8 ";
- getSymbol(GVar)->print(O, MAI);
- O << "[";
- if (ElementSize) {
- O << ElementSize;
- }
- O << "]";
- break;
- default:
- llvm_unreachable("type not supported yet");
- }
-}
-
void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
const DataLayout &DL = getDataLayout();
const NVPTXSubtarget &STI = TM.getSubtarget<NVPTXSubtarget>(*F);
@@ -1899,7 +1745,6 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
const NVPTXMachineFunctionInfo *MFI =
MF ? MF->getInfo<NVPTXMachineFunctionInfo>() : nullptr;
- bool IsFirst = true;
const bool IsKernelFunc = isKernelFunction(*F);
// Zero-sized arguments (e.g. empty structs) do not produce a parameter.
@@ -1918,138 +1763,57 @@ void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
O << "(\n";
- for (const auto &[ParamIndex, Arg] : enumerate(NonEmptyArgs)) {
+ auto MakeParam = [&](const auto &IndexedArg) {
+ const auto &[ParamIndex, Arg] = IndexedArg;
Type *Ty = Arg.getType();
MCSymbol *const ParamSym = TLI->getParamSymbol(OutContext, F, ParamIndex);
- if (!IsFirst)
- O << ",\n";
-
- IsFirst = false;
-
- // Handle image/sampler parameters
- if (IsKernelFunc) {
- const PTXOpaqueType ArgOpaqueType = getPTXOpaqueType(Arg);
- if (ArgOpaqueType != PTXOpaqueType::None) {
- const bool EmitImgPtr = !MFI || !MFI->checkImageHandleSymbol(ParamSym);
- O << "\t.param ";
- if (EmitImgPtr)
+ O << "\t";
+
+ // A byval param is passed as a copy of the pointee and an aggregate is
+ // passed as a blob of bytes; both are declared as a byte array.
+ const bool IsByVal = Arg.hasByValAttr();
+ const bool AsArray = IsByVal || shouldPassAsArray(Ty);
+
+ // Kernels declare image/sampler handles and the address space of a
+ // pointee. Both of those are scalar handles, so a byte-array param is
+ // neither.
+ if (IsKernelFunc && !AsArray) {
+ const StringRef OpaqueType = getPTXOpaqueTypeName(getPTXOpaqueType(Arg));
+ if (!OpaqueType.empty()) {
+ O << ".param ";
+ if (!MFI || !MFI->checkImageHandleSymbol(ParamSym))
O << ".u64 .ptr ";
- switch (ArgOpaqueType) {
- case PTXOpaqueType::Sampler:
- O << ".samplerref ";
- break;
- case PTXOpaqueType::Texture:
- O << ".texref ";
- break;
- case PTXOpaqueType::Surface:
- O << ".surfref ";
- break;
- case PTXOpaqueType::None:
- llvm_unreachable("handled above");
- }
- O << *ParamSym;
- continue;
+ O << OpaqueType << " " << *ParamSym;
+ return;
}
- }
- if (Arg.hasByValAttr()) {
- // param has byVal attribute.
- Type *ETy = Arg.getParamByValType();
- assert(ETy && "Param should have byval type");
-
- // Print .param .align <a> .b8 .param[size];
- // <a> = optimal alignment for the element type; always multiple of
- // PAL.getParamAlignment
- // size = typeallocsize of element type
- const unsigned ParamIdx = Arg.getArgNo() + AttributeList::FirstArgIndex;
- const Align OptimalAlign =
- IsKernelFunc ? getPTXParamAlign(F, ETy, ParamIdx, DL)
- : getDeviceByValParamAlign(F, ETy, ParamIdx, DL);
-
- O << "\t.param .align " << OptimalAlign.value() << " .b8 " << *ParamSym
- << "[" << DL.getTypeAllocSize(ETy) << "]";
- continue;
- }
-
- if (shouldPassAsArray(Ty)) {
- // Just print .param .align <a> .b8 .param[size];
- // <a> = optimal alignment for the element type; always multiple of
- // PAL.getParamAlignment
- // size = typeallocsize of element type
- Align OptimalAlign = getPTXParamAlign(
- F, Ty, Arg.getArgNo() + AttributeList::FirstArgIndex, DL);
+ if (auto *PTy = dyn_cast<PointerType>(Ty)) {
+ const unsigned AS = PTy->getAddressSpace();
+ O << ".param .u" << DL.getPointerSizeInBits(AS) << " .ptr";
- O << "\t.param .align " << OptimalAlign.value() << " .b8 " << *ParamSym
- << "[" << DL.getTypeAllocSize(Ty) << "]";
-
- continue;
- }
- // Just a scalar
- auto *PTy = dyn_cast<PointerType>(Ty);
- unsigned PTySizeInBits = 0;
- if (PTy) {
- PTySizeInBits =
- TLI->getPointerTy(DL, PTy->getAddressSpace()).getSizeInBits();
- assert(PTySizeInBits && "Invalid pointer size");
- }
-
- if (IsKernelFunc) {
- if (PTy) {
- O << "\t.param .u" << PTySizeInBits << " .ptr";
-
- switch (PTy->getAddressSpace()) {
- default:
- break;
- case ADDRESS_SPACE_GLOBAL:
- O << " .global";
- break;
- case ADDRESS_SPACE_SHARED:
- O << " .shared";
- break;
- case ADDRESS_SPACE_CONST:
- O << " .const";
- break;
- case ADDRESS_SPACE_LOCAL:
- O << " .local";
- break;
- }
+ const StringRef Space = getPTXAddressSpaceName(AS);
+ if (!Space.empty())
+ O << " " << Space;
O << " .align " << Arg.getParamAlign().valueOrOne().value() << " "
<< *ParamSym;
- continue;
+ return;
}
-
- // non-pointer scalar to kernel func
- O << "\t.param .";
- // Special case: predicate operands become .u8 types
- if (Ty->isIntegerTy(1))
- O << "u8";
- else
- O << getPTXFundamentalTypeStr(Ty);
- O << " " << *ParamSym;
- continue;
}
- // Non-kernel function, just print .param .b<size> for ABI
- // and .reg .b<size> for non-ABI
- unsigned Size;
- if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
- Size = promoteScalarArgumentSize(ITy->getBitWidth());
- } else if (PTy) {
- assert(PTySizeInBits && "Invalid pointer size");
- Size = PTySizeInBits;
- } else
- Size = Ty->getPrimitiveSizeInBits();
- O << "\t.param .b" << Size << " " << *ParamSym;
- }
-
- if (F->isVarArg()) {
- if (!IsFirst)
- O << ",\n";
- O << "\t.param .align " << STI.getMaxRequiredAlignment() << " .b8 "
+
+ printParam(F, IsByVal ? Arg.getParamByValType() : Ty,
+ Arg.getArgNo() + AttributeList::FirstArgIndex, IsByVal,
+ IsKernelFunc, ParamSym->getName(), DL, O);
+ };
+
+ interleave(enumerate(NonEmptyArgs), O, MakeParam, ",\n");
+
+ if (F->isVarArg())
+ O << (NonEmptyArgs.empty() ? "" : ",\n") << "\t.param .align "
+ << STI.getMaxRequiredAlignment() << " .b8 "
<< *TLI->getParamSymbol(OutContext, F, /* vararg */ -1) << "[]";
- }
O << "\n)";
}
@@ -2121,24 +1885,16 @@ void NVPTXAsmPrinter::encodeDebugInfoRegisterNumbers(
void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp,
raw_ostream &O) const {
- APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
- bool ignored;
- unsigned int numHex;
- const char *lead;
-
- if (Fp->getType()->getTypeID() == Type::FloatTyID) {
- numHex = 8;
- lead = "0f";
- APF.convert(APFloat::IEEEsingle(), APFloat::rmNearestTiesToEven, &ignored);
- } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
- numHex = 16;
- lead = "0d";
- APF.convert(APFloat::IEEEdouble(), APFloat::rmNearestTiesToEven, &ignored);
- } else
+ if (Fp->getType()->isFloatTy())
+ O << "0f";
+ else if (Fp->getType()->isDoubleTy())
+ O << "0d";
+ else
llvm_unreachable("unsupported fp type");
- APInt API = APF.bitcastToAPInt();
- O << lead << format_hex_no_prefix(API.getZExtValue(), numHex, /*Upper=*/true);
+ const APInt API = Fp->getValueAPF().bitcastToAPInt();
+ O << format_hex_no_prefix(API.getZExtValue(), API.getBitWidth() / 4,
+ /*Upper=*/true);
}
void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
@@ -2147,7 +1903,10 @@ void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
return;
}
if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
- printFPConstant(CFP, O);
+ const APInt API = CFP->getValueAPF().bitcastToAPInt();
+ O << "0x"
+ << format_hex_no_prefix(API.getZExtValue(), API.getBitWidth() / 4,
+ /*Upper=*/true);
return;
}
if (isa<ConstantPointerNull>(CPV)) {
diff --git a/llvm/lib/Target/NVPTX/NVPTXUtilities.h b/llvm/lib/Target/NVPTX/NVPTXUtilities.h
index ae3692fecafae..2567ea793095b 100644
--- a/llvm/lib/Target/NVPTX/NVPTXUtilities.h
+++ b/llvm/lib/Target/NVPTX/NVPTXUtilities.h
@@ -58,14 +58,14 @@ Align getPTXParamAlign(const CallBase *CB, Type *Ty, unsigned AttrIdx,
// PTX ABI requires all scalar argument/return values to have
// bit-size as a power of two of at least 32 bits.
-inline unsigned promoteScalarArgumentSize(unsigned size) {
- if (size <= 32)
- return 32;
- if (size <= 64)
- return 64;
- if (size <= 128)
- return 128;
- return size;
+inline unsigned promoteScalarArgumentSize(unsigned Size) {
+ assert(Size < 128 && "Size should be less than 128 (shouldPassAsArray)");
+ return PowerOf2Ceil(std::max(Size, 32U));
+}
+
+inline unsigned promoteScalarKernelArgumentSize(unsigned Size) {
+ assert(Size < 128 && "Size should be less than 128 (shouldPassAsArray)");
+ return PowerOf2Ceil(std::max(Size, 8U));
}
inline bool shouldPassAsArray(Type *Ty) {
diff --git a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
index 2a36c4162ba02..11b12f9b07394 100644
--- a/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
+++ b/llvm/test/CodeGen/NVPTX/call_bitcast_byval.ll
@@ -39,7 +39,7 @@ define void @boom() {
; CHECK-DAG: .param .align 2 .b8 param0[4];
; CHECK-DAG: st.param.b16 [param0], %rs{{[0-9]+}};
; CHECK-DAG: st.param.b16 [param0+2], %rs{{[0-9]+}};
- ; CHECK-DAG: .callprototype ()_ (.param .align 2 .b8 _[4]);
+ ; CHECK-DAG: .callprototype _ (.param .align 2 .b8 _[4]);
call void %fp(ptr byval(%"class.complex") null)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/callchain.ll b/llvm/test/CodeGen/NVPTX/callchain.ll
index 847d8b80ad38a..7b4fcb33f161d 100644
--- a/llvm/test/CodeGen/NVPTX/callchain.ll
+++ b/llvm/test/CodeGen/NVPTX/callchain.ll
@@ -4,7 +4,7 @@
target triple = "nvptx"
define void @foo(ptr %ptr) {
-; CHECK: prototype_0 : .callprototype ()_ ()
+; CHECK: prototype_0 : .callprototype _ ()
tail call void %ptr()
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/empty-type.ll b/llvm/test/CodeGen/NVPTX/empty-type.ll
index f29c3d7c2e21b..960aad7d8d46f 100644
--- a/llvm/test/CodeGen/NVPTX/empty-type.ll
+++ b/llvm/test/CodeGen/NVPTX/empty-type.ll
@@ -132,7 +132,7 @@ define void @indirect_caller(ptr %fp) {
; CHECK-LABEL: indirect_caller(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-NEXT: prototype_4 : .callprototype ()_ (.param .b32 _);
+; CHECK-NEXT: prototype_4 : .callprototype _ (.param .b32 _);
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [indirect_caller_param_0];
; CHECK-NEXT: { // callseq 4, 0
diff --git a/llvm/test/CodeGen/NVPTX/fp-global.ll b/llvm/test/CodeGen/NVPTX/fp-global.ll
new file mode 100644
index 0000000000000..8282538eb970a
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/fp-global.ll
@@ -0,0 +1,55 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
+
+; Floating-point globals are declared with a bit type and initialized with
+; their bit pattern. PTX has no 16-bit floating-point literal, so this is the
+; only form available for half and bfloat, and the .f32 and .f64 literals are
+; themselves just bit patterns in hex.
+
+; CHECK-DAG: .visible .global .align 2 .b16 h = 0x3C00;
+ at h = addrspace(1) global half 0xH3C00
+
+; CHECK-DAG: .visible .global .align 2 .b16 bf = 0x3F80;
+ at bf = addrspace(1) global bfloat 0xR3F80
+
+; CHECK-DAG: .visible .global .align 4 .b32 f = 0x3F800000;
+ at f = addrspace(1) global float 1.0
+
+; CHECK-DAG: .visible .global .align 8 .b64 d = 0x4000000000000000;
+ at d = addrspace(1) global double 2.0
+
+; The bit pattern is padded out to the width of the type.
+; CHECK-DAG: .visible .global .align 2 .b16 h_denormal = 0x0001;
+ at h_denormal = addrspace(1) global half 0xH0001
+
+; Signed zeros and NaNs round-trip as the bits they are.
+; CHECK-DAG: .visible .global .align 4 .b32 f_negzero = 0x80000000;
+ at f_negzero = addrspace(1) global float -0.0
+
+; CHECK-DAG: .visible .global .align 8 .b64 d_nan = 0x7FF8000000000000;
+ at d_nan = addrspace(1) global double 0x7FF8000000000000
+
+; A zero initializer is treated as no value specified.
+; CHECK-DAG: .visible .global .align 2 .b16 h_zero;
+ at h_zero = addrspace(1) global half 0xH0000
+
+; A declaration has no initializer, but must still agree on the type.
+; CHECK-DAG: .extern .global .align 4 .b32 f_decl;
+ at f_decl = external addrspace(1) global float
+
+; CHECK-DAG: .visible .const .align 4 .b32 f_const = 0x40600000;
+ at f_const = addrspace(4) global float 3.5
+
+; Aggregates keep going through the byte-buffered path.
+; CHECK-DAG: .visible .global .align 2 .b8 h_arr[4] = {0, 60, 0, 64};
+ at h_arr = addrspace(1) global [2 x half] [half 0xH3C00, half 0xH4000]
+
+define ptx_kernel void @use(ptr %p) {
+ %v = load float, ptr addrspace(1) @f_decl
+ ; Instruction immediates still need a floating-point literal; ptxas rejects
+ ; an integer where an .f32 immediate is expected.
+ ; CHECK: add.rn.f32 %r{{[0-9]+}}, %r{{[0-9]+}}, 0f3FC00000;
+ %a = fadd float %v, 1.5
+ store float %a, ptr %p
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/i1-param.ll b/llvm/test/CodeGen/NVPTX/i1-param.ll
index 14d417bca459d..bfecde95786f1 100644
--- a/llvm/test/CodeGen/NVPTX/i1-param.ll
+++ b/llvm/test/CodeGen/NVPTX/i1-param.ll
@@ -7,7 +7,7 @@ target triple = "nvptx-nvidia-cuda"
; Make sure predicate (i1) operands to kernels get expanded out to .u8
; CHECK: .entry foo
-; CHECK: .param .u8 foo_param_0
+; CHECK: .param .b8 foo_param_0
; CHECK: .param .u64 .ptr .align 1 foo_param_1
define ptx_kernel void @foo(i1 %p, ptr %out) {
%val = zext i1 %p to i32
diff --git a/llvm/test/CodeGen/NVPTX/managed.ll b/llvm/test/CodeGen/NVPTX/managed.ll
index 7a246b7ccbed9..0fffccc6352d1 100644
--- a/llvm/test/CodeGen/NVPTX/managed.ll
+++ b/llvm/test/CodeGen/NVPTX/managed.ll
@@ -11,7 +11,7 @@
; CHECK: .extern .global .align 4 .u32 decl_g;
@decl_g = external addrspace(1) global i32, align 4
-; CHECK: .extern .global .attribute(.managed) .align 8 .b64 managed_decl_g;
+; CHECK: .extern .global .attribute(.managed) .align 8 .u64 managed_decl_g;
@managed_decl_g = external addrspace(1) global ptr, align 8
; A cyclic managed definition needs an exactly matching managed declaration.
diff --git a/llvm/test/CodeGen/NVPTX/noreturn.ll b/llvm/test/CodeGen/NVPTX/noreturn.ll
index 0062e62756d36..3d1149fa4dc33 100644
--- a/llvm/test/CodeGen/NVPTX/noreturn.ll
+++ b/llvm/test/CodeGen/NVPTX/noreturn.ll
@@ -32,7 +32,7 @@ define ptx_kernel void @ignore_kernel_noreturn() #0 {
}
; CHECK-LABEL: .entry callprototype_noreturn(
-; CHECK: prototype_{{[0-9]+}} : .callprototype ()_ (.param .b32 _) .noreturn;
+; CHECK: prototype_{{[0-9]+}} : .callprototype _ (.param .b32 _) .noreturn;
; CHECK: prototype_{{[0-9]+}} : .callprototype (.param .b32 _) _ (.param .b32 _);
define ptx_kernel void @callprototype_noreturn(i32) {
diff --git a/llvm/test/CodeGen/NVPTX/odd-integer-width.ll b/llvm/test/CodeGen/NVPTX/odd-integer-width.ll
new file mode 100644
index 0000000000000..1ccb272d97667
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/odd-integer-width.ll
@@ -0,0 +1,37 @@
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -verify-machineinstrs | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -verify-machineinstrs | %ptxas-verify -arch=sm_90 %}
+
+; PTX only has 8, 16, 32, and 64-bit integer types, so a value whose width is
+; not one of those is named by the type of the storage it occupies.
+
+; CHECK: .visible .global .align 1 .u8 g2;
+ at g2 = addrspace(1) global i2 0
+
+; CHECK: .visible .global .align 4 .u32 g24;
+ at g24 = addrspace(1) global i24 0
+
+; An i1 is stored as a byte. It must not be declared as .pred, which is only
+; valid in the register state space.
+
+; CHECK: .extern .global .align 1 .u8 g1;
+ at g1 = external addrspace(1) global i1
+
+; Kernel parameters keep their storage size, which is part of the launch ABI.
+
+; CHECK-LABEL: .visible .entry kernel(
+; CHECK-NEXT: .param .b8 kernel_param_0,
+; CHECK-NEXT: .param .b32 kernel_param_1,
+; CHECK-NEXT: .param .b64 kernel_param_2
+define ptx_kernel void @kernel(i2 %a, i24 %b, i48 %c) {
+ ret void
+}
+
+; Device function parameters are promoted to the PTX ABI's 32-bit minimum.
+
+; CHECK-LABEL: .visible .func device(
+; CHECK-NEXT: .param .b32 device_param_0,
+; CHECK-NEXT: .param .b32 device_param_1,
+; CHECK-NEXT: .param .b64 device_param_2
+define void @device(i2 %a, i24 %b, i48 %c) {
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/param-align.ll b/llvm/test/CodeGen/NVPTX/param-align.ll
index c85080fdf295a..993e5b673c7ad 100644
--- a/llvm/test/CodeGen/NVPTX/param-align.ll
+++ b/llvm/test/CodeGen/NVPTX/param-align.ll
@@ -55,17 +55,17 @@ declare ptr @getfp(i32 %n)
define ptx_device void @t6() {
; CHECK: .func t6
%fp = call ptr @getfp(i32 0)
-; CHECK: prototype_2 : .callprototype ()_ (.param .align 8 .b8 _[8]);
+; CHECK: prototype_2 : .callprototype _ (.param .align 8 .b8 _[8]);
call void %fp(ptr byval(double) null);
%fp2 = call ptr @getfp(i32 1)
-; NOALIGN4: prototype_4 : .callprototype ()_ (.param .align 2 .b8 _[4]);
-; ALIGN4: prototype_4 : .callprototype ()_ (.param .align 4 .b8 _[4]);
+; NOALIGN4: prototype_4 : .callprototype _ (.param .align 2 .b8 _[4]);
+; ALIGN4: prototype_4 : .callprototype _ (.param .align 4 .b8 _[4]);
call void %fp(ptr byval(%struct.half2) null);
%fp3 = call ptr @getfp(i32 2)
-; NOALIGN4: prototype_6 : .callprototype ()_ (.param .align 1 .b8 _[1]);
-; ALIGN4: prototype_6 : .callprototype ()_ (.param .align 4 .b8 _[1]);
+; NOALIGN4: prototype_6 : .callprototype _ (.param .align 1 .b8 _[1]);
+; ALIGN4: prototype_6 : .callprototype _ (.param .align 4 .b8 _[1]);
call void %fp(ptr byval(i8) null);
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/vector-global.ll b/llvm/test/CodeGen/NVPTX/vector-global.ll
index bb8af20ba4fac..9c93d5e18e7aa 100644
--- a/llvm/test/CodeGen/NVPTX/vector-global.ll
+++ b/llvm/test/CodeGen/NVPTX/vector-global.ll
@@ -8,3 +8,10 @@ target triple = "nvptx64-nvidia-cuda"
; CHECK: .extern .global .align 16 .b8 g1[16];
@g2 = global <4 x i32> zeroinitializer ; module-level global variable
; CHECK: .visible .global .align 16 .b8 g2[16];
+
+; A vector is sized by the whole vector, not by its element type, even when the
+; element is 128 bits wide. The declaration must agree with the definition.
+ at g3 = external global <2 x i128>
+; CHECK: .extern .global .align 32 .b8 g3[32];
+ at g4 = global <2 x i128> zeroinitializer
+; CHECK: .visible .global .align 32 .b8 g4[32];
diff --git a/llvm/test/DebugInfo/NVPTX/debug-info.ll b/llvm/test/DebugInfo/NVPTX/debug-info.ll
index 2f7763526f344..277e80ad00996 100644
--- a/llvm/test/DebugInfo/NVPTX/debug-info.ll
+++ b/llvm/test/DebugInfo/NVPTX/debug-info.ll
@@ -13,8 +13,8 @@
; CHECK: .target sm_{{[0-9]+}}, debug
; CHECK: .visible .entry _Z5saxpyifPfS_(
-; CHECK: .param .u32 {{.+}},
-; CHECK: .param .f32 {{.+}},
+; CHECK: .param .b32 {{.+}},
+; CHECK: .param .b32 {{.+}},
; CHECK: .param .u64 {{.+}},
; CHECK: .param .u64 {{.+}}
; CHECK: )
>From 115158a932ae0291cb0d6bc7fd11f907c76eb563 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Mon, 28 Sep 2026 12:49:18 -0700
Subject: [PATCH 2/2] [NVPTX] Add AsmPrinter test coverage requested in review
---
llvm/test/CodeGen/NVPTX/fp-global.ll | 22 +++++++++++++++++++++-
llvm/test/CodeGen/NVPTX/ptr-global.ll | 22 ++++++++++++++++++++++
2 files changed, 43 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/NVPTX/ptr-global.ll
diff --git a/llvm/test/CodeGen/NVPTX/fp-global.ll b/llvm/test/CodeGen/NVPTX/fp-global.ll
index 8282538eb970a..507f6e5826ec5 100644
--- a/llvm/test/CodeGen/NVPTX/fp-global.ll
+++ b/llvm/test/CodeGen/NVPTX/fp-global.ll
@@ -22,10 +22,16 @@
; CHECK-DAG: .visible .global .align 2 .b16 h_denormal = 0x0001;
@h_denormal = addrspace(1) global half 0xH0001
-; Signed zeros and NaNs round-trip as the bits they are.
+; Signed zeros, infinities, and NaNs round-trip as the bits they are.
; CHECK-DAG: .visible .global .align 4 .b32 f_negzero = 0x80000000;
@f_negzero = addrspace(1) global float -0.0
+; CHECK-DAG: .visible .global .align 4 .b32 f_inf = 0x7F800000;
+ at f_inf = addrspace(1) global float +inf
+
+; CHECK-DAG: .visible .global .align 4 .b32 f_neginf = 0xFF800000;
+ at f_neginf = addrspace(1) global float -inf
+
; CHECK-DAG: .visible .global .align 8 .b64 d_nan = 0x7FF8000000000000;
@d_nan = addrspace(1) global double 0x7FF8000000000000
@@ -33,6 +39,9 @@
; CHECK-DAG: .visible .global .align 2 .b16 h_zero;
@h_zero = addrspace(1) global half 0xH0000
+; CHECK-DAG: .visible .global .align 2 .b16 bf_zero;
+ at bf_zero = addrspace(1) global bfloat 0xR0000
+
; A declaration has no initializer, but must still agree on the type.
; CHECK-DAG: .extern .global .align 4 .b32 f_decl;
@f_decl = external addrspace(1) global float
@@ -44,6 +53,7 @@
; CHECK-DAG: .visible .global .align 2 .b8 h_arr[4] = {0, 60, 0, 64};
@h_arr = addrspace(1) global [2 x half] [half 0xH3C00, half 0xH4000]
+; CHECK-LABEL: .entry use(
define ptx_kernel void @use(ptr %p) {
%v = load float, ptr addrspace(1) @f_decl
; Instruction immediates still need a floating-point literal; ptxas rejects
@@ -53,3 +63,13 @@ define ptx_kernel void @use(ptr %p) {
store float %a, ptr %p
ret void
}
+
+; CHECK-LABEL: .entry use_double(
+define ptx_kernel void @use_double(ptr %p) {
+ %v = load double, ptr %p
+ ; Double instruction immediates likewise need an .f64 literal.
+ ; CHECK: add.rn.f64 %rd{{[0-9]+}}, %rd{{[0-9]+}}, 0d3FF8000000000000;
+ %a = fadd double %v, 1.5
+ store double %a, ptr %p
+ ret void
+}
diff --git a/llvm/test/CodeGen/NVPTX/ptr-global.ll b/llvm/test/CodeGen/NVPTX/ptr-global.ll
new file mode 100644
index 0000000000000..7f729a0e4022d
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/ptr-global.ll
@@ -0,0 +1,22 @@
+; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=PTX32
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s --check-prefix=PTX64
+; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
+
+; Pointer globals use the same unsigned integer type in declarations and
+; definitions, with or without an initializer.
+; PTX32-DAG: .extern .global .align 4 .u32 ptr_decl;
+; PTX64-DAG: .extern .global .align 8 .u64 ptr_decl;
+ at ptr_decl = external addrspace(1) global ptr
+
+; PTX32-DAG: .visible .global .align 4 .u32 ptr_null;
+; PTX64-DAG: .visible .global .align 8 .u64 ptr_null;
+ at ptr_null = addrspace(1) global ptr null
+
+; PTX32-DAG: .visible .global .align 4 .u32 value;
+; PTX64-DAG: .visible .global .align 4 .u32 value;
+ at value = addrspace(1) global i32 0
+
+; PTX32-DAG: .visible .global .align 4 .u32 ptr_init = value;
+; PTX64-DAG: .visible .global .align 8 .u64 ptr_init = value;
+ at ptr_init = addrspace(1) global ptr addrspace(1) @value
More information about the llvm-commits
mailing list