[llvm-branch-commits] [llvm] [7/7][PISA] Add PISA AsmPrinter (PR #214647)
Michal Paszkowski via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Sep 22 04:51:49 PDT 2026
https://github.com/michalpaszkowski updated https://github.com/llvm/llvm-project/pull/214647
>From cfb0ce00abefd4991ead2b4d83c9a70e8f26c5fe Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Mon, 3 Aug 2026 04:18:55 -0700
Subject: [PATCH 01/11] Add PISA AsmPrinter
Add PISAAsmPrinter to complete the codegen pipeline with end-to-end code
emission.
---
llvm/lib/Target/PISA/CMakeLists.txt | 1 +
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 933 ++++++++++++++++++++++++
llvm/test/CodeGen/PISA/load.ll | 350 +++++++++
llvm/test/CodeGen/PISA/return.ll | 37 +
llvm/test/CodeGen/PISA/store.ll | 352 +++++++++
llvm/test/CodeGen/PISA/unreachable.ll | 14 +
6 files changed, 1687 insertions(+)
create mode 100644 llvm/lib/Target/PISA/PISAAsmPrinter.cpp
create mode 100644 llvm/test/CodeGen/PISA/load.ll
create mode 100644 llvm/test/CodeGen/PISA/return.ll
create mode 100644 llvm/test/CodeGen/PISA/store.ll
create mode 100644 llvm/test/CodeGen/PISA/unreachable.ll
diff --git a/llvm/lib/Target/PISA/CMakeLists.txt b/llvm/lib/Target/PISA/CMakeLists.txt
index 94f0f3ecc6489..399052b3d3a74 100644
--- a/llvm/lib/Target/PISA/CMakeLists.txt
+++ b/llvm/lib/Target/PISA/CMakeLists.txt
@@ -18,6 +18,7 @@ tablegen(LLVM PISAGenPostLegalizeGICombiner.inc -gen-global-isel-combiner
add_public_tablegen_target(PISACommonTableGen)
add_llvm_target(PISACodeGen
+ PISAAsmPrinter.cpp
PISACacheCtrlMMRA.cpp
PISACacheHintSelector.cpp
PISACallLowering.cpp
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
new file mode 100644
index 0000000000000..512d1d2efafc3
--- /dev/null
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -0,0 +1,933 @@
+//===-- PISAAsmPrinter.cpp - PISA LLVM assembly writer --------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "MCTargetDesc/PISAInstPrinter.h"
+#include "MCTargetDesc/PISARegEncoder.h"
+#include "MCTargetDesc/PISATargetStreamer.h"
+#include "PISA.h"
+#include "PISAInstrInfo.h"
+#include "PISAMCInstLower.h"
+#include "PISAMachineFunctionInfo.h"
+#include "PISARegManager.h"
+#include "PISASubtarget.h"
+#include "PISATargetMachine.h"
+#include "PISAUtils.h"
+#include "TargetInfo/PISATargetInfo.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/STLExtras.h"
+#include "llvm/ADT/SmallSet.h"
+#include "llvm/Analysis/ConstantFolding.h"
+#include "llvm/Analysis/ValueTracking.h"
+#include "llvm/CodeGen/AsmPrinter.h"
+#include "llvm/CodeGen/MachineConstantPool.h"
+#include "llvm/CodeGen/MachineFrameInfo.h"
+#include "llvm/CodeGen/MachineFunction.h"
+#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/MachineInstr.h"
+#include "llvm/CodeGen/MachineModuleInfo.h"
+#include "llvm/CodeGen/MachineRegisterInfo.h"
+#include "llvm/CodeGen/TargetLoweringObjectFileImpl.h"
+#include "llvm/IR/IRPrintingPasses.h"
+#include "llvm/IR/Module.h"
+#include "llvm/MC/MCAsmInfo.h"
+#include "llvm/MC/MCInst.h"
+#include "llvm/MC/MCStreamer.h"
+#include "llvm/MC/MCSymbol.h"
+#include "llvm/MC/MCValue.h"
+#include "llvm/MC/TargetRegistry.h"
+#include "llvm/Support/Compiler.h"
+#include "llvm/Support/Endian.h"
+#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/FileSystem.h"
+#include "llvm/Support/PISAAddrSpace.h"
+#include "llvm/Support/Regex.h"
+#include "llvm/Support/raw_ostream.h"
+#include "llvm/TargetParser/PISATargetParser.h"
+#include <llvm/IR/DiagnosticInfo.h>
+
+using namespace llvm;
+
+#define DEBUG_TYPE "asm-printer"
+
+namespace {
+class PISAAsmPrinter : public AsmPrinter {
+
+public:
+ PISATargetStreamer &getTargetStreamer() const {
+ return static_cast<PISATargetStreamer &>(*OutStreamer->getTargetStreamer());
+ }
+
+private:
+ void collectRegDcls(PISA::RegDcls &);
+ void collectLocalVariableDcls(PISA::LocalVariableDcls &);
+ void updateFuncParamIdxs(PISA::DataTypes &DTs);
+
+ void outputInstruction(const MachineInstr *MI);
+ void printOperand(const MachineInstr *MI, int OpNum, raw_ostream &O);
+
+ std::string getVirtualRegisterName(Register R) const;
+
+ void collectFunctionDeclaration(PISA::FunctionDeclaration &,
+ const Function &F);
+ void collectFunctionSignature(PISA::FunctionSignature &);
+ void collectFunctionParameters(PISA::FunctionSignature &);
+ void collectKernelParameters(PISA::FunctionSignature &);
+ void collectFunctionDirectiveAndName(PISA::FunctionDirectiveAndName &DN,
+ const Function &F);
+ PISA::LinkageTy collectLinkage(const GlobalValue &V);
+ void collectGlobalVariable(PISA::GlobalVariableDcl &PGV,
+ const GlobalVariable &GV);
+
+ void emitGlobalsAndFuncDecls(Module &M);
+
+ const PISASubtarget *ST = nullptr;
+ const PISAInstrInfo *TII = nullptr;
+ const PISARegisterInfo *TRI = nullptr;
+ PISA::RegManager *RegMgr = nullptr;
+ PISA::DataTypes *DTs = nullptr;
+ bool GlobalsEmitted = false;
+
+ class FlattenGlobal {
+ public:
+ FlattenGlobal(const Constant *C, PISA::VariableInit &VI,
+ const DataLayout &DL, AsmPrinter &AP)
+ : DL(DL), VI(VI), AP(AP) {
+ process(C);
+ dischargeZeros();
+ assert(computeSize(C) == DL.getTypeAllocSize(C->getType()) &&
+ "size mismatch?");
+ }
+
+ private:
+ bool isZero(const Constant *C) const {
+ if (isa<ConstantPointerNull>(C))
+ return false;
+
+ return C->isNullValue() || isa<UndefValue>(C);
+ }
+ void pad(const Constant *C, unsigned NumElts = 0) {
+ unsigned Size = DL.getTypeAllocSize(C->getType());
+ if (NumElts == 0) {
+ ZeroCnt += Size;
+ return;
+ }
+ unsigned EmittedSize =
+ DL.getTypeAllocSize(C->getType()->getContainedType(0)) * NumElts;
+ assert(EmittedSize <= Size && "Size cannot be less than EmittedSize!");
+ if (unsigned Padding = Size - EmittedSize)
+ ZeroCnt += Padding;
+ }
+ void pad(uint64_t NumBytes) { ZeroCnt += NumBytes; }
+ void dischargeZeros() {
+ if (ZeroCnt == 0)
+ return;
+ // Insert dummy slot
+ VI.Initializer.push_back({LLT{}, 0});
+ uint64_t Idx = VI.Initializer.size() - 1;
+ VI.Exprs.insert({Idx, PISA::VariableInit::Zeros{ZeroCnt}});
+ ZeroCnt = 0;
+ }
+ void addVal(LLT Ty, uint64_t Val) {
+ dischargeZeros();
+ VI.Initializer.push_back({Ty, Val});
+ }
+ void addGlobal(LLT Ty, const PISA::VariableInit::GlobalExpr &GE) {
+ // Insert dummy slot
+ addVal(Ty, 0);
+ uint64_t Idx = VI.Initializer.size() - 1;
+ VI.Exprs.insert({Idx, GE});
+ }
+ void lowerConstant(const Constant *C) {
+ auto *Expr = AP.lowerConstant(C);
+ MCValue Res;
+ if (!Expr->evaluateAsRelocatable(Res, nullptr))
+ llvm_unreachable("unhandled expression!");
+ LLT Ty = getLLTForType(*C->getType(), DL);
+ if (!Res.getAddSym() && !Res.getSubSym()) {
+ if (Res.getConstant() == 0)
+ pad(C);
+ else
+ addVal(Ty, static_cast<uint64_t>(Res.getConstant()));
+ return;
+ }
+ assert(!Res.getSubSym() && "unhandled expression!");
+ std::string Name = Res.getAddSym()->getName().str();
+ PISA::VariableInit::GlobalExpr E{std::move(Name), Res.getConstant()};
+ addGlobal(Ty, E);
+ }
+ uint64_t computeSize(const Constant *C) const {
+ uint64_t Total = 0;
+ for (auto [i, Elt] : llvm::enumerate(VI.Initializer)) {
+ if (auto Iter = VI.Exprs.find(i); Iter != VI.Exprs.end()) {
+ auto &Entry = Iter->second;
+ if (auto *Z = std::get_if<PISA::VariableInit::Zeros>(&Entry)) {
+ Total += Z->N;
+ continue;
+ }
+ }
+ Total += Elt.Type.getSizeInBytes();
+ }
+ return Total;
+ }
+ void emitGlobalConstantLargeInt(const ConstantInt *CI) {
+ unsigned BitWidth = CI->getBitWidth();
+
+ // Copy the value as we may massage the layout for constants whose bit
+ // width is not a multiple of 64-bits.
+ APInt Realigned(CI->getValue());
+ uint64_t ExtraBits = 0;
+ unsigned ExtraBitsSize = BitWidth & 63;
+
+ if (ExtraBitsSize) {
+ // The bit width of the data is not a multiple of 64-bits.
+ // The extra bits are expected to be at the end of the chunk of the
+ // memory. Little endian:
+ // * Nothing to be done, just record the extra bits to emit.
+ ExtraBits = Realigned.getRawData()[BitWidth / 64];
+ }
+
+ // We don't expect assemblers to support integer data directives
+ // for more than 64 bits, so we emit the data in at most 64-bit
+ // quantities at a time.
+ const uint64_t *RawData = Realigned.getRawData();
+ for (unsigned I = 0, E = BitWidth / 64; I != E; ++I)
+ addVal(LLT::integer(64), RawData[I]);
+
+ if (ExtraBitsSize) {
+ // Emit the extra bits after the 64-bits chunks.
+ // Emit a directive that fills the expected size.
+ uint64_t Size = DL.getTypeStoreSize(CI->getType());
+ Size -= (BitWidth / 64) * 8;
+ assert(Size && Size * 8 >= ExtraBitsSize &&
+ (ExtraBits & (((uint64_t)-1) >> (64 - ExtraBitsSize))) ==
+ ExtraBits &&
+ "Directive too small for extra bits.");
+ addVal(LLT::integer(Size * 8), ExtraBits);
+ }
+ }
+ const DataLayout &DL;
+ PISA::VariableInit &VI;
+ AsmPrinter &AP;
+ void process(const Constant *C);
+ unsigned ZeroCnt = 0;
+ };
+
+protected:
+ bool doInitialization(Module &M) override;
+ bool doFinalization(Module &M) override;
+
+public:
+ explicit PISAAsmPrinter(TargetMachine &TM,
+ std::unique_ptr<MCStreamer> Streamer)
+ : AsmPrinter(TM, std::move(Streamer)) {}
+
+ StringRef getPassName() const override { return "PISA Assembly Printer"; }
+ bool PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
+ const char *ExtraCode, raw_ostream &O) override;
+
+ void emitInstruction(const MachineInstr *MI) override;
+ void emitFunctionHeader() override;
+ void emitFunctionBodyStart() override;
+ void emitFunctionBodyEnd() override;
+ void emitEndOfAsmFile(Module &) override;
+
+ void emitFunctionEntryLabel() override {}
+ void emitBasicBlockEnd(const MachineBasicBlock &MBB) override {}
+ void emitGlobalVariable(const GlobalVariable *GV) override {}
+
+ bool runOnMachineFunction(MachineFunction &MF) override;
+};
+} // namespace
+
+void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
+ uint64_t Size = DL.getTypeAllocSize(C->getType());
+ if (isZero(C))
+ return pad(C);
+ auto AddSplatVector = [&](LLT ScalarTy, const APInt &EltVal) {
+ assert(EltVal.getBitWidth() <= 64 && "Splat element too wide for uint64_t");
+ auto *VTy = cast<FixedVectorType>(C->getType());
+ unsigned NumElts = VTy->getNumElements();
+ uint64_t Val = EltVal.getZExtValue();
+ for (unsigned I = 0; I < NumElts; ++I)
+ addVal(ScalarTy, Val);
+ pad(C, NumElts);
+ };
+
+ if (auto *CI = dyn_cast<ConstantInt>(C)) {
+ if (C->getType()->isVectorTy()) {
+ auto *VTy = cast<FixedVectorType>(C->getType());
+ uint64_t EltAllocSize = DL.getTypeAllocSize(VTy->getElementType());
+ LLT ScalarTy = LLT::integer(EltAllocSize * 8);
+ AddSplatVector(ScalarTy, CI->getValue());
+ } else {
+ // We don't use the LLT type of `C` directly here because `C` could be,
+ // for example, a s1. The allocation size is 1, so we want to give it
+ // a type of s8 to reflect that.
+ addVal(LLT::integer(Size * 8), CI->getZExtValue());
+ }
+ } else if (auto *FP = dyn_cast<ConstantFP>(C)) {
+ if (C->getType()->isVectorTy()) {
+ auto *VTy = cast<FixedVectorType>(C->getType());
+ LLT ScalarTy = getLLTForType(*VTy->getElementType(), DL);
+ AddSplatVector(ScalarTy, FP->getValueAPF().bitcastToAPInt());
+ } else {
+ LLT Ty = getLLTForType(*C->getType(), DL);
+ addVal(Ty, FP->getValueAPF().bitcastToAPInt().getZExtValue());
+ }
+ } else if (isa<ConstantPointerNull>(C)) {
+ LLT Ty = getLLTForType(*C->getType(), DL);
+ unsigned AS = C->getType()->getPointerAddressSpace();
+ if (uint64_t Val = PISATargetMachine::getNullPointerValue(AS))
+ addVal(Ty, Val);
+ else
+ pad(C);
+ } else if (auto *CV = dyn_cast<ConstantVector>(C)) {
+ Type *ElementType = CV->getType()->getElementType();
+ uint64_t ElementSizeInBits = DL.getTypeSizeInBits(ElementType);
+ uint64_t ElementAllocSizeInBits = DL.getTypeAllocSizeInBits(ElementType);
+ if (ElementSizeInBits != ElementAllocSizeInBits) {
+ // If the allocation size of an element is different from the size in
+ // bits, printing each element separately will insert incorrect padding.
+ //
+ // The general algorithm here is complicated; instead of writing it out
+ // here, just use the existing code in ConstantFolding.
+ Type *IntT = IntegerType::get(CV->getContext(),
+ DL.getTypeSizeInBits(CV->getType()));
+ ConstantInt *CI = dyn_cast_or_null<ConstantInt>(ConstantFoldConstant(
+ ConstantExpr::getBitCast(const_cast<ConstantVector *>(CV), IntT),
+ DL));
+ if (!CI) {
+ report_fatal_error(
+ "Cannot lower vector global with unusual element type");
+ }
+ emitGlobalConstantLargeInt(CI);
+ uint64_t EmittedSize = DL.getTypeStoreSize(CV->getType());
+ if (unsigned Padding = Size - EmittedSize)
+ pad(Padding);
+ } else {
+ for (unsigned I = 0; I < CV->getNumOperands(); I++)
+ process(CV->getAggregateElement(I));
+ pad(C, CV->getNumOperands());
+ }
+ } else if (auto *CA = dyn_cast<ConstantArray>(C)) {
+ for (unsigned I = 0; I < CA->getNumOperands(); I++)
+ process(CA->getAggregateElement(I));
+ } else if (auto *CS = dyn_cast<ConstantStruct>(C)) {
+ auto *StructTy = cast<StructType>(CS->getType());
+ auto *Layout = DL.getStructLayout(StructTy);
+ for (unsigned I = 0, E = CS->getNumOperands(); I != E; ++I) {
+ const Constant *Field = CS->getOperand(I);
+ // Print the actual field value.
+ process(Field);
+ // Check if padding is needed and insert one or more 0s.
+ uint64_t FieldSize = DL.getTypeAllocSize(Field->getType());
+ uint64_t PadSize =
+ ((I == E - 1 ? Size : Layout->getElementOffset(I + 1)) -
+ Layout->getElementOffset(I)) -
+ FieldSize;
+ // Insert padding - this may include padding to increase the size of the
+ // current field up to the ABI size (if the struct is not packed) as well
+ // as padding to ensure that the next field starts at the right offset.
+ pad(PadSize);
+ }
+ } else if (auto *CDS = dyn_cast<ConstantDataSequential>(C)) {
+ for (unsigned I = 0; I < CDS->getNumElements(); I++)
+ process(CDS->getElementAsConstant(I));
+ pad(C, CDS->getNumElements());
+ } else if (const ConstantExpr *CE = dyn_cast<ConstantExpr>(C)) {
+ // Look through bitcasts, which might not be able to be MCExpr'ized (e.g.
+ // of vectors).
+ if (CE->getOpcode() == Instruction::BitCast)
+ return process(CE->getOperand(0));
+ if (Size > 8) {
+ // If the constant expression's size is greater than 64-bits, then we
+ // have to emit the value in chunks. Try to constant fold the value and
+ // emit it that way.
+ Constant *New = ConstantFoldConstant(CE, DL);
+ if (New != CE)
+ return process(New);
+ }
+ lowerConstant(C);
+ } else if (isa<GlobalVariable>(C) || isa<Function>(C)) {
+ assert(Size == 8 && "global symbol with non 64-bit size?");
+ lowerConstant(C);
+ } else {
+ llvm_unreachable("unhandled constant!");
+ }
+}
+
+static bool isIgnoredIntrinsicGlobal(const GlobalVariable &GV) {
+ if (GV.getName() == "llvm.used")
+ return true;
+
+ // Ignore debug and non-emitted data. This handles llvm.compiler.used.
+ if (GV.getSection() == "llvm.metadata")
+ return true;
+
+ // Skip globals only used as annotation strings by llvm.ptr.annotation.
+ // These are metadata for the annotation intrinsic, not real data.
+ if (GV.hasPrivateLinkage() && GV.isConstant() &&
+ all_of(GV.users(), [](const User *U) {
+ if (auto *CE = dyn_cast<ConstantExpr>(U))
+ return all_of(CE->users(), [](const User *UU) {
+ auto *CI = dyn_cast<CallInst>(UU);
+ return CI && CI->getCalledFunction() &&
+ CI->getCalledFunction()->getIntrinsicID() ==
+ Intrinsic::ptr_annotation;
+ });
+ auto *CI = dyn_cast<CallInst>(U);
+ return CI && CI->getCalledFunction() &&
+ CI->getCalledFunction()->getIntrinsicID() ==
+ Intrinsic::ptr_annotation;
+ }))
+ return true;
+
+ if (!GV.hasAppendingLinkage())
+ return false;
+
+ if (GV.getName() == "llvm.global_ctors")
+ report_fatal_error(
+ "llvm.global_ctors is not supported by the PISA backend");
+
+ if (GV.getName() == "llvm.global_dtors")
+ report_fatal_error(
+ "llvm.global_ctors is not supported by the PISA backend");
+
+ report_fatal_error("unknown special variable with appending linkage");
+}
+
+void PISAAsmPrinter::emitGlobalsAndFuncDecls(Module &M) {
+ PISATargetStreamer &TS = getTargetStreamer();
+
+ // emit header info
+ // - we always emit in latest PISA syntax
+ auto GetHdrTarget = [&]() -> SmallString<16> {
+ return ST ? ST->getPISATargetName() : "";
+ };
+ PISA::HeaderDcl HD = {PISA::LatestPISAVersion, GetHdrTarget()};
+ TS.emitHeader(HD);
+ OutStreamer->addBlankLine();
+
+ // Emit Module level function decl
+ for (auto &F : M) {
+ if (!F.isDeclaration() || F.isIntrinsic()) // avoid llvm builtins
+ continue;
+
+ PISA::FunctionDeclaration Dcl;
+
+ collectFunctionDeclaration(Dcl, F);
+ TS.emitFunctionDeclaration(Dcl);
+ OutStreamer->addBlankLine();
+ }
+
+ // Translate global variables
+ for (auto &GV : M.globals()) {
+ if (isIgnoredIntrinsicGlobal(GV))
+ continue;
+
+ PISA::GlobalVariableDcl PGV;
+ collectGlobalVariable(PGV, GV);
+ TS.emitGlobalVariable(PGV);
+ }
+}
+
+bool PISAAsmPrinter::doInitialization(Module &M) {
+ GlobalsEmitted = false;
+ return AsmPrinter::doInitialization(M);
+}
+
+bool PISAAsmPrinter::doFinalization(Module &M) {
+ // If we did not emit any functions, then the global declarations have not
+ // yet been emitted.
+ if (!GlobalsEmitted) {
+ emitGlobalsAndFuncDecls(M);
+ GlobalsEmitted = true;
+ }
+ return AsmPrinter::doFinalization(M);
+}
+
+bool PISAAsmPrinter::runOnMachineFunction(MachineFunction &MF) {
+ ST = &MF.getSubtarget<PISASubtarget>();
+ TII = ST->getInstrInfo();
+ TRI = ST->getRegisterInfo();
+
+ if (!GlobalsEmitted) {
+ emitGlobalsAndFuncDecls(*MF.getFunction().getParent());
+ GlobalsEmitted = true;
+ }
+
+ PISA::RegManager Mgr{MF};
+ RegMgr = &Mgr;
+
+ return AsmPrinter::runOnMachineFunction(MF);
+}
+
+void PISAAsmPrinter::emitFunctionHeader() {
+ const Function &F = MF->getFunction();
+
+ auto *Section = getObjFileLowering().SectionForGlobal(&F, TM);
+ MF->setSection(Section);
+}
+
+void PISAAsmPrinter::updateFuncParamIdxs(PISA::DataTypes &DTs) {
+ // Update DataTypes records of RegStart for body register
+ // declarations (vs the already-processed func param dcls)
+ DTs.finalizeFuncParams();
+
+ llvm::DenseMap<std::tuple</*NumElts=*/unsigned, /*BitWidth=*/unsigned,
+ /*Type=*/unsigned>,
+ /*Index=*/unsigned>
+ ParamIdxs;
+
+ auto &MRI = MF->getRegInfo();
+ for (auto &[CurReg, Info] : RegMgr->mapping()) {
+ // We are only trying to update indices for function parameters
+ if (!(Info.Flags & PISA::RegManager::NoEmissionDef))
+ continue;
+
+ auto *RC = MRI.getRegClass(CurReg);
+ unsigned BitWidth = TRI->getBitSizeFromRegClass(RC);
+ unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
+ auto [It, Inserted] =
+ ParamIdxs.try_emplace(std::make_tuple(NumElts, BitWidth, Info.Type), 0);
+ RegMgr->setRegIdx(CurReg, It->second++);
+
+ // Sanity check that all function parameter indexes are < the total
+ // number of function parameters of that type (recorded in DTs)
+ [[maybe_unused]] bool ValidIdx =
+ Info.Idx < DTs.getInfo(NumElts, BitWidth, Info.Type).RegCounter;
+ assert(ValidIdx && "function parameter index out of range!");
+ }
+}
+
+void PISAAsmPrinter::collectRegDcls(PISA::RegDcls &Dcls) {
+ auto &MRI = MF->getRegInfo();
+ for (auto &[CurReg, Info] : RegMgr->mapping()) {
+ if (Info.Flags & PISA::RegManager::NoEmissionDef)
+ continue;
+ auto *RC = MRI.getRegClass(CurReg);
+ unsigned BitWidth = TRI->getBitSizeFromRegClass(RC);
+ unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
+ TypeInfo &TI = DTs->emplaceInfo(NumElts, BitWidth, Info.Type);
+ auto Bank = RegMgr->getRegBank(NumElts, BitWidth);
+ const char *Prefix =
+ RegMgr->getPrefixFromBank(static_cast<PISA::RegManager::RegBank>(Bank));
+ Dcls.Regs[std::make_tuple(NumElts, BitWidth, Info.Type)].push_back(
+ std::make_pair(Prefix, TI.RegCounter));
+ RegMgr->setRegIdx(CurReg, TI.RegCounter);
+ TI.RegCounter++;
+ }
+}
+
+void PISAAsmPrinter::collectLocalVariableDcls(PISA::LocalVariableDcls &Dcls) {
+ auto &MFI = MF->getFrameInfo();
+ for (int Idx = MFI.getObjectIndexBegin(), EndIdx = MFI.getObjectIndexEnd();
+ Idx != EndIdx; ++Idx) {
+ if (MFI.isDeadObjectIndex(Idx))
+ continue;
+ // translation of 'alloca' creates a local 0-sized object within current
+ // frame. Since PISA does not create real frame, omit such objects.
+ if (!MFI.getObjectSize(Idx))
+ continue;
+ PISA::VariableDcl VarDecl;
+ VarDecl.Linkage = PISA::LinkageTy::DEFAULT;
+ switch (MFI.getStackID(Idx)) {
+ case TargetStackID::Default:
+ VarDecl.SS = PISA::StorageSpace::PRIVATE;
+ break;
+ case TargetStackID::PISAShared:
+ VarDecl.SS = PISA::StorageSpace::SHARED;
+ break;
+ default:
+ llvm_unreachable("unknown stack ID!");
+ }
+
+ VarDecl.Size = MFI.getObjectSize(Idx);
+ VarDecl.Alignment = MFI.getObjectAlign(Idx);
+ VarDecl.StackIndex = Idx;
+
+ Dcls.Vars.push_back(std::move(VarDecl));
+ }
+}
+
+PISA::LinkageTy PISAAsmPrinter::collectLinkage(const GlobalValue &V) {
+ if (V.hasLocalLinkage())
+ return PISA::LinkageTy::DEFAULT;
+
+ // global variable linkage
+ if (auto *GVar = dyn_cast<GlobalVariable>(&V)) {
+ // External GV with no initializer must be .import. In llvm, global
+ // variable definitions must be initialized. Though PISA allows
+ // a GV definition with no initializer, we can safely determine the
+ // linkage by having initializer or not here
+ return GVar->hasInitializer() ? PISA::LinkageTy::EXPORT
+ : PISA::LinkageTy::IMPORT;
+ }
+
+ // function variable linkage
+ return V.isDeclaration() ? PISA::LinkageTy::IMPORT : PISA::LinkageTy::EXPORT;
+}
+
+static void collectIntelHostAccessMetadata(PISA::GlobalVariableDcl &PGV,
+ const GlobalVariable &GV) {
+ // !intel_host_access !{i32 <HostAccessQualifier>, !"<Name>"}
+ // -> .host_access("Name")
+ MDNode *MD = GV.getMetadata("intel_host_access");
+ if (!MD || MD->getNumOperands() < 2)
+ return;
+
+ auto *NameMD = dyn_cast<MDString>(MD->getOperand(1));
+ if (!NameMD)
+ return;
+
+ PGV.Dcl.HostAccessName = NameMD->getString().str();
+}
+
+void PISAAsmPrinter::collectGlobalVariable(PISA::GlobalVariableDcl &PGV,
+ const GlobalVariable &GV) {
+ auto &DL = GV.getParent()->getDataLayout();
+ PGV.Dcl.Linkage = collectLinkage(GV);
+ PGV.Dcl.SS =
+ PISA::mapAddrSpaceToStorageSpace(GV.getType()->getAddressSpace());
+ PGV.Dcl.Alignment = DL.getPreferredAlign(&GV);
+ PGV.Dcl.Name = getSymbol(&GV)->getName();
+ PGV.Dcl.Size = DL.getTypeAllocSize(GV.getValueType());
+ PGV.Dcl.Section = GV.getSection();
+
+ collectIntelHostAccessMetadata(PGV, GV);
+
+ if (GV.hasInitializer() && !isa<UndefValue>(GV.getInitializer()))
+ FlattenGlobal FG{GV.getInitializer(), PGV.Init, DL, *this};
+}
+
+void PISAAsmPrinter::collectFunctionDeclaration(PISA::FunctionDeclaration &Dcl,
+ const Function &F) {
+ assert(F.isDeclaration());
+ assert(F.getCallingConv() != CallingConv::PISA_KERNEL);
+
+ collectFunctionDirectiveAndName(Dcl.DN, F);
+ for (auto &P : F.args()) {
+ PISA::FunctionDeclParam Param;
+ if (P.getType()->getScalarSizeInBits() == 1) {
+ Param.Ty = LLT::integer(8);
+ } else {
+ Param.Ty = getLLTForType(*P.getType(), F.getParent()->getDataLayout());
+ }
+ Dcl.FunctionParams.push_back(Param);
+ }
+}
+
+void PISAAsmPrinter::collectKernelParameters(PISA::FunctionSignature &Sig) {
+ const PISAMachineFunctionInfo *MFInfo =
+ MF->getInfo<PISAMachineFunctionInfo>();
+
+ // print params
+ auto &DL = MF->getFunction().getParent()->getDataLayout();
+ for (unsigned Index = 0; Index < MF->getFunction().arg_size(); ++Index) {
+ auto [Size, IsByRef] = MFInfo->getArgInfo(Index);
+ PISA::KernelParameter Param;
+ Param.Size = Size;
+ auto *ArgTy = MF->getFunction().getArg(Index)->getType()->getScalarType();
+ auto Align = alignTo(PowerOf2Ceil(DL.getABITypeAlign(ArgTy).value()), 4);
+ if (Align != 8) {
+ // Kernel parameters are aligned to 8 bytes by default.
+ Param.Align = Align;
+ }
+ if (ArgTy->isPointerTy() && !IsByRef) {
+ auto AS = ArgTy->getPointerAddressSpace();
+ if ((AS == (unsigned)PISAAS::AddressSpace::CONSTANT) ||
+ (AS == (unsigned)PISAAS::AddressSpace::GLOBAL) ||
+ (AS == (unsigned)PISAAS::AddressSpace::SHARED))
+ Param.AS = ArgTy->getPointerAddressSpace();
+ if (auto PtrAlign = MF->getFunction().getParamAlign(Index))
+ Param.PtrAlign = PtrAlign->value();
+ }
+
+ // Read OpenCL kernel arg metadata (emitted by the OpenCL frontend with
+ // -cl-kernel-arg-info).
+ const Function &F = MF->getFunction();
+ if (MDNode *MD = F.getMetadata("kernel_arg_name"))
+ if (Index < MD->getNumOperands())
+ if (auto *S = dyn_cast<MDString>(MD->getOperand(Index)))
+ if (!S->getString().empty())
+ Param.ArgName = S->getString().str();
+
+ Sig.KernelParams.push_back(std::move(Param));
+ }
+}
+
+void PISAAsmPrinter::collectFunctionParameters(PISA::FunctionSignature &Sig) {
+ llvm::SmallVector<const MachineInstr *, 8> FuncParamInsts;
+ for (auto &MBB : *MF) {
+ // FunctionParam must be contiguous and in the same BB
+ // Find the iterator of the first FunctionParam inst and iterate from it
+ // to collect all FunctionParam insts
+ auto MIIt = find_if(
+ MBB, [&](MachineInstr &MI) { return TII->isFunctionParamInstr(MI); });
+
+ for (; MIIt != MBB.end(); ++MIIt) {
+ if (!TII->isFunctionParamInstr(*MIIt))
+ break;
+ FuncParamInsts.push_back(&*MIIt);
+ }
+ }
+
+ // Sort FunctionParam by param index
+ llvm::sort(FuncParamInsts, [](const MachineInstr *L, const MachineInstr *R) {
+ return L->getOperand(1).getImm() < R->getOperand(1).getImm();
+ });
+
+ // Collect params
+ for (auto *MI : FuncParamInsts) {
+ const MachineOperand &MO = MI->getOperand(0);
+ assert(MO.getSubReg() == 0 && "no swizzle allowed on args!");
+ const auto *RC = MF->getRegInfo().getRegClass(MO.getReg());
+ PISA::FunctionParameter Param;
+ unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
+ unsigned EltSize = TRI->getBitSizeFromRegClass(RC);
+ TypeInfo &TI = DTs->emplaceInfo(NumElts, EltSize, PISA::RegEncoder::REG);
+ Param.Ty = TI.Ty;
+ Param.Prefix = TI.Prefix;
+ Param.Idx = TI.RegCounter++;
+ Sig.FunctionParams.push_back(std::move(Param));
+ }
+}
+
+static std::string getNameFromType(Type *Ty, bool IsSigned) {
+ std::string Name = "unknown";
+ switch (Ty->getTypeID()) {
+ default:
+ llvm_unreachable("unsupported type");
+ break;
+ case Type::IntegerTyID: {
+ switch (Ty->getIntegerBitWidth()) {
+ default:
+ llvm_unreachable("unsupported integer type");
+ break;
+ case 8:
+ Name = IsSigned ? "char" : "uchar";
+ break;
+ case 16:
+ Name = IsSigned ? "short" : "ushort";
+ break;
+ case 32:
+ Name = IsSigned ? "int" : "uint";
+ break;
+ case 64:
+ Name = IsSigned ? "long" : "ulong";
+ break;
+ }
+ } break;
+ case Type::HalfTyID:
+ Name = "half";
+ break;
+ case Type::FloatTyID:
+ Name = "float";
+ break;
+ case Type::DoubleTyID:
+ Name = "double";
+ break;
+ case Type::FixedVectorTyID: {
+ auto *VecTy = cast<FixedVectorType>(Ty);
+ Name = getNameFromType(VecTy->getElementType(), IsSigned) +
+ std::to_string(VecTy->getNumElements());
+ } break;
+ }
+ return Name;
+}
+
+void PISAAsmPrinter::collectFunctionDirectiveAndName(
+ PISA::FunctionDirectiveAndName &DN, const Function &F) {
+
+ DN.CC = F.getCallingConv();
+ if (DN.CC != CallingConv::PISA_KERNEL)
+ DN.Linkage = collectLinkage(F);
+ DN.Name = getSymbol(&F)->getName();
+ if (DN.CC != CallingConv::PISA_KERNEL) {
+ Type *RetType = F.getReturnType();
+ if (!RetType->isVoidTy()) {
+ if (RetType->getScalarSizeInBits() == 1) {
+ DN.RetLLT = LLT::integer(8);
+ } else {
+ DN.RetLLT =
+ llvm::getLLTForType(*RetType, F.getParent()->getDataLayout());
+ }
+ }
+ }
+
+ std::vector<std::pair<StringRef, PISA::KernelAttributeType>>
+ AvailableKernelMetadataNodeTypes = {
+ {"reqd_work_group_size",
+ PISA::KernelAttributeType::REQD_WORK_GROUP_SIZE},
+ {"vec_type_hint", PISA::KernelAttributeType::VEC_TYPE_HINT}};
+
+ for (auto [MetadataName, EnumVal] : AvailableKernelMetadataNodeTypes) {
+ MDNode *Node = dyn_cast_or_null<MDNode>(F.getMetadata(MetadataName));
+ if (!Node)
+ continue;
+ auto &KernelAttr = DN.KernelAttrs.emplace_back();
+ KernelAttr.KernelAttrType = EnumVal;
+ switch (EnumVal) {
+ case llvm::PISA::KernelAttributeType::REQD_WORK_GROUP_SIZE: {
+ KernelAttr.KernelAttrValues.emplace<std::vector<uint32_t>>();
+ std::transform(Node->op_begin(), Node->op_end(),
+ std::back_inserter(std::get<std::vector<uint32_t>>(
+ KernelAttr.KernelAttrValues)),
+ [](const MDOperand &Operand) -> uint32_t {
+ const ValueAsMetadata *OperandAsVal =
+ cast<ValueAsMetadata>(Operand);
+ ConstantInt *OperandVal =
+ cast<ConstantInt>(OperandAsVal->getValue());
+ return static_cast<uint32_t>(OperandVal->getZExtValue());
+ });
+ } break;
+ case llvm::PISA::KernelAttributeType::VEC_TYPE_HINT: {
+ Metadata *Op0 = Node->getOperand(0);
+ Metadata *Op1 = Node->getOperand(1);
+ ConstantInt *CI =
+ cast<ConstantInt>(cast<ValueAsMetadata>(Op1)->getValue());
+ auto TypeName =
+ getNameFromType(cast<ValueAsMetadata>(Op0)->getType(), CI->isOne());
+ KernelAttr.KernelAttrValues.emplace<std::string>(TypeName);
+ } break;
+ }
+ }
+}
+
+void PISAAsmPrinter::collectFunctionSignature(PISA::FunctionSignature &Sig) {
+ Function &F = MF->getFunction();
+ collectFunctionDirectiveAndName(Sig.DN, F);
+ if (Sig.DN.CC == CallingConv::PISA_KERNEL)
+ collectKernelParameters(Sig);
+ else
+ collectFunctionParameters(Sig);
+}
+
+void PISAAsmPrinter::emitFunctionBodyStart() {
+ PISATargetStreamer &TS = getTargetStreamer();
+ PISA::FunctionSignature Sig;
+ PISA::DataTypes FuncDTs;
+ DTs = &FuncDTs;
+ // Collect function signature.
+ collectFunctionSignature(Sig);
+ TS.emitFunctionSignature(Sig);
+ updateFuncParamIdxs(*DTs);
+ // Emit the function body start.
+ TS.emitFuncBodyStart();
+ // Collect local registers.
+ PISA::RegDcls Regs;
+ collectRegDcls(Regs);
+ TS.emitRegDcls(Regs, *DTs);
+ // Collect local variables.
+ PISA::LocalVariableDcls Vars;
+ collectLocalVariableDcls(Vars);
+ TS.emitLocalVariableDcls(Vars);
+}
+
+void PISAAsmPrinter::emitFunctionBodyEnd() {
+ PISATargetStreamer &TS = getTargetStreamer();
+ // Emit the function body end.
+ TS.emitFuncBodyEnd();
+}
+
+void PISAAsmPrinter::emitEndOfAsmFile(llvm::Module &) {}
+
+void PISAAsmPrinter::printOperand(const MachineInstr *MI, int OpNum,
+ raw_ostream &O) {
+ const MachineOperand &MO = MI->getOperand(OpNum);
+
+ switch (MO.getType()) {
+ case MachineOperand::MO_Register: {
+ auto Reg = MO.getReg();
+ if (Reg.isPhysical())
+ O << PISAInstPrinter::getRegisterName(Reg);
+ else {
+ O << getVirtualRegisterName(Reg);
+ O << TRI->getSwizzleName(MO.getSubReg());
+ }
+ } break;
+
+ case MachineOperand::MO_Immediate:
+ O << MO.getImm();
+ break;
+
+ case MachineOperand::MO_FPImmediate:
+ O << MO.getFPImm();
+ break;
+
+ case MachineOperand::MO_MachineBasicBlock:
+ O << *MO.getMBB()->getSymbol();
+ break;
+
+ case MachineOperand::MO_GlobalAddress:
+ O << *getSymbol(MO.getGlobal());
+ break;
+
+ case MachineOperand::MO_BlockAddress: {
+ MCSymbol *BA = GetBlockAddressSymbol(MO.getBlockAddress());
+ O << BA->getName();
+ break;
+ }
+
+ case MachineOperand::MO_ExternalSymbol:
+ O << *GetExternalSymbolSymbol(MO.getSymbolName());
+ break;
+
+ case MachineOperand::MO_JumpTableIndex:
+ case MachineOperand::MO_ConstantPoolIndex:
+ default:
+ llvm_unreachable("<unknown operand type>");
+ }
+}
+
+bool PISAAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
+ const char *ExtraCode, raw_ostream &O) {
+ if (ExtraCode && ExtraCode[0])
+ return true; // Invalid instruction - PISA does not have special
+ // modifiers
+
+ printOperand(MI, OpNo, O);
+ return false;
+}
+
+std::string PISAAsmPrinter::getVirtualRegisterName(Register R) const {
+ auto &MRI = MF->getRegInfo();
+ const auto *RC = MRI.getRegClass(R);
+
+ std::string Name;
+ raw_string_ostream O(Name);
+
+ unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
+ unsigned EltSize = TRI->getBitSizeFromRegClass(RC);
+ RegEncoder::RegBank Bank = RegMgr->getRegBank(NumElts, EltSize);
+ O << RegMgr->getPrefixFromBank(Bank) << RegMgr->getRegIdx(R);
+ return Name;
+}
+
+void PISAAsmPrinter::outputInstruction(const MachineInstr *MI) {
+ PISAMCInstLower MCInstLowering{OutContext, *TRI, *RegMgr, *this};
+ PISAMCInst Inst;
+ MCInstLowering.lower(MI, Inst);
+ if (MI->getOpcode() == PISA::DBG_VALUE)
+ return;
+ OutStreamer->emitInstruction(Inst, *OutContext.getSubtargetInfo());
+}
+
+void PISAAsmPrinter::emitInstruction(const MachineInstr *MI) {
+ PISA_MC::verifyInstructionPredicates(MI->getOpcode(),
+ getSubtargetInfo().getFeatureBits());
+
+ if (!TII->isNoEmissionInstr(*MI))
+ outputInstruction(MI);
+}
+
+// Force static initialization.
+// NOLINTNEXTLINE(readability-identifier-naming)
+extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializePISAAsmPrinter() {
+ RegisterAsmPrinter<PISAAsmPrinter> Y(getThePISATarget());
+}
diff --git a/llvm/test/CodeGen/PISA/load.ll b/llvm/test/CodeGen/PISA/load.ll
new file mode 100644
index 0000000000000..b3fa8ad1ac728
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/load.ll
@@ -0,0 +1,350 @@
+; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -O0 < %s -march=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O0 %s
+; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O2 %s
+; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs
+
+define i8 @i8(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .8b @i8(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.8b [[R8_B0]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R8_B0]];
+ %2 = load i8, ptr addrspace(4) %0, align 1
+ ret i8 %2
+}
+
+define i16 @i16(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .16b @i16(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.16b [[R16_H0]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R16_H0]];
+ %2 = load i16, ptr addrspace(4) %0, align 2
+ ret i16 %2
+}
+
+define i32 @i32(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .32b @i32(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.32b [[R32_W1]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R32_W1]];
+ %2 = load i32, ptr addrspace(4) %0, align 4
+ ret i32 %2
+}
+
+define i64 @i64(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .64b @i64(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.64b [[R64_D0]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R64_D0]];
+ %2 = load i64, ptr addrspace(4) %0, align 8
+ ret i64 %2
+}
+
+define float @f32(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .32b @f32(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.32b [[R32_W1]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R32_W1]];
+ %2 = load float, ptr addrspace(4) %0, align 4
+ ret float %2
+}
+
+define double @f64(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .64b @f64(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.64b [[R64_D0]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R64_D0]];
+ %2 = load double, ptr addrspace(4) %0, align 8
+ ret double %2
+}
+
+define i32 @load_global_reg(ptr addrspace(1) %arg) {
+; CHECK-LABEL: .32b @load_global_reg(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.global.32b [[R32_W0]], [[[R64_D0]]];
+; CHECK-NEXT: return [[R32_W0]];
+ %1 = load i32, ptr addrspace(1) %arg, align 8
+ ret i32 %1
+}
+
+define i32 @load_shared_reg(ptr addrspace(3) %arg) {
+; CHECK-LABEL: .32b @load_shared_reg(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.shared.32b [[R32_W1]], [[[R32_W0]]];
+; CHECK-NEXT: return [[R32_W1]];
+ %1 = load i32, ptr addrspace(3) %arg, align 8
+ ret i32 %1
+}
+
+define i32 @load_global_reg_imm(ptr addrspace(1) %arg) {
+; CHECK-LABEL: .32b @load_global_reg_imm(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.global.32b [[R32_W0]], [[[R64_D0]] + 8];
+; CHECK-NEXT: return [[R32_W0]];
+ %1 = getelementptr [4 x i32], ptr addrspace(1) %arg, i64 0, i64 2
+ %2 = load i32, ptr addrspace(1) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_shared_reg_imm(ptr addrspace(3) %arg) {
+; CHECK-LABEL: .32b @load_shared_reg_imm(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.shared.32b [[R32_W1]], [[[R32_W0]] + 8];
+; CHECK-NEXT: return [[R32_W1]];
+ %1 = getelementptr [4 x i32], ptr addrspace(3) %arg, i32 0, i32 2
+ %2 = load i32, ptr addrspace(3) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_global_reg_reg(ptr addrspace(1) %arg, i64 %idx) {
+; CHECK-O0-LABEL: .32b @load_global_reg_reg(
+; CHECK-O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-O0-SAME: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]]
+; CHECK-O0: .reg .64b [[R64_D2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0: smul.64b [[R64_D2]], [[R64_D1]], 4;
+; CHECK-O0-NEXT: ld.global.32b [[R32_W0]], [[[R64_D0]] + [[R64_D2]]];
+; CHECK-O0-NEXT: return [[R32_W0]];
+;
+; CHECK-O2-LABEL: .32b @load_global_reg_reg(
+; CHECK-O2-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-O2-SAME: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]]
+; CHECK-O2: .reg .64b [[R64_D2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2: shl.64b [[R64_D2]], [[R64_D1]], 2;
+; CHECK-O2-NEXT: ld.global.32b [[R32_W0]], [[[R64_D0]] + [[R64_D2]]];
+; CHECK-O2-NEXT: return [[R32_W0]];
+ %1 = getelementptr [4 x i32], ptr addrspace(1) %arg, i64 0, i64 %idx
+ %2 = load i32, ptr addrspace(1) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_shared_reg_reg(ptr addrspace(3) %arg, i32 %idx) {
+; CHECK-O0-LABEL: .32b @load_shared_reg_reg(
+; CHECK-O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK-O0: .reg .32b [[R32_W2:%[-a-zA-Z$._0-9]+]], [[R32_W3:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0: smul.32b [[R32_W2]], [[R32_W1]], 4;
+; CHECK-O0-NEXT: ld.shared.32b [[R32_W3]], [[[R32_W0]] + [[R32_W2]]];
+; CHECK-O0-NEXT: return [[R32_W3]];
+;
+; CHECK-O2-LABEL: .32b @load_shared_reg_reg(
+; CHECK-O2-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-O2-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK-O2: .reg .32b [[R32_W2:%[-a-zA-Z$._0-9]+]], [[R32_W3:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2: shl.32b [[R32_W2]], [[R32_W1]], 2;
+; CHECK-O2-NEXT: ld.shared.32b [[R32_W3]], [[[R32_W0]] + [[R32_W2]]];
+; CHECK-O2-NEXT: return [[R32_W3]];
+ %1 = getelementptr [4 x i32], ptr addrspace(3) %arg, i32 0, i32 %idx
+ %2 = load i32, ptr addrspace(3) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_global_reg_immneg(ptr addrspace(1) %arg) {
+; CHECK-LABEL: .32b @load_global_reg_immneg(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.global.32b [[R32_W0]], [[[R64_D0]] - 8];
+; CHECK-NEXT: return [[R32_W0]];
+ %1 = getelementptr [4 x i32], ptr addrspace(1) %arg, i64 0, i64 -2
+ %2 = load i32, ptr addrspace(1) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_shared_reg_immneg(ptr addrspace(3) %arg) {
+; CHECK-LABEL: .32b @load_shared_reg_immneg(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.shared.32b [[R32_W1]], [[[R32_W0]] - 8];
+; CHECK-NEXT: return [[R32_W1]];
+ %1 = getelementptr [4 x i32], ptr addrspace(3) %arg, i32 0, i32 -2
+ %2 = load i32, ptr addrspace(3) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_global_reg_immneg_limit(ptr addrspace(1) %arg) {
+; CHECK-LABEL: .32b @load_global_reg_immneg_limit(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: mov.64b [[R64_D1]], -9223372036854775808;
+; CHECK-NEXT: ld.global.32b [[R32_W0]], [[[R64_D0]] + [[R64_D1]]];
+; CHECK-NEXT: return [[R32_W0]];
+ %1 = getelementptr i8, ptr addrspace(1) %arg, i64 -9223372036854775808
+ %2 = load i32, ptr addrspace(1) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_shared_reg_immneg_limit(ptr addrspace(3) %arg) {
+; CHECK-LABEL: .32b @load_shared_reg_immneg_limit(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.shared.32b [[R32_W1]], [[[R32_W0]] - 2147483648];
+; CHECK-NEXT: return [[R32_W1]];
+ %1 = getelementptr i8, ptr addrspace(3) %arg, i32 -2147483648
+ %2 = load i32, ptr addrspace(3) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_global_reg_immpos_limit(ptr addrspace(1) %arg) {
+; CHECK-LABEL: .32b @load_global_reg_immpos_limit(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: mov.64b [[R64_D1]], 9223372036854775807;
+; CHECK-NEXT: ld.global.32b [[R32_W0]], [[[R64_D0]] + [[R64_D1]]];
+; CHECK-NEXT: return [[R32_W0]];
+ %1 = getelementptr i8, ptr addrspace(1) %arg, i64 9223372036854775807
+ %2 = load i32, ptr addrspace(1) %1, align 4
+ ret i32 %2
+}
+
+define i32 @load_shared_reg_immpos_limit(ptr addrspace(3) %arg) {
+; CHECK-LABEL: .32b @load_shared_reg_immpos_limit(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.shared.32b [[R32_W1]], [[[R32_W0]] + 2147483647];
+; CHECK-NEXT: return [[R32_W1]];
+ %1 = getelementptr i8, ptr addrspace(3) %arg, i32 2147483647
+ %2 = load i32, ptr addrspace(3) %1, align 4
+ ret i32 %2
+}
+
+define void @load_extend() {
+; CHECK-O0-LABEL: void @load_extend(
+; CHECK-O0: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]], [[R32_W1:%[-a-zA-Z$._0-9]+]], [[R32_W2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0: mov.64b [[R64_D0]], 0;
+; CHECK-O0-NEXT: ld.global.32b [[R32_W0]], [[[R64_D0]]];
+; CHECK-O0-NEXT: shl.32b [[R32_W1]], [[R32_W0]], 16;
+; CHECK-O0-NEXT: asr.32b [[R32_W2]], [[R32_W1]], 16;
+; CHECK-O0-NEXT: st.global.32b [[[R64_D0]]], [[R32_W2]];
+; CHECK-O0-NEXT: return;
+;
+; CHECK-O2-LABEL: void @load_extend(
+; CHECK-O2: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2-NEXT: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2: mov.64b [[R64_D0]], 0;
+; CHECK-O2-NEXT: ld.global.16b [[R16_H0]], [[[R64_D0]]];
+; CHECK-O2-NEXT: sext.32b.16b [[R32_W0]], [[R16_H0]];
+; CHECK-O2-NEXT: st.global.32b [[[R64_D0]]], [[R32_W0]];
+; CHECK-O2-NEXT: return;
+entry:
+ %0 = load i32, ptr addrspace(1) null, align 4
+ %1 = shl i32 %0, 16
+ %conv9 = ashr i32 %1, 16
+ store i32 %conv9, ptr addrspace(1) null, align 4
+ ret void
+}
+
+define void @load_extendZ() {
+; CHECK-LABEL: void @load_extendZ(
+; CHECK: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]], [[R32_W1:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]];
+; CHECK: mov.64b [[R64_D0]], 0;
+; CHECK-NEXT: ld.global.32b [[R32_W0]], [[[R64_D0]]];
+; CHECK-NEXT: and.32b [[R32_W1]], [[R32_W0]], 65535;
+; CHECK-NEXT: trunc.16b.32b [[R16_H0]], [[R32_W1]];
+; CHECK-NEXT: st.global.16b [[[R64_D0]]], [[R16_H0]];
+; CHECK-NEXT: return;
+entry:
+ %0 = load i32, ptr addrspace(1) null, align 4
+ %1 = and i32 %0, u0xffff
+ %conv9 = trunc i32 %1 to i16
+ store i16 %conv9, ptr addrspace(1) null, align 4
+ ret void
+}
+
+define void @load_addr_with_swizzle(ptr addrspace(2) %data) {
+; CHECK-O0-LABEL: void @load_addr_with_swizzle(
+; CHECK-O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-O0: .reg .64b %d<1~7>;
+; CHECK-O0-NEXT: .reg .v2.64b [[VR2_64_V2D0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0-NEXT: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]], [[R8_B1:%[-a-zA-Z$._0-9]+]], [[R8_B2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0-NEXT: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]], [[R16_H1:%[-a-zA-Z$._0-9]+]], [[R16_H2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O0: mov.64b %d1, 0;
+; CHECK-O0-NEXT: mov.64b %d2, 0;
+; CHECK-O0-NEXT: ld.const.v2.64b [[VR2_64_V2D0]], [%d1];
+; CHECK-O0-NEXT: mov.64b %d3, [[VR2_64_V2D0]].x;
+; CHECK-O0-NEXT: mov.64b %d4, [[VR2_64_V2D0]].y;
+; CHECK-O0-NEXT: mov.64b %d5, [[VR2_64_V2D0]].x;
+; CHECK-O0-NEXT: mov.64b %d6, [[VR2_64_V2D0]].y;
+; CHECK-O0-NEXT: ld.const.8b [[R8_B0]], [[[R64_D0]] + %d3];
+; CHECK-O0-NEXT: ld.const.8b [[R8_B1]], [[[R64_D0]] + %d6];
+; CHECK-O0-NEXT: zext.16b.8b [[R16_H0]], [[R8_B1]];
+; CHECK-O0-NEXT: zext.16b.8b [[R16_H1]], [[R8_B0]];
+; CHECK-O0-NEXT: iadd.16b [[R16_H2]], [[R16_H0]], [[R16_H1]];
+; CHECK-O0-NEXT: trunc.8b.16b [[R8_B2]], [[R16_H2]];
+; CHECK-O0-NEXT: st.global.8b [%d2], [[R8_B2]];
+; CHECK-O0-NEXT: return;
+;
+; CHECK-O2-LABEL: void @load_addr_with_swizzle(
+; CHECK-O2-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-O2: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2-NEXT: .reg .v2.64b [[VR2_64_V2D0:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2-NEXT: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]], [[R8_B1:%[-a-zA-Z$._0-9]+]], [[R8_B2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2-NEXT: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]], [[R16_H1:%[-a-zA-Z$._0-9]+]], [[R16_H2:%[-a-zA-Z$._0-9]+]];
+; CHECK-O2: mov.64b [[R64_D1]], 0;
+; CHECK-O2-NEXT: ld.const.v2.64b [[VR2_64_V2D0]], [[[R64_D1]]];
+; CHECK-O2-NEXT: ld.const.8b [[R8_B0]], [[[R64_D0]] + [[VR2_64_V2D0]].x];
+; CHECK-O2-NEXT: ld.const.8b [[R8_B1]], [[[R64_D0]] + [[VR2_64_V2D0]].y];
+; CHECK-O2-NEXT: zext.16b.8b [[R16_H0]], [[R8_B1]];
+; CHECK-O2-NEXT: zext.16b.8b [[R16_H1]], [[R8_B0]];
+; CHECK-O2-NEXT: iadd.16b [[R16_H2]], [[R16_H0]], [[R16_H1]];
+; CHECK-O2-NEXT: trunc.8b.16b [[R8_B2]], [[R16_H2]];
+; CHECK-O2-NEXT: st.global.8b [[[R64_D1]]], [[R8_B2]];
+; CHECK-O2-NEXT: return;
+entry:
+ %0 = load <2 x i64>, ptr addrspace(2) null
+ %1 = extractelement <2 x i64> %0, i32 0
+ %2 = extractelement <2 x i64> %0, i32 1
+ %arrayidx2 = getelementptr inbounds i8, ptr addrspace(2) %data, i64 %1
+ %3 = load i8, ptr addrspace(2) %arrayidx2
+ %arrayidx3 = getelementptr inbounds i8, ptr addrspace(2) %data, i64 %2
+ %4 = load i8, ptr addrspace(2) %arrayidx3
+ %add = add i8 %4, %3
+ store i8 %add, ptr addrspace(1) null
+ ret void
+}
+
+define <4 x float> @load_redundant_offset(ptr addrspace(4) %addr, i32 %i) {
+; CHECK-LABEL: .v4.32b @load_redundant_offset(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W2:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .v4.32b [[VR4_32_V4W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: iadd.32b [[R32_W2]], [[R32_W0]], [[R32_W1]];
+; CHECK-NEXT: ld.private.v4.32b [[VR4_32_V4W0]], [[[R32_W2]] + 16];
+; CHECK-NEXT: return [[VR4_32_V4W0]];
+entry:
+ %gep0 = getelementptr inbounds i8, ptr addrspace(4) %addr, i32 %i
+ %gep1 = getelementptr inbounds i8, ptr addrspace(4) %gep0, i32 16
+ %fvec = load <4 x float>, ptr addrspace(4) %gep1, align 4
+ ret <4 x float> %fvec
+}
+
+define i8 @i4(ptr addrspace(4) noundef %0) {
+; CHECK-LABEL: .8b @i4(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]], [[R8_B1:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]], [[R16_H1:%[-a-zA-Z$._0-9]+]];
+; CHECK: ld.private.8b [[R8_B0]], [[[R32_W0]]];
+; CHECK-NEXT: zext.16b.8b [[R16_H0]], [[R8_B0]];
+; CHECK-NEXT: and.16b [[R16_H1]], [[R16_H0]], 15;
+; CHECK-NEXT: trunc.8b.16b [[R8_B1]], [[R16_H1]];
+; CHECK-NEXT: return [[R8_B1]];
+ %v = load i4, ptr addrspace(4) %0, align 1
+ %r = zext i4 %v to i8
+ ret i8 %r
+}
diff --git a/llvm/test/CodeGen/PISA/return.ll b/llvm/test/CodeGen/PISA/return.ll
new file mode 100644
index 0000000000000..258143595bbf9
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/return.ll
@@ -0,0 +1,37 @@
+; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
+
+define i64 @return_value(i64 %a, i64 %b) {
+; O0-LABEL: .64b @return_value(
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]]
+; O0: .reg .64b [[R64_D2:%[-a-zA-Z$._0-9]+]];
+; O0: iadd.64b [[R64_D2]], [[R64_D0]], [[R64_D1]];
+; O0-NEXT: return [[R64_D2]];
+;
+; CHECK-LABEL: .64b @return_value(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D2:%[-a-zA-Z$._0-9]+]];
+; CHECK: iadd.64b [[R64_D2]], [[R64_D0]], [[R64_D1]];
+; CHECK-NEXT: return [[R64_D2]];
+ %result = add i64 %a, %b
+ ret i64 %result
+}
+
+define void @return_void(ptr addrspace(4) noundef %0, i8 %1) {
+; O0-LABEL: void @return_void(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.8b [[[R32_W0]]], [[R8_B0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @return_void(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.8b [[[R32_W0]]], [[R8_B0]];
+; CHECK-NEXT: return;
+ store i8 %1, ptr addrspace(4) %0, align 1
+ ret void
+}
diff --git a/llvm/test/CodeGen/PISA/store.ll b/llvm/test/CodeGen/PISA/store.ll
new file mode 100644
index 0000000000000..687ebdab1aba5
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/store.ll
@@ -0,0 +1,352 @@
+; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
+
+define void @i8(ptr addrspace(4) noundef %0, i8 %1) {
+; O0-LABEL: void @i8(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.8b [[[R32_W0]]], [[R8_B0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @i8(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.8b [[[R32_W0]]], [[R8_B0]];
+; CHECK-NEXT: return;
+ store i8 %1, ptr addrspace(4) %0, align 1
+ ret void
+}
+
+define void @i16(ptr addrspace(4) noundef %0, i16 %1) {
+; O0-LABEL: void @i16(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.16b [[[R32_W0]]], [[R16_H0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @i16(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.16b [[[R32_W0]]], [[R16_H0]];
+; CHECK-NEXT: return;
+ store i16 %1, ptr addrspace(4) %0, align 2
+ ret void
+}
+
+define void @i32(ptr addrspace(4) noundef %0, i32 %1) {
+; O0-LABEL: void @i32(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.32b [[[R32_W0]]], [[R32_W1]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @i32(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.32b [[[R32_W0]]], [[R32_W1]];
+; CHECK-NEXT: return;
+ store i32 %1, ptr addrspace(4) %0, align 4
+ ret void
+}
+
+define void @i64(ptr addrspace(4) noundef %0, i64 %1) {
+; O0-LABEL: void @i64(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.64b [[[R32_W0]]], [[R64_D0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @i64(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.64b [[[R32_W0]]], [[R64_D0]];
+; CHECK-NEXT: return;
+ store i64 %1, ptr addrspace(4) %0, align 8
+ ret void
+}
+
+define void @f16(ptr addrspace(4) noundef %0, half %1) {
+; O0-LABEL: void @f16(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.16b [[[R32_W0]]], [[R16_H0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @f16(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .16b [[R16_H0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.16b [[[R32_W0]]], [[R16_H0]];
+; CHECK-NEXT: return;
+ store half %1, ptr addrspace(4) %0, align 4
+ ret void
+}
+
+define void @f32(ptr addrspace(4) noundef %0, float %1) {
+; O0-LABEL: void @f32(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.32b [[[R32_W0]]], [[R32_W1]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @f32(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.32b [[[R32_W0]]], [[R32_W1]];
+; CHECK-NEXT: return;
+ store float %1, ptr addrspace(4) %0, align 4
+ ret void
+}
+
+define void @f64(ptr addrspace(4) noundef %0, double %1) {
+; O0-LABEL: void @f64(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.64b [[[R32_W0]]], [[R64_D0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @f64(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.64b [[[R32_W0]]], [[R64_D0]];
+; CHECK-NEXT: return;
+ store double %1, ptr addrspace(4) %0, align 8
+ ret void
+}
+
+define void @store_global_reg_imm(ptr addrspace(1) %arg, i32 %data) {
+; O0-LABEL: void @store_global_reg_imm(
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0: st.global.32b [[[R64_D0]] + 8], [[R32_W0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_global_reg_imm(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.global.32b [[[R64_D0]] + 8], [[R32_W0]];
+; CHECK-NEXT: return;
+ %1 = getelementptr [4 x i32], ptr addrspace(1) %arg, i64 0, i64 2
+ store i32 %data, ptr addrspace(1) %1, align 4
+ ret void
+}
+
+define void @store_shared_reg_imm(ptr addrspace(3) %arg, i32 %data) {
+; O0-LABEL: void @store_shared_reg_imm(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0: st.shared.32b [[[R32_W0]] + 8], [[R32_W1]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_shared_reg_imm(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.shared.32b [[[R32_W0]] + 8], [[R32_W1]];
+; CHECK-NEXT: return;
+ %1 = getelementptr [4 x i32], ptr addrspace(3) %arg, i32 0, i32 2
+ store i32 %data, ptr addrspace(3) %1, align 4
+ ret void
+}
+
+define void @store_global_reg_reg(ptr addrspace(1) %arg, i64 %idx, i32 %data) {
+; O0-LABEL: void @store_global_reg_reg(
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0: .reg .64b [[R64_D2:%[-a-zA-Z$._0-9]+]];
+; O0: smul.64b [[R64_D2]], [[R64_D1]], 4;
+; O0-NEXT: st.global.32b [[[R64_D0]] + [[R64_D2]]], [[R32_W0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_global_reg_reg(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D2:%[-a-zA-Z$._0-9]+]];
+; CHECK: shl.64b [[R64_D2]], [[R64_D1]], 2;
+; CHECK-NEXT: st.global.32b [[[R64_D0]] + [[R64_D2]]], [[R32_W0]];
+; CHECK-NEXT: return;
+ %1 = getelementptr [4 x i32], ptr addrspace(1) %arg, i64 0, i64 %idx
+ store i32 %data, ptr addrspace(1) %1, align 4
+ ret void
+}
+
+define void @store_shared_reg_reg(ptr addrspace(3) %arg, i32 %idx, i32 %data) {
+; O0-LABEL: void @store_shared_reg_reg(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W2:%[-a-zA-Z$._0-9]+]]
+; O0: .reg .32b [[R32_W3:%[-a-zA-Z$._0-9]+]];
+; O0: smul.32b [[R32_W3]], [[R32_W1]], 4;
+; O0-NEXT: st.shared.32b [[[R32_W0]] + [[R32_W3]]], [[R32_W2]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_shared_reg_reg(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W2:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .32b [[R32_W3:%[-a-zA-Z$._0-9]+]];
+; CHECK: shl.32b [[R32_W3]], [[R32_W1]], 2;
+; CHECK-NEXT: st.shared.32b [[[R32_W0]] + [[R32_W3]]], [[R32_W2]];
+; CHECK-NEXT: return;
+ %1 = getelementptr [4 x i32], ptr addrspace(3) %arg, i32 0, i32 %idx
+ store i32 %data, ptr addrspace(3) %1, align 4
+ ret void
+}
+
+define void @store_global_reg_immneg(ptr addrspace(1) %arg, i32 %data) {
+; O0-LABEL: void @store_global_reg_immneg(
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0: st.global.32b [[[R64_D0]] - 8], [[R32_W0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_global_reg_immneg(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.global.32b [[[R64_D0]] - 8], [[R32_W0]];
+; CHECK-NEXT: return;
+ %1 = getelementptr [4 x i32], ptr addrspace(1) %arg, i64 0, i64 -2
+ store i32 %data, ptr addrspace(1) %1, align 4
+ ret void
+}
+
+define void @store_shared_reg_immneg(ptr addrspace(3) %arg, i32 %data) {
+; O0-LABEL: void @store_shared_reg_immneg(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0: st.shared.32b [[[R32_W0]] - 8], [[R32_W1]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_shared_reg_immneg(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.shared.32b [[[R32_W0]] - 8], [[R32_W1]];
+; CHECK-NEXT: return;
+ %1 = getelementptr [4 x i32], ptr addrspace(3) %arg, i32 0, i32 -2
+ store i32 %data, ptr addrspace(3) %1, align 4
+ ret void
+}
+
+define void @store_global_reg_immneg_limit(ptr addrspace(1) %arg, i32 %data) {
+; O0-LABEL: void @store_global_reg_immneg_limit(
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; O0: mov.64b [[R64_D1]], -9223372036854775808;
+; O0-NEXT: st.global.32b [[[R64_D0]] + [[R64_D1]]], [[R32_W0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_global_reg_immneg_limit(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; CHECK: mov.64b [[R64_D1]], -9223372036854775808;
+; CHECK: st.global.32b [[[R64_D0]] + [[R64_D1]]], [[R32_W0]];
+; CHECK-NEXT: return;
+ %1 = getelementptr i8, ptr addrspace(1) %arg, i64 -9223372036854775808
+ store i32 %data, ptr addrspace(1) %1, align 4
+ ret void
+}
+
+define void @store_shared_reg_immneg_limit(ptr addrspace(3) %arg, i32 %data) {
+; O0-LABEL: void @store_shared_reg_immneg_limit(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0: st.shared.32b [[[R32_W0]] - 2147483648], [[R32_W1]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_shared_reg_immneg_limit(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.shared.32b [[[R32_W0]] - 2147483648], [[R32_W1]];
+; CHECK-NEXT: return;
+ %1 = getelementptr i8, ptr addrspace(3) %arg, i32 -2147483648
+ store i32 %data, ptr addrspace(3) %1, align 4
+ ret void
+}
+
+define void @store_global_reg_immpos_limit(ptr addrspace(1) %arg, i32 %data) {
+; O0-LABEL: void @store_global_reg_immpos_limit(
+; O0-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; O0: mov.64b [[R64_D1]], 9223372036854775807;
+; O0-NEXT: st.global.32b [[[R64_D0]] + [[R64_D1]]], [[R32_W0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_global_reg_immpos_limit(
+; CHECK-SAME: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .64b [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; CHECK: mov.64b [[R64_D1]], 9223372036854775807;
+; CHECK: st.global.32b [[[R64_D0]] + [[R64_D1]]], [[R32_W0]];
+; CHECK-NEXT: return;
+ %1 = getelementptr i8, ptr addrspace(1) %arg, i64 9223372036854775807
+ store i32 %data, ptr addrspace(1) %1, align 4
+ ret void
+}
+
+define void @store_shared_reg_immpos_limit(ptr addrspace(3) %arg, i32 %data) {
+; O0-LABEL: void @store_shared_reg_immpos_limit(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; O0: st.shared.32b [[[R32_W0]] + 2147483647], [[R32_W1]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @store_shared_reg_immpos_limit(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .32b [[R32_W1:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.shared.32b [[[R32_W0]] + 2147483647], [[R32_W1]];
+; CHECK-NEXT: return;
+ %1 = getelementptr i8, ptr addrspace(3) %arg, i32 2147483647
+ store i32 %data, ptr addrspace(3) %1, align 4
+ ret void
+}
+
+define void @i128(ptr addrspace(4) %0) {
+; O0-LABEL: void @i128(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0: .reg .64b [[R64_D0:%[-a-zA-Z$._0-9]+]], [[R64_D1:%[-a-zA-Z$._0-9]+]];
+; O0-NEXT: .reg .v2.64b [[VR2_64_V2D0:%[-a-zA-Z$._0-9]+]], [[VR2_64_V2D1:%[-a-zA-Z$._0-9]+]];
+; O0-NEXT: .reg .128b [[R128_Q0:%[-a-zA-Z$._0-9]+]];
+; O0-NEXT: .reg .v4.32b [[VR4_32_V4W0:%[-a-zA-Z$._0-9]+]];
+; O0: mov.64b [[R64_D0]], 0;
+; O0-NEXT: mov.64b [[R64_D1]], 0;
+; O0-NEXT: mov.64b [[VR2_64_V2D0]].x, [[R64_D0]];
+; O0-NEXT: mov.128b [[VR2_64_V2D1]].xy, [[VR2_64_V2D0]].xy;
+; O0-NEXT: mov.64b [[VR2_64_V2D1]].y, [[R64_D1]];
+; O0-NEXT: mov.128b [[R128_Q0]], [[VR2_64_V2D1]].xy;
+; O0-NEXT: mov.128b [[VR4_32_V4W0]].xyzw, [[R128_Q0]];
+; O0-NEXT: st.private.v4.32b [[[R32_W0]]], [[VR4_32_V4W0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @i128(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK: .reg .v2.64b [[VR2_64_V2D0:%[-a-zA-Z$._0-9]+]];
+; CHECK-NEXT: .reg .v4.32b [[VR4_32_V4W0:%[-a-zA-Z$._0-9]+]];
+; CHECK: mov.64b [[VR2_64_V2D0]].x, 0;
+; CHECK-NEXT: mov.64b [[VR2_64_V2D0]].y, [[VR2_64_V2D0]].x;
+; CHECK-NEXT: mov.128b [[VR4_32_V4W0]].xyzw, [[VR2_64_V2D0]].xy;
+; CHECK-NEXT: st.private.v4.32b [[[R32_W0]]], [[VR4_32_V4W0]];
+; CHECK-NEXT: return;
+ store i128 0, ptr addrspace(4) %0, align 64
+ ret void
+}
+
+define void @i4(ptr addrspace(4) noundef %0, i8 %1) {
+; O0-LABEL: void @i4(
+; O0-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]]
+; O0: st.private.8b [[[R32_W0]]], [[R8_B0]];
+; O0-NEXT: return;
+;
+; CHECK-LABEL: void @i4(
+; CHECK-SAME: .reg .32b [[R32_W0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .reg .8b [[R8_B0:%[-a-zA-Z$._0-9]+]]
+; CHECK: st.private.8b [[[R32_W0]]], [[R8_B0]];
+; CHECK-NEXT: return;
+ %v = trunc i8 %1 to i4
+ store i4 %v, ptr addrspace(4) %0, align 1
+ ret void
+}
diff --git a/llvm/test/CodeGen/PISA/unreachable.ll b/llvm/test/CodeGen/PISA/unreachable.ll
new file mode 100644
index 0000000000000..0e9fb95e4de91
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/unreachable.ll
@@ -0,0 +1,14 @@
+; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
+
+define pisa_kernel void @test_unreachable(i32 %a, i32 %b) {
+; O0-LABEL: @test_unreachable(
+; O0-SAME: .param[4] .align(4) [[PARAM_ARG0:%[-a-zA-Z$._0-9]+]]
+; O0-SAME: .param[4] .align(4) [[PARAM_ARG1:%[-a-zA-Z$._0-9]+]]
+;
+; CHECK-LABEL: @test_unreachable(
+; CHECK-SAME: .param[4] .align(4) [[PARAM_ARG0:%[-a-zA-Z$._0-9]+]]
+; CHECK-SAME: .param[4] .align(4) [[PARAM_ARG1:%[-a-zA-Z$._0-9]+]]
+ unreachable
+}
>From 7af293ea986b0188edc51e5e0086d95e3a7478df Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Thu, 6 Aug 2026 23:22:44 -0700
Subject: [PATCH 02/11] Formatting fix
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index 512d1d2efafc3..bcae661568b8a 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -928,6 +928,7 @@ void PISAAsmPrinter::emitInstruction(const MachineInstr *MI) {
// Force static initialization.
// NOLINTNEXTLINE(readability-identifier-naming)
-extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void LLVMInitializePISAAsmPrinter() {
+extern "C" LLVM_ABI LLVM_EXTERNAL_VISIBILITY void
+LLVMInitializePISAAsmPrinter() {
RegisterAsmPrinter<PISAAsmPrinter> Y(getThePISATarget());
}
>From 12cc3fbeda0a01128ccc2ebbc2b068778fa89936 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:02:29 +0000
Subject: [PATCH 03/11] Fix issues raised in earlier PRs
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 206 ++++++++++--------
.../asm-printer-unsupported-vector-hint.ll | 11 +
llvm/test/CodeGen/PISA/load.ll | 6 +-
llvm/test/CodeGen/PISA/return.ll | 4 +-
llvm/test/CodeGen/PISA/store.ll | 4 +-
llvm/test/CodeGen/PISA/unreachable.ll | 4 +-
6 files changed, 132 insertions(+), 103 deletions(-)
create mode 100644 llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint.ll
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index bcae661568b8a..6f4ec4267923b 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -35,6 +35,7 @@
#include "llvm/IR/IRPrintingPasses.h"
#include "llvm/IR/Module.h"
#include "llvm/MC/MCAsmInfo.h"
+#include "llvm/MC/MCExpr.h"
#include "llvm/MC/MCInst.h"
#include "llvm/MC/MCStreamer.h"
#include "llvm/MC/MCSymbol.h"
@@ -142,11 +143,22 @@ class PISAAsmPrinter : public AsmPrinter {
uint64_t Idx = VI.Initializer.size() - 1;
VI.Exprs.insert({Idx, GE});
}
+ void addSplatVector(const Constant *C, LLT ScalarTy,
+ const APInt &EltVal) {
+ assert(EltVal.getBitWidth() <= 64 &&
+ "Splat element too wide for uint64_t");
+ FixedVectorType *VTy = cast<FixedVectorType>(C->getType());
+ unsigned NumElts = VTy->getNumElements();
+ uint64_t Val = EltVal.getZExtValue();
+ for (unsigned I = 0; I < NumElts; ++I)
+ addVal(ScalarTy, Val);
+ pad(C, NumElts);
+ }
void lowerConstant(const Constant *C) {
- auto *Expr = AP.lowerConstant(C);
+ const MCExpr *Expr = AP.lowerConstant(C);
MCValue Res;
if (!Expr->evaluateAsRelocatable(Res, nullptr))
- llvm_unreachable("unhandled expression!");
+ reportFatalUsageError("unsupported PISA global initializer expression");
LLT Ty = getLLTForType(*C->getType(), DL);
if (!Res.getAddSym() && !Res.getSubSym()) {
if (Res.getConstant() == 0)
@@ -155,22 +167,26 @@ class PISAAsmPrinter : public AsmPrinter {
addVal(Ty, static_cast<uint64_t>(Res.getConstant()));
return;
}
- assert(!Res.getSubSym() && "unhandled expression!");
+ if (Res.getSubSym())
+ reportFatalUsageError("unsupported PISA global initializer expression");
std::string Name = Res.getAddSym()->getName().str();
PISA::VariableInit::GlobalExpr E{std::move(Name), Res.getConstant()};
addGlobal(Ty, E);
}
uint64_t computeSize(const Constant *C) const {
uint64_t Total = 0;
- for (auto [i, Elt] : llvm::enumerate(VI.Initializer)) {
- if (auto Iter = VI.Exprs.find(i); Iter != VI.Exprs.end()) {
- auto &Entry = Iter->second;
- if (auto *Z = std::get_if<PISA::VariableInit::Zeros>(&Entry)) {
+ for (uint64_t I = 0; I < VI.Initializer.size(); ++I) {
+ DenseMap<uint64_t, PISA::VariableInit::SpecialEntry>::const_iterator
+ Iter = VI.Exprs.find(I);
+ if (Iter != VI.Exprs.end()) {
+ const PISA::VariableInit::SpecialEntry &Entry = Iter->second;
+ if (const PISA::VariableInit::Zeros *Z =
+ std::get_if<PISA::VariableInit::Zeros>(&Entry)) {
Total += Z->N;
continue;
}
}
- Total += Elt.Type.getSizeInBytes();
+ Total += VI.Initializer[I].Type.getSizeInBytes();
}
return Total;
}
@@ -248,33 +264,23 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
uint64_t Size = DL.getTypeAllocSize(C->getType());
if (isZero(C))
return pad(C);
- auto AddSplatVector = [&](LLT ScalarTy, const APInt &EltVal) {
- assert(EltVal.getBitWidth() <= 64 && "Splat element too wide for uint64_t");
- auto *VTy = cast<FixedVectorType>(C->getType());
- unsigned NumElts = VTy->getNumElements();
- uint64_t Val = EltVal.getZExtValue();
- for (unsigned I = 0; I < NumElts; ++I)
- addVal(ScalarTy, Val);
- pad(C, NumElts);
- };
-
- if (auto *CI = dyn_cast<ConstantInt>(C)) {
+ if (const ConstantInt *CI = dyn_cast<ConstantInt>(C)) {
if (C->getType()->isVectorTy()) {
- auto *VTy = cast<FixedVectorType>(C->getType());
+ FixedVectorType *VTy = cast<FixedVectorType>(C->getType());
uint64_t EltAllocSize = DL.getTypeAllocSize(VTy->getElementType());
LLT ScalarTy = LLT::integer(EltAllocSize * 8);
- AddSplatVector(ScalarTy, CI->getValue());
+ addSplatVector(C, ScalarTy, CI->getValue());
} else {
// We don't use the LLT type of `C` directly here because `C` could be,
// for example, a s1. The allocation size is 1, so we want to give it
// a type of s8 to reflect that.
addVal(LLT::integer(Size * 8), CI->getZExtValue());
}
- } else if (auto *FP = dyn_cast<ConstantFP>(C)) {
+ } else if (const ConstantFP *FP = dyn_cast<ConstantFP>(C)) {
if (C->getType()->isVectorTy()) {
- auto *VTy = cast<FixedVectorType>(C->getType());
+ FixedVectorType *VTy = cast<FixedVectorType>(C->getType());
LLT ScalarTy = getLLTForType(*VTy->getElementType(), DL);
- AddSplatVector(ScalarTy, FP->getValueAPF().bitcastToAPInt());
+ addSplatVector(C, ScalarTy, FP->getValueAPF().bitcastToAPInt());
} else {
LLT Ty = getLLTForType(*C->getType(), DL);
addVal(Ty, FP->getValueAPF().bitcastToAPInt().getZExtValue());
@@ -286,7 +292,7 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
addVal(Ty, Val);
else
pad(C);
- } else if (auto *CV = dyn_cast<ConstantVector>(C)) {
+ } else if (const ConstantVector *CV = dyn_cast<ConstantVector>(C)) {
Type *ElementType = CV->getType()->getElementType();
uint64_t ElementSizeInBits = DL.getTypeSizeInBits(ElementType);
uint64_t ElementAllocSizeInBits = DL.getTypeAllocSizeInBits(ElementType);
@@ -302,7 +308,7 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
ConstantExpr::getBitCast(const_cast<ConstantVector *>(CV), IntT),
DL));
if (!CI) {
- report_fatal_error(
+ reportFatalUsageError(
"Cannot lower vector global with unusual element type");
}
emitGlobalConstantLargeInt(CI);
@@ -314,12 +320,12 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
process(CV->getAggregateElement(I));
pad(C, CV->getNumOperands());
}
- } else if (auto *CA = dyn_cast<ConstantArray>(C)) {
+ } else if (const ConstantArray *CA = dyn_cast<ConstantArray>(C)) {
for (unsigned I = 0; I < CA->getNumOperands(); I++)
process(CA->getAggregateElement(I));
- } else if (auto *CS = dyn_cast<ConstantStruct>(C)) {
- auto *StructTy = cast<StructType>(CS->getType());
- auto *Layout = DL.getStructLayout(StructTy);
+ } else if (const ConstantStruct *CS = dyn_cast<ConstantStruct>(C)) {
+ StructType *StructTy = cast<StructType>(CS->getType());
+ const StructLayout *Layout = DL.getStructLayout(StructTy);
for (unsigned I = 0, E = CS->getNumOperands(); I != E; ++I) {
const Constant *Field = CS->getOperand(I);
// Print the actual field value.
@@ -335,7 +341,8 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
// as padding to ensure that the next field starts at the right offset.
pad(PadSize);
}
- } else if (auto *CDS = dyn_cast<ConstantDataSequential>(C)) {
+ } else if (const ConstantDataSequential *CDS =
+ dyn_cast<ConstantDataSequential>(C)) {
for (unsigned I = 0; I < CDS->getNumElements(); I++)
process(CDS->getElementAsConstant(I));
pad(C, CDS->getNumElements());
@@ -357,7 +364,7 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
assert(Size == 8 && "global symbol with non 64-bit size?");
lowerConstant(C);
} else {
- llvm_unreachable("unhandled constant!");
+ reportFatalUsageError("unsupported PISA global initializer constant");
}
}
@@ -373,14 +380,14 @@ static bool isIgnoredIntrinsicGlobal(const GlobalVariable &GV) {
// These are metadata for the annotation intrinsic, not real data.
if (GV.hasPrivateLinkage() && GV.isConstant() &&
all_of(GV.users(), [](const User *U) {
- if (auto *CE = dyn_cast<ConstantExpr>(U))
+ if (const ConstantExpr *CE = dyn_cast<ConstantExpr>(U))
return all_of(CE->users(), [](const User *UU) {
- auto *CI = dyn_cast<CallInst>(UU);
+ const CallInst *CI = dyn_cast<CallInst>(UU);
return CI && CI->getCalledFunction() &&
CI->getCalledFunction()->getIntrinsicID() ==
Intrinsic::ptr_annotation;
});
- auto *CI = dyn_cast<CallInst>(U);
+ const CallInst *CI = dyn_cast<CallInst>(U);
return CI && CI->getCalledFunction() &&
CI->getCalledFunction()->getIntrinsicID() ==
Intrinsic::ptr_annotation;
@@ -391,14 +398,14 @@ static bool isIgnoredIntrinsicGlobal(const GlobalVariable &GV) {
return false;
if (GV.getName() == "llvm.global_ctors")
- report_fatal_error(
+ reportFatalUsageError(
"llvm.global_ctors is not supported by the PISA backend");
if (GV.getName() == "llvm.global_dtors")
- report_fatal_error(
- "llvm.global_ctors is not supported by the PISA backend");
+ reportFatalUsageError(
+ "llvm.global_dtors is not supported by the PISA backend");
- report_fatal_error("unknown special variable with appending linkage");
+ reportFatalUsageError("unknown special variable with appending linkage");
}
void PISAAsmPrinter::emitGlobalsAndFuncDecls(Module &M) {
@@ -406,15 +413,15 @@ void PISAAsmPrinter::emitGlobalsAndFuncDecls(Module &M) {
// emit header info
// - we always emit in latest PISA syntax
- auto GetHdrTarget = [&]() -> SmallString<16> {
- return ST ? ST->getPISATargetName() : "";
- };
- PISA::HeaderDcl HD = {PISA::LatestPISAVersion, GetHdrTarget()};
+ SmallString<16> HeaderTarget;
+ if (ST)
+ HeaderTarget = ST->getPISATargetName();
+ PISA::HeaderDcl HD = {PISA::LatestPISAVersion, HeaderTarget};
TS.emitHeader(HD);
OutStreamer->addBlankLine();
// Emit Module level function decl
- for (auto &F : M) {
+ for (Function &F : M) {
if (!F.isDeclaration() || F.isIntrinsic()) // avoid llvm builtins
continue;
@@ -426,7 +433,7 @@ void PISAAsmPrinter::emitGlobalsAndFuncDecls(Module &M) {
}
// Translate global variables
- for (auto &GV : M.globals()) {
+ for (GlobalVariable &GV : M.globals()) {
if (isIgnoredIntrinsicGlobal(GV))
continue;
@@ -470,7 +477,7 @@ bool PISAAsmPrinter::runOnMachineFunction(MachineFunction &MF) {
void PISAAsmPrinter::emitFunctionHeader() {
const Function &F = MF->getFunction();
- auto *Section = getObjFileLowering().SectionForGlobal(&F, TM);
+ MCSection *Section = getObjFileLowering().SectionForGlobal(&F, TM);
MF->setSection(Section);
}
@@ -479,22 +486,27 @@ void PISAAsmPrinter::updateFuncParamIdxs(PISA::DataTypes &DTs) {
// declarations (vs the already-processed func param dcls)
DTs.finalizeFuncParams();
- llvm::DenseMap<std::tuple</*NumElts=*/unsigned, /*BitWidth=*/unsigned,
- /*Type=*/unsigned>,
- /*Index=*/unsigned>
- ParamIdxs;
-
- auto &MRI = MF->getRegInfo();
- for (auto &[CurReg, Info] : RegMgr->mapping()) {
+ using ParamIndexMap =
+ DenseMap<std::tuple</*NumElts=*/unsigned, /*BitWidth=*/unsigned,
+ /*Type=*/unsigned>,
+ /*Index=*/unsigned>;
+ ParamIndexMap ParamIdxs;
+
+ MachineRegisterInfo &MRI = MF->getRegInfo();
+ for (const PISA::RegManager::MappingTy::value_type &Entry :
+ RegMgr->mapping()) {
+ Register CurReg = Entry.first;
+ const PISA::RegManager::RegInfo &Info = Entry.second;
// We are only trying to update indices for function parameters
if (!(Info.Flags & PISA::RegManager::NoEmissionDef))
continue;
- auto *RC = MRI.getRegClass(CurReg);
+ const TargetRegisterClass *RC = MRI.getRegClass(CurReg);
unsigned BitWidth = TRI->getBitSizeFromRegClass(RC);
unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
- auto [It, Inserted] =
- ParamIdxs.try_emplace(std::make_tuple(NumElts, BitWidth, Info.Type), 0);
+ ParamIndexMap::iterator It =
+ ParamIdxs.try_emplace(std::make_tuple(NumElts, BitWidth, Info.Type), 0)
+ .first;
RegMgr->setRegIdx(CurReg, It->second++);
// Sanity check that all function parameter indexes are < the total
@@ -506,26 +518,30 @@ void PISAAsmPrinter::updateFuncParamIdxs(PISA::DataTypes &DTs) {
}
void PISAAsmPrinter::collectRegDcls(PISA::RegDcls &Dcls) {
- auto &MRI = MF->getRegInfo();
- for (auto &[CurReg, Info] : RegMgr->mapping()) {
+ MachineRegisterInfo &MRI = MF->getRegInfo();
+ for (const PISA::RegManager::MappingTy::value_type &Entry :
+ RegMgr->mapping()) {
+ Register CurReg = Entry.first;
+ const PISA::RegManager::RegInfo &Info = Entry.second;
if (Info.Flags & PISA::RegManager::NoEmissionDef)
continue;
- auto *RC = MRI.getRegClass(CurReg);
+ const TargetRegisterClass *RC = MRI.getRegClass(CurReg);
unsigned BitWidth = TRI->getBitSizeFromRegClass(RC);
unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
TypeInfo &TI = DTs->emplaceInfo(NumElts, BitWidth, Info.Type);
- auto Bank = RegMgr->getRegBank(NumElts, BitWidth);
+ PISA::RegEncoder::RegBank Bank =
+ RegMgr->getRegBank(NumElts, BitWidth);
const char *Prefix =
RegMgr->getPrefixFromBank(static_cast<PISA::RegManager::RegBank>(Bank));
Dcls.Regs[std::make_tuple(NumElts, BitWidth, Info.Type)].push_back(
- std::make_pair(Prefix, TI.RegCounter));
+ {Prefix, TI.RegCounter});
RegMgr->setRegIdx(CurReg, TI.RegCounter);
TI.RegCounter++;
}
}
void PISAAsmPrinter::collectLocalVariableDcls(PISA::LocalVariableDcls &Dcls) {
- auto &MFI = MF->getFrameInfo();
+ MachineFrameInfo &MFI = MF->getFrameInfo();
for (int Idx = MFI.getObjectIndexBegin(), EndIdx = MFI.getObjectIndexEnd();
Idx != EndIdx; ++Idx) {
if (MFI.isDeadObjectIndex(Idx))
@@ -560,7 +576,7 @@ PISA::LinkageTy PISAAsmPrinter::collectLinkage(const GlobalValue &V) {
return PISA::LinkageTy::DEFAULT;
// global variable linkage
- if (auto *GVar = dyn_cast<GlobalVariable>(&V)) {
+ if (const GlobalVariable *GVar = dyn_cast<GlobalVariable>(&V)) {
// External GV with no initializer must be .import. In llvm, global
// variable definitions must be initialized. Though PISA allows
// a GV definition with no initializer, we can safely determine the
@@ -581,7 +597,7 @@ static void collectIntelHostAccessMetadata(PISA::GlobalVariableDcl &PGV,
if (!MD || MD->getNumOperands() < 2)
return;
- auto *NameMD = dyn_cast<MDString>(MD->getOperand(1));
+ MDString *NameMD = dyn_cast<MDString>(MD->getOperand(1));
if (!NameMD)
return;
@@ -590,7 +606,7 @@ static void collectIntelHostAccessMetadata(PISA::GlobalVariableDcl &PGV,
void PISAAsmPrinter::collectGlobalVariable(PISA::GlobalVariableDcl &PGV,
const GlobalVariable &GV) {
- auto &DL = GV.getParent()->getDataLayout();
+ const DataLayout &DL = GV.getParent()->getDataLayout();
PGV.Dcl.Linkage = collectLinkage(GV);
PGV.Dcl.SS =
PISA::mapAddrSpaceToStorageSpace(GV.getType()->getAddressSpace());
@@ -611,7 +627,7 @@ void PISAAsmPrinter::collectFunctionDeclaration(PISA::FunctionDeclaration &Dcl,
assert(F.getCallingConv() != CallingConv::PISA_KERNEL);
collectFunctionDirectiveAndName(Dcl.DN, F);
- for (auto &P : F.args()) {
+ for (const Argument &P : F.args()) {
PISA::FunctionDeclParam Param;
if (P.getType()->getScalarSizeInBits() == 1) {
Param.Ty = LLT::integer(8);
@@ -627,24 +643,25 @@ void PISAAsmPrinter::collectKernelParameters(PISA::FunctionSignature &Sig) {
MF->getInfo<PISAMachineFunctionInfo>();
// print params
- auto &DL = MF->getFunction().getParent()->getDataLayout();
+ const DataLayout &DL = MF->getFunction().getParent()->getDataLayout();
for (unsigned Index = 0; Index < MF->getFunction().arg_size(); ++Index) {
- auto [Size, IsByRef] = MFInfo->getArgInfo(Index);
+ std::pair<unsigned, bool> ArgInfo = MFInfo->getArgInfo(Index);
PISA::KernelParameter Param;
- Param.Size = Size;
- auto *ArgTy = MF->getFunction().getArg(Index)->getType()->getScalarType();
- auto Align = alignTo(PowerOf2Ceil(DL.getABITypeAlign(ArgTy).value()), 4);
+ Param.Size = ArgInfo.first;
+ Type *ArgTy = MF->getFunction().getArg(Index)->getType()->getScalarType();
+ uint64_t Align =
+ alignTo(PowerOf2Ceil(DL.getABITypeAlign(ArgTy).value()), 4);
if (Align != 8) {
// Kernel parameters are aligned to 8 bytes by default.
Param.Align = Align;
}
- if (ArgTy->isPointerTy() && !IsByRef) {
- auto AS = ArgTy->getPointerAddressSpace();
+ if (ArgTy->isPointerTy() && !ArgInfo.second) {
+ unsigned AS = ArgTy->getPointerAddressSpace();
if ((AS == (unsigned)PISAAS::AddressSpace::CONSTANT) ||
(AS == (unsigned)PISAAS::AddressSpace::GLOBAL) ||
(AS == (unsigned)PISAAS::AddressSpace::SHARED))
Param.AS = ArgTy->getPointerAddressSpace();
- if (auto PtrAlign = MF->getFunction().getParamAlign(Index))
+ if (MaybeAlign PtrAlign = MF->getFunction().getParamAlign(Index))
Param.PtrAlign = PtrAlign->value();
}
@@ -653,7 +670,7 @@ void PISAAsmPrinter::collectKernelParameters(PISA::FunctionSignature &Sig) {
const Function &F = MF->getFunction();
if (MDNode *MD = F.getMetadata("kernel_arg_name"))
if (Index < MD->getNumOperands())
- if (auto *S = dyn_cast<MDString>(MD->getOperand(Index)))
+ if (MDString *S = dyn_cast<MDString>(MD->getOperand(Index)))
if (!S->getString().empty())
Param.ArgName = S->getString().str();
@@ -663,11 +680,11 @@ void PISAAsmPrinter::collectKernelParameters(PISA::FunctionSignature &Sig) {
void PISAAsmPrinter::collectFunctionParameters(PISA::FunctionSignature &Sig) {
llvm::SmallVector<const MachineInstr *, 8> FuncParamInsts;
- for (auto &MBB : *MF) {
+ for (MachineBasicBlock &MBB : *MF) {
// FunctionParam must be contiguous and in the same BB
// Find the iterator of the first FunctionParam inst and iterate from it
// to collect all FunctionParam insts
- auto MIIt = find_if(
+ MachineBasicBlock::iterator MIIt = find_if(
MBB, [&](MachineInstr &MI) { return TII->isFunctionParamInstr(MI); });
for (; MIIt != MBB.end(); ++MIIt) {
@@ -683,10 +700,11 @@ void PISAAsmPrinter::collectFunctionParameters(PISA::FunctionSignature &Sig) {
});
// Collect params
- for (auto *MI : FuncParamInsts) {
+ for (const MachineInstr *MI : FuncParamInsts) {
const MachineOperand &MO = MI->getOperand(0);
assert(MO.getSubReg() == 0 && "no swizzle allowed on args!");
- const auto *RC = MF->getRegInfo().getRegClass(MO.getReg());
+ const TargetRegisterClass *RC =
+ MF->getRegInfo().getRegClass(MO.getReg());
PISA::FunctionParameter Param;
unsigned NumElts = TRI->getNumEltsFromRegClass(RC);
unsigned EltSize = TRI->getBitSizeFromRegClass(RC);
@@ -702,13 +720,11 @@ static std::string getNameFromType(Type *Ty, bool IsSigned) {
std::string Name = "unknown";
switch (Ty->getTypeID()) {
default:
- llvm_unreachable("unsupported type");
- break;
+ reportFatalUsageError("unsupported PISA kernel vector type hint");
case Type::IntegerTyID: {
switch (Ty->getIntegerBitWidth()) {
default:
- llvm_unreachable("unsupported integer type");
- break;
+ reportFatalUsageError("unsupported PISA kernel vector type hint integer");
case 8:
Name = IsSigned ? "char" : "uchar";
break;
@@ -733,7 +749,7 @@ static std::string getNameFromType(Type *Ty, bool IsSigned) {
Name = "double";
break;
case Type::FixedVectorTyID: {
- auto *VecTy = cast<FixedVectorType>(Ty);
+ FixedVectorType *VecTy = cast<FixedVectorType>(Ty);
Name = getNameFromType(VecTy->getElementType(), IsSigned) +
std::to_string(VecTy->getNumElements());
} break;
@@ -766,13 +782,15 @@ void PISAAsmPrinter::collectFunctionDirectiveAndName(
PISA::KernelAttributeType::REQD_WORK_GROUP_SIZE},
{"vec_type_hint", PISA::KernelAttributeType::VEC_TYPE_HINT}};
- for (auto [MetadataName, EnumVal] : AvailableKernelMetadataNodeTypes) {
- MDNode *Node = dyn_cast_or_null<MDNode>(F.getMetadata(MetadataName));
+ for (const std::pair<StringRef, PISA::KernelAttributeType> &MetadataType :
+ AvailableKernelMetadataNodeTypes) {
+ MDNode *Node =
+ dyn_cast_or_null<MDNode>(F.getMetadata(MetadataType.first));
if (!Node)
continue;
- auto &KernelAttr = DN.KernelAttrs.emplace_back();
- KernelAttr.KernelAttrType = EnumVal;
- switch (EnumVal) {
+ PISA::KernelAttribute &KernelAttr = DN.KernelAttrs.emplace_back();
+ KernelAttr.KernelAttrType = MetadataType.second;
+ switch (MetadataType.second) {
case llvm::PISA::KernelAttributeType::REQD_WORK_GROUP_SIZE: {
KernelAttr.KernelAttrValues.emplace<std::vector<uint32_t>>();
std::transform(Node->op_begin(), Node->op_end(),
@@ -791,7 +809,7 @@ void PISAAsmPrinter::collectFunctionDirectiveAndName(
Metadata *Op1 = Node->getOperand(1);
ConstantInt *CI =
cast<ConstantInt>(cast<ValueAsMetadata>(Op1)->getValue());
- auto TypeName =
+ std::string TypeName =
getNameFromType(cast<ValueAsMetadata>(Op0)->getType(), CI->isOne());
KernelAttr.KernelAttrValues.emplace<std::string>(TypeName);
} break;
@@ -843,7 +861,7 @@ void PISAAsmPrinter::printOperand(const MachineInstr *MI, int OpNum,
switch (MO.getType()) {
case MachineOperand::MO_Register: {
- auto Reg = MO.getReg();
+ Register Reg = MO.getReg();
if (Reg.isPhysical())
O << PISAInstPrinter::getRegisterName(Reg);
else {
@@ -881,7 +899,7 @@ void PISAAsmPrinter::printOperand(const MachineInstr *MI, int OpNum,
case MachineOperand::MO_JumpTableIndex:
case MachineOperand::MO_ConstantPoolIndex:
default:
- llvm_unreachable("<unknown operand type>");
+ reportFatalUsageError("unsupported PISA assembly operand type");
}
}
@@ -896,8 +914,8 @@ bool PISAAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
}
std::string PISAAsmPrinter::getVirtualRegisterName(Register R) const {
- auto &MRI = MF->getRegInfo();
- const auto *RC = MRI.getRegClass(R);
+ const MachineRegisterInfo &MRI = MF->getRegInfo();
+ const TargetRegisterClass *RC = MRI.getRegClass(R);
std::string Name;
raw_string_ostream O(Name);
diff --git a/llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint.ll b/llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint.ll
new file mode 100644
index 0000000000000..901af7e28738e
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint.ll
@@ -0,0 +1,11 @@
+; RUN: not llc -mtriple=pisa -filetype=asm %s -o /dev/null 2>&1 | FileCheck %s
+
+target triple = "pisa"
+
+define pisa_kernel void @unsupported_vector_hint() !vec_type_hint !0 {
+ ret void
+}
+
+!0 = !{<2 x i2> zeroinitializer, i32 0}
+
+; CHECK: LLVM ERROR: unsupported PISA kernel vector type hint integer
diff --git a/llvm/test/CodeGen/PISA/load.ll b/llvm/test/CodeGen/PISA/load.ll
index b3fa8ad1ac728..127cbccb1fab1 100644
--- a/llvm/test/CodeGen/PISA/load.ll
+++ b/llvm/test/CodeGen/PISA/load.ll
@@ -1,7 +1,7 @@
; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
-; RUN: llc -O0 < %s -march=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O0 %s
-; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O2 %s
-; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs
+; RUN: llc -O0 < %s -mtriple=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O0 %s
+; RUN: llc < %s -mtriple=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O2 %s
+; RUN: llc < %s -mtriple=pisa -O0 -verify-machineinstrs
define i8 @i8(ptr addrspace(4) noundef %0) {
; CHECK-LABEL: .8b @i8(
diff --git a/llvm/test/CodeGen/PISA/return.ll b/llvm/test/CodeGen/PISA/return.ll
index 258143595bbf9..06ca44d0d809b 100644
--- a/llvm/test/CodeGen/PISA/return.ll
+++ b/llvm/test/CodeGen/PISA/return.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
-; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck %s
-; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
+; RUN: llc < %s -mtriple=pisa -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -mtriple=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
define i64 @return_value(i64 %a, i64 %b) {
; O0-LABEL: .64b @return_value(
diff --git a/llvm/test/CodeGen/PISA/store.ll b/llvm/test/CodeGen/PISA/store.ll
index 687ebdab1aba5..cf0abe31a8e61 100644
--- a/llvm/test/CodeGen/PISA/store.ll
+++ b/llvm/test/CodeGen/PISA/store.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
-; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck %s
-; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
+; RUN: llc < %s -mtriple=pisa -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -mtriple=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
define void @i8(ptr addrspace(4) noundef %0, i8 %1) {
; O0-LABEL: void @i8(
diff --git a/llvm/test/CodeGen/PISA/unreachable.ll b/llvm/test/CodeGen/PISA/unreachable.ll
index 0e9fb95e4de91..3acc08219afc2 100644
--- a/llvm/test/CodeGen/PISA/unreachable.ll
+++ b/llvm/test/CodeGen/PISA/unreachable.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 5
-; RUN: llc < %s -march=pisa -verify-machineinstrs | FileCheck %s
-; RUN: llc < %s -march=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
+; RUN: llc < %s -mtriple=pisa -verify-machineinstrs | FileCheck %s
+; RUN: llc < %s -mtriple=pisa -O0 -verify-machineinstrs | FileCheck --check-prefix=O0 %s
define pisa_kernel void @test_unreachable(i32 %a, i32 %b) {
; O0-LABEL: @test_unreachable(
>From 3d8e05b96cc4890714099a9b8c67c33fd496ea22 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:06:26 +0000
Subject: [PATCH 04/11] Inline comparison into the assertion
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index 6f4ec4267923b..87b734d427516 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -511,9 +511,9 @@ void PISAAsmPrinter::updateFuncParamIdxs(PISA::DataTypes &DTs) {
// Sanity check that all function parameter indexes are < the total
// number of function parameters of that type (recorded in DTs)
- [[maybe_unused]] bool ValidIdx =
- Info.Idx < DTs.getInfo(NumElts, BitWidth, Info.Type).RegCounter;
- assert(ValidIdx && "function parameter index out of range!");
+ assert((Info.Idx <
+ DTs.getInfo(NumElts, BitWidth, Info.Type).RegCounter) &&
+ "function parameter index out of range!");
}
}
>From e17a94ecfd3ba004b7f202f1812dd6b7ee9c351c Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:13:43 +0000
Subject: [PATCH 05/11] Fix printing FP immediates
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 11 +++++++--
.../PISA/asm-printer-fp-inline-asm.mir | 23 +++++++++++++++++++
2 files changed, 32 insertions(+), 2 deletions(-)
create mode 100644 llvm/test/CodeGen/PISA/asm-printer-fp-inline-asm.mir
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index 87b734d427516..b049a12d91992 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -18,6 +18,7 @@
#include "PISATargetMachine.h"
#include "PISAUtils.h"
#include "TargetInfo/PISATargetInfo.h"
+#include "llvm/ADT/APInt.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallSet.h"
@@ -45,6 +46,7 @@
#include "llvm/Support/Endian.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/FileSystem.h"
+#include "llvm/Support/Format.h"
#include "llvm/Support/PISAAddrSpace.h"
#include "llvm/Support/Regex.h"
#include "llvm/Support/raw_ostream.h"
@@ -874,9 +876,14 @@ void PISAAsmPrinter::printOperand(const MachineInstr *MI, int OpNum,
O << MO.getImm();
break;
- case MachineOperand::MO_FPImmediate:
- O << MO.getFPImm();
+ case MachineOperand::MO_FPImmediate: {
+ const APInt Bits = MO.getFPImm()->getValueAPF().bitcastToAPInt();
+ if (Bits.getBitWidth() > 64)
+ reportFatalUsageError("unsupported PISA floating-point immediate");
+ O << format_hex(Bits.getZExtValue(), Bits.getBitWidth() / 4 + 2,
+ /*Upper=*/true);
break;
+ }
case MachineOperand::MO_MachineBasicBlock:
O << *MO.getMBB()->getSymbol();
diff --git a/llvm/test/CodeGen/PISA/asm-printer-fp-inline-asm.mir b/llvm/test/CodeGen/PISA/asm-printer-fp-inline-asm.mir
new file mode 100644
index 0000000000000..8dca5145eb8e8
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/asm-printer-fp-inline-asm.mir
@@ -0,0 +1,23 @@
+# RUN: llc -mtriple=pisa -start-after=irtranslator -verify-machineinstrs \
+# RUN: -filetype=asm %s -o - | FileCheck %s
+
+# Inline assembly operands must use the same exact-width hexadecimal format as
+# PISA floating-point instruction immediates, not ConstantFP pointer addresses.
+
+---
+name: fp_inline_asm
+selected: true
+tracksRegLiveness: true
+body: |
+ bb.0:
+ ; CHECK-LABEL: @fp_inline_asm()
+ ; CHECK: // half 0x3C00
+ ; CHECK: // bfloat 0x3F80
+ ; CHECK: // float 0x3F800000
+ ; CHECK: // double 0x3FF0000000000000
+ INLINEASM &"// half $0", sideeffect, imm, half 1.000000e+00
+ INLINEASM &"// bfloat $0", sideeffect, imm, bfloat 1.000000e+00
+ INLINEASM &"// float $0", sideeffect, imm, float 1.000000e+00
+ INLINEASM &"// double $0", sideeffect, imm, double 1.000000e+00
+ ret
+...
>From 140f0f15e80cd4cda2062a78490a544aa7ab6ce3 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:15:05 +0000
Subject: [PATCH 06/11] Remove -verify-machineinstr
---
llvm/test/CodeGen/PISA/load.ll | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/llvm/test/CodeGen/PISA/load.ll b/llvm/test/CodeGen/PISA/load.ll
index 127cbccb1fab1..3887ca51efa47 100644
--- a/llvm/test/CodeGen/PISA/load.ll
+++ b/llvm/test/CodeGen/PISA/load.ll
@@ -1,7 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_pisa_test_checks.py UTC_ARGS: --version 6
-; RUN: llc -O0 < %s -mtriple=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O0 %s
-; RUN: llc < %s -mtriple=pisa -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-O2 %s
-; RUN: llc < %s -mtriple=pisa -O0 -verify-machineinstrs
+; RUN: llc -O0 < %s -mtriple=pisa | FileCheck --check-prefixes=CHECK,CHECK-O0 %s
+; RUN: llc < %s -mtriple=pisa | FileCheck --check-prefixes=CHECK,CHECK-O2 %s
define i8 @i8(ptr addrspace(4) noundef %0) {
; CHECK-LABEL: .8b @i8(
>From d2edd3e137b429da28d96cfbc5fab5df0a6c2a82 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:18:35 +0000
Subject: [PATCH 07/11] Add a test for reportFatalUsageError
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 5 ++---
.../CodeGen/PISA/asm-printer-unusual-vector-global.ll | 8 ++++++++
2 files changed, 10 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/CodeGen/PISA/asm-printer-unusual-vector-global.ll
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index b049a12d91992..6cc652ad30d9e 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -309,10 +309,9 @@ void PISAAsmPrinter::FlattenGlobal::process(const Constant *C) {
ConstantInt *CI = dyn_cast_or_null<ConstantInt>(ConstantFoldConstant(
ConstantExpr::getBitCast(const_cast<ConstantVector *>(CV), IntT),
DL));
- if (!CI) {
+ if (!CI)
reportFatalUsageError(
- "Cannot lower vector global with unusual element type");
- }
+ "cannot lower vector global with unusual element type");
emitGlobalConstantLargeInt(CI);
uint64_t EmittedSize = DL.getTypeStoreSize(CV->getType());
if (unsigned Padding = Size - EmittedSize)
diff --git a/llvm/test/CodeGen/PISA/asm-printer-unusual-vector-global.ll b/llvm/test/CodeGen/PISA/asm-printer-unusual-vector-global.ll
new file mode 100644
index 0000000000000..eef240a2f9b5b
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/asm-printer-unusual-vector-global.ll
@@ -0,0 +1,8 @@
+; RUN: not llc -mtriple=pisa -filetype=asm %s -o /dev/null 2>&1 | FileCheck %s
+
+target triple = "pisa"
+
+ at target = addrspace(1) global i8 0
+ at bad = addrspace(1) global <2 x i24> <i24 ptrtoint (ptr addrspace(1) @target to i24), i24 1>
+
+; CHECK: LLVM ERROR: cannot lower vector global with unusual element type
>From 0c660ac6dda98b62394fc41c9adcbdf210b88f42 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:24:11 +0000
Subject: [PATCH 08/11] Document why AsmPrinter::emitGlobalConstant is not used
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index 6cc652ad30d9e..07be17c6dfa33 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -95,6 +95,10 @@ class PISAAsmPrinter : public AsmPrinter {
PISA::DataTypes *DTs = nullptr;
bool GlobalsEmitted = false;
+ // AsmPrinter::emitGlobalConstant writes values and padding directly to
+ // MCStreamer. PISA initializers belong inside a single .global or .const
+ // declaration, so collect their typed values, symbols, and zero runs for
+ // PISATargetStreamer instead.
class FlattenGlobal {
public:
FlattenGlobal(const Constant *C, PISA::VariableInit &VI,
>From 207eb29e083a5fd7bfca872d15702570c588e6c1 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:29:11 +0000
Subject: [PATCH 09/11] Remove the user scan and add a test
---
llvm/lib/Target/PISA/PISAAsmPrinter.cpp | 18 ------------------
.../PISA/asm-printer-annotation-globals.ll | 19 +++++++++++++++++++
2 files changed, 19 insertions(+), 18 deletions(-)
create mode 100644 llvm/test/CodeGen/PISA/asm-printer-annotation-globals.ll
diff --git a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
index 07be17c6dfa33..aa8656c0fa7e2 100644
--- a/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
+++ b/llvm/lib/Target/PISA/PISAAsmPrinter.cpp
@@ -381,24 +381,6 @@ static bool isIgnoredIntrinsicGlobal(const GlobalVariable &GV) {
if (GV.getSection() == "llvm.metadata")
return true;
- // Skip globals only used as annotation strings by llvm.ptr.annotation.
- // These are metadata for the annotation intrinsic, not real data.
- if (GV.hasPrivateLinkage() && GV.isConstant() &&
- all_of(GV.users(), [](const User *U) {
- if (const ConstantExpr *CE = dyn_cast<ConstantExpr>(U))
- return all_of(CE->users(), [](const User *UU) {
- const CallInst *CI = dyn_cast<CallInst>(UU);
- return CI && CI->getCalledFunction() &&
- CI->getCalledFunction()->getIntrinsicID() ==
- Intrinsic::ptr_annotation;
- });
- const CallInst *CI = dyn_cast<CallInst>(U);
- return CI && CI->getCalledFunction() &&
- CI->getCalledFunction()->getIntrinsicID() ==
- Intrinsic::ptr_annotation;
- }))
- return true;
-
if (!GV.hasAppendingLinkage())
return false;
diff --git a/llvm/test/CodeGen/PISA/asm-printer-annotation-globals.ll b/llvm/test/CodeGen/PISA/asm-printer-annotation-globals.ll
new file mode 100644
index 0000000000000..c911fabe06b0f
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/asm-printer-annotation-globals.ll
@@ -0,0 +1,19 @@
+; RUN: llc -mtriple=pisa -filetype=asm %s -o - | FileCheck %s
+
+target triple = "pisa"
+
+ at metadata_string = private addrspace(2) constant [5 x i8] c"meta\00", section "llvm.metadata"
+ at ordinary_string = private addrspace(2) constant [6 x i8] c"plain\00"
+ at data = addrspace(1) global i32 42
+
+define ptr @annotate(ptr %p) {
+ %annotated = call ptr @llvm.ptr.annotation.p0.p2(ptr %p, ptr addrspace(2) @metadata_string, ptr addrspace(2) @ordinary_string, i32 1, ptr addrspace(2) null)
+ ret ptr %annotated
+}
+
+declare ptr @llvm.ptr.annotation.p0.p2(ptr, ptr addrspace(2), ptr addrspace(2), i32, ptr addrspace(2))
+
+; CHECK-NOT: @metadata_string
+; CHECK: @ordinary_string =
+; CHECK: @data =
+; CHECK-NOT: @metadata_string
>From 32605944e507d74943011bfcd3b1a980de60ab82 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:33:23 +0000
Subject: [PATCH 10/11] Add asm-printer-special-globals.ll to test
isIgnoredIntrinsicGlobal
---
.../PISA/asm-printer-special-globals.ll | 48 +++++++++++++++++++
1 file changed, 48 insertions(+)
create mode 100644 llvm/test/CodeGen/PISA/asm-printer-special-globals.ll
diff --git a/llvm/test/CodeGen/PISA/asm-printer-special-globals.ll b/llvm/test/CodeGen/PISA/asm-printer-special-globals.ll
new file mode 100644
index 0000000000000..2e9daef884abd
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/asm-printer-special-globals.ll
@@ -0,0 +1,48 @@
+; RUN: split-file %s %t
+; RUN: llc -mtriple=pisa -filetype=asm %t/used.ll -o - | FileCheck --check-prefix=USED %s
+; RUN: not llc -mtriple=pisa -filetype=asm %t/ctors.ll -o /dev/null 2>&1 | FileCheck --check-prefix=CTORS %s
+; RUN: not llc -mtriple=pisa -filetype=asm %t/dtors.ll -o /dev/null 2>&1 | FileCheck --check-prefix=DTORS %s
+; RUN: not llc -mtriple=pisa -filetype=asm %t/unknown.ll -o /dev/null 2>&1 | FileCheck --check-prefix=UNKNOWN %s
+
+; USED-NOT: @llvm.used
+; USED-NOT: @llvm.compiler.used
+; USED: @data =
+; USED-NOT: @llvm.used
+; USED-NOT: @llvm.compiler.used
+; CTORS: LLVM ERROR: llvm.global_ctors is not supported by the PISA backend
+; DTORS: LLVM ERROR: llvm.global_dtors is not supported by the PISA backend
+; UNKNOWN: LLVM ERROR: unknown special variable with appending linkage
+
+;--- used.ll
+target triple = "pisa"
+
+ at data = addrspace(1) global i32 7
+ at llvm.used = appending global [1 x ptr addrspace(1)] [ptr addrspace(1) @data]
+ at llvm.compiler.used = appending global [1 x ptr addrspace(1)] [ptr addrspace(1) @data], section "llvm.metadata"
+
+;--- ctors.ll
+target triple = "pisa"
+
+ at llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [
+ { i32, ptr, ptr } { i32 65535, ptr @init, ptr null }
+]
+
+define void @init() {
+ ret void
+}
+
+;--- dtors.ll
+target triple = "pisa"
+
+ at llvm.global_dtors = appending global [1 x { i32, ptr, ptr }] [
+ { i32, ptr, ptr } { i32 65535, ptr @fini, ptr null }
+]
+
+define void @fini() {
+ ret void
+}
+
+;--- unknown.ll
+target triple = "pisa"
+
+ at mystery = appending global [1 x i8] c"x"
>From b5b593013a5ac818870091329dae0a461bde32a2 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal.paszkowski at intel.com>
Date: Tue, 22 Sep 2026 11:36:02 +0000
Subject: [PATCH 11/11] Add asm-printer-unsupported-vector-hint-fp128.ll test
---
.../PISA/asm-printer-unsupported-vector-hint-fp128.ll | 11 +++++++++++
1 file changed, 11 insertions(+)
create mode 100644 llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint-fp128.ll
diff --git a/llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint-fp128.ll b/llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint-fp128.ll
new file mode 100644
index 0000000000000..c30d61ddb6b21
--- /dev/null
+++ b/llvm/test/CodeGen/PISA/asm-printer-unsupported-vector-hint-fp128.ll
@@ -0,0 +1,11 @@
+; RUN: not llc -mtriple=pisa -filetype=asm %s -o /dev/null 2>&1 | FileCheck %s
+
+target triple = "pisa"
+
+define void @k_fp128() !vec_type_hint !0 {
+ ret void
+}
+
+!0 = !{fp128 poison, i32 0}
+
+; CHECK: LLVM ERROR: unsupported PISA kernel vector type hint{{$}}
More information about the llvm-branch-commits
mailing list