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 <[email protected]> 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 <[email protected]> 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()); } _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
