[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
Thu Aug 20 08:50:54 PDT 2026
https://github.com/michalpaszkowski updated https://github.com/llvm/llvm-project/pull/214647
>From 489da4889b8d816d2326a412e2403c34874f7f2d 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 1/2] 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 8c85277a615ca..5d6ab740a3eb9 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
PISACacheHintSelector.cpp
PISACallLowering.cpp
PISAConstProp.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 8b3b867492a231fa45dbceeebb60ea70b204541d 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 2/2] 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());
}
More information about the llvm-branch-commits
mailing list