[llvm] f352ce3 - [SPIR-V] Emit SPIR-V generator magic number and version (#87951)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 22 10:47:50 PDT 2024
Author: Michal Paszkowski
Date: 2024-04-22T10:47:46-07:00
New Revision: f352ce368af39e57d337495d7ca3a21975ede8e6
URL: https://github.com/llvm/llvm-project/commit/f352ce368af39e57d337495d7ca3a21975ede8e6
DIFF: https://github.com/llvm/llvm-project/commit/f352ce368af39e57d337495d7ca3a21975ede8e6.diff
LOG: [SPIR-V] Emit SPIR-V generator magic number and version (#87951)
This patch:
- Adds SPIR-V backend's registered generator magic number to the emitted
binary. The magic number consists of the generator ID (43) and LLVM
major version.
- Adds SPIR-V version to the binary.
- Allows reading the expected (maximum supported) SPIR-V version from
the target triple.
- Uses VersionTuple for representing versions throughout the backend's
codebase.
- Registers v1.6 for spirv32 and spirv64 triple.
See more: https://github.com/KhronosGroup/SPIRV-Headers/commit/7d500c
Added:
llvm/test/CodeGen/SPIRV/physical-layout/generator-magic-number.ll
llvm/test/CodeGen/SPIRV/physical-layout/spirv-version.ll
Modified:
llvm/lib/MC/SPIRVObjectWriter.cpp
llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.cpp
llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.h
llvm/lib/Target/SPIRV/SPIRV.td
llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
llvm/lib/Target/SPIRV/SPIRVSubtarget.h
llvm/lib/TargetParser/Triple.cpp
llvm/test/CodeGen/SPIRV/ComparePointers.ll
llvm/test/CodeGen/SPIRV/empty-opencl32.ll
llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
Removed:
################################################################################
diff --git a/llvm/lib/MC/SPIRVObjectWriter.cpp b/llvm/lib/MC/SPIRVObjectWriter.cpp
index d72d6e07f2e6fd..5d85c5de4e4e10 100644
--- a/llvm/lib/MC/SPIRVObjectWriter.cpp
+++ b/llvm/lib/MC/SPIRVObjectWriter.cpp
@@ -43,10 +43,10 @@ class SPIRVObjectWriter : public MCObjectWriter {
void SPIRVObjectWriter::writeHeader(const MCAssembler &Asm) {
constexpr uint32_t MagicNumber = 0x07230203;
- constexpr uint32_t GeneratorMagicNumber = 0;
+ constexpr uint32_t GeneratorID = 43;
+ constexpr uint32_t GeneratorMagicNumber =
+ (GeneratorID << 16) | (LLVM_VERSION_MAJOR);
constexpr uint32_t Schema = 0;
-
- // Construct SPIR-V version and Bound
const MCAssembler::VersionInfoType &VIT = Asm.getVersionInfo();
uint32_t VersionNumber = 0 | (VIT.Major << 16) | (VIT.Minor << 8);
uint32_t Bound = VIT.Update;
diff --git a/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.cpp b/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.cpp
index b69031adb16730..d96d2bf31b6204 100644
--- a/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.cpp
+++ b/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.cpp
@@ -88,28 +88,28 @@ getSymbolicOperandMnemonic(SPIRV::OperandCategory::OperandCategory Category,
return Name;
}
-uint32_t
+VersionTuple
getSymbolicOperandMinVersion(SPIRV::OperandCategory::OperandCategory Category,
uint32_t Value) {
const SPIRV::SymbolicOperand *Lookup =
SPIRV::lookupSymbolicOperandByCategoryAndValue(Category, Value);
if (Lookup)
- return Lookup->MinVersion;
+ return VersionTuple(Lookup->MinVersion / 10, Lookup->MinVersion % 10);
- return 0;
+ return VersionTuple(0);
}
-uint32_t
+VersionTuple
getSymbolicOperandMaxVersion(SPIRV::OperandCategory::OperandCategory Category,
uint32_t Value) {
const SPIRV::SymbolicOperand *Lookup =
SPIRV::lookupSymbolicOperandByCategoryAndValue(Category, Value);
if (Lookup)
- return Lookup->MaxVersion;
+ return VersionTuple(Lookup->MaxVersion / 10, Lookup->MaxVersion % 10);
- return 0;
+ return VersionTuple();
}
CapabilityList
diff --git a/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.h b/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.h
index 616d2ea71b39b1..990eb1d230bcc7 100644
--- a/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.h
+++ b/llvm/lib/Target/SPIRV/MCTargetDesc/SPIRVBaseInfo.h
@@ -17,6 +17,7 @@
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringRef.h"
+#include "llvm/Support/VersionTuple.h"
#include <string>
namespace llvm {
@@ -214,10 +215,10 @@ using ExtensionList = SmallVector<SPIRV::Extension::Extension, 8>;
std::string
getSymbolicOperandMnemonic(SPIRV::OperandCategory::OperandCategory Category,
int32_t Value);
-uint32_t
+VersionTuple
getSymbolicOperandMinVersion(SPIRV::OperandCategory::OperandCategory Category,
uint32_t Value);
-uint32_t
+VersionTuple
getSymbolicOperandMaxVersion(SPIRV::OperandCategory::OperandCategory Category,
uint32_t Value);
CapabilityList
diff --git a/llvm/lib/Target/SPIRV/SPIRV.td b/llvm/lib/Target/SPIRV/SPIRV.td
index beb55d05307ca9..108c7e6d3861f0 100644
--- a/llvm/lib/Target/SPIRV/SPIRV.td
+++ b/llvm/lib/Target/SPIRV/SPIRV.td
@@ -20,19 +20,6 @@ class Proc<string Name, list<SubtargetFeature> Features>
def : Proc<"generic", []>;
-def SPIRV10 : SubtargetFeature<"spirv1.0", "SPIRVVersion", "10",
- "Use SPIR-V version 1.0">;
-def SPIRV11 : SubtargetFeature<"spirv1.1", "SPIRVVersion", "11",
- "Use SPIR-V version 1.1">;
-def SPIRV12 : SubtargetFeature<"spirv1.2", "SPIRVVersion", "12",
- "Use SPIR-V version 1.2">;
-def SPIRV13 : SubtargetFeature<"spirv1.3", "SPIRVVersion", "13",
- "Use SPIR-V version 1.3">;
-def SPIRV14 : SubtargetFeature<"spirv1.4", "SPIRVVersion", "14",
- "Use SPIR-V version 1.4">;
-def SPIRV15 : SubtargetFeature<"spirv1.5", "SPIRVVersion", "15",
- "Use SPIR-V version 1.5">;
-
def SPIRVInstPrinter : AsmWriter {
string AsmWriterClassName = "InstPrinter";
bit isMCAsmWriter = 1;
diff --git a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
index 1de4616fd5b774..2ebe5bdc47715b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
@@ -108,9 +108,9 @@ void SPIRVAsmPrinter::emitEndOfAsmFile(Module &M) {
}
ST = static_cast<const SPIRVTargetMachine &>(TM).getSubtargetImpl();
- uint32_t DecSPIRVVersion = ST->getSPIRVVersion();
- uint32_t Major = DecSPIRVVersion / 10;
- uint32_t Minor = DecSPIRVVersion - Major * 10;
+ VersionTuple SPIRVVersion = ST->getSPIRVVersion();
+ uint32_t Major = SPIRVVersion.getMajor();
+ uint32_t Minor = SPIRVVersion.getMinor().value_or(0);
// Bound is an approximation that accounts for the maximum used register
// number and number of generated OpLabels
unsigned Bound = 2 * (ST->getBound() + 1) + NLabels;
@@ -321,8 +321,8 @@ void SPIRVAsmPrinter::outputEntryPoints() {
// the Input and Output storage classes. Starting with version 1.4,
// the interface's storage classes are all storage classes used in
// declaring all global variables referenced by the entry point call tree.
- if (ST->getSPIRVVersion() >= 14 || SC == SPIRV::StorageClass::Input ||
- SC == SPIRV::StorageClass::Output) {
+ if (ST->isAtLeastSPIRVVer(VersionTuple(1, 4)) ||
+ SC == SPIRV::StorageClass::Input || SC == SPIRV::StorageClass::Output) {
MachineFunction *MF = MI->getMF();
Register Reg = MAI->getRegisterAlias(MF, MI->getOperand(0).getReg());
InterfaceIDs.insert(Reg);
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 8395d4b2bf66b5..235f947901d837 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -76,18 +76,20 @@ getSymbolicOperandRequirements(SPIRV::OperandCategory::OperandCategory Category,
SPIRV::RequirementHandler &Reqs) {
static AvoidCapabilitiesSet
AvoidCaps; // contains capabilities to avoid if there is another option
- unsigned ReqMinVer = getSymbolicOperandMinVersion(Category, i);
- unsigned ReqMaxVer = getSymbolicOperandMaxVersion(Category, i);
- unsigned TargetVer = ST.getSPIRVVersion();
- bool MinVerOK = !ReqMinVer || !TargetVer || TargetVer >= ReqMinVer;
- bool MaxVerOK = !ReqMaxVer || !TargetVer || TargetVer <= ReqMaxVer;
+
+ VersionTuple ReqMinVer = getSymbolicOperandMinVersion(Category, i);
+ VersionTuple ReqMaxVer = getSymbolicOperandMaxVersion(Category, i);
+ VersionTuple SPIRVVersion = ST.getSPIRVVersion();
+ bool MinVerOK = SPIRVVersion.empty() || SPIRVVersion >= ReqMinVer;
+ bool MaxVerOK =
+ ReqMaxVer.empty() || SPIRVVersion.empty() || SPIRVVersion <= ReqMaxVer;
CapabilityList ReqCaps = getSymbolicOperandCapabilities(Category, i);
ExtensionList ReqExts = getSymbolicOperandExtensions(Category, i);
if (ReqCaps.empty()) {
if (ReqExts.empty()) {
if (MinVerOK && MaxVerOK)
return {true, {}, {}, ReqMinVer, ReqMaxVer};
- return {false, {}, {}, 0, 0};
+ return {false, {}, {}, VersionTuple(), VersionTuple()};
}
} else if (MinVerOK && MaxVerOK) {
if (ReqCaps.size() == 1) {
@@ -118,9 +120,13 @@ getSymbolicOperandRequirements(SPIRV::OperandCategory::OperandCategory Category,
if (llvm::all_of(ReqExts, [&ST](const SPIRV::Extension::Extension &Ext) {
return ST.canUseExtension(Ext);
})) {
- return {true, {}, ReqExts, 0, 0}; // TODO: add versions to extensions.
+ return {true,
+ {},
+ ReqExts,
+ VersionTuple(),
+ VersionTuple()}; // TODO: add versions to extensions.
}
- return {false, {}, {}, 0, 0};
+ return {false, {}, {}, VersionTuple(), VersionTuple()};
}
void SPIRVModuleAnalysis::setBaseInfo(const Module &M) {
@@ -510,25 +516,25 @@ void SPIRV::RequirementHandler::addRequirements(
addExtensions(Req.Exts);
- if (Req.MinVer) {
- if (MaxVersion && Req.MinVer > MaxVersion) {
+ if (!Req.MinVer.empty()) {
+ if (!MaxVersion.empty() && Req.MinVer > MaxVersion) {
LLVM_DEBUG(dbgs() << "Conflicting version requirements: >= " << Req.MinVer
<< " and <= " << MaxVersion << "\n");
report_fatal_error("Adding SPIR-V requirements that can't be satisfied.");
}
- if (MinVersion == 0 || Req.MinVer > MinVersion)
+ if (MinVersion.empty() || Req.MinVer > MinVersion)
MinVersion = Req.MinVer;
}
- if (Req.MaxVer) {
- if (MinVersion && Req.MaxVer < MinVersion) {
+ if (!Req.MaxVer.empty()) {
+ if (!MinVersion.empty() && Req.MaxVer < MinVersion) {
LLVM_DEBUG(dbgs() << "Conflicting version requirements: <= " << Req.MaxVer
<< " and >= " << MinVersion << "\n");
report_fatal_error("Adding SPIR-V requirements that can't be satisfied.");
}
- if (MaxVersion == 0 || Req.MaxVer < MaxVersion)
+ if (MaxVersion.empty() || Req.MaxVer < MaxVersion)
MaxVersion = Req.MaxVer;
}
}
@@ -539,7 +545,7 @@ void SPIRV::RequirementHandler::checkSatisfiable(
bool IsSatisfiable = true;
auto TargetVer = ST.getSPIRVVersion();
- if (MaxVersion && TargetVer && MaxVersion < TargetVer) {
+ if (!MaxVersion.empty() && !TargetVer.empty() && MaxVersion < TargetVer) {
LLVM_DEBUG(
dbgs() << "Target SPIR-V version too high for required features\n"
<< "Required max version: " << MaxVersion << " target version "
@@ -547,14 +553,14 @@ void SPIRV::RequirementHandler::checkSatisfiable(
IsSatisfiable = false;
}
- if (MinVersion && TargetVer && MinVersion > TargetVer) {
+ if (!MinVersion.empty() && !TargetVer.empty() && MinVersion > TargetVer) {
LLVM_DEBUG(dbgs() << "Target SPIR-V version too low for required features\n"
<< "Required min version: " << MinVersion
<< " target version " << TargetVer << "\n");
IsSatisfiable = false;
}
- if (MinVersion && MaxVersion && MinVersion > MaxVersion) {
+ if (!MinVersion.empty() && !MaxVersion.empty() && MinVersion > MaxVersion) {
LLVM_DEBUG(
dbgs()
<< "Version is too low for some features and too high for others.\n"
@@ -632,12 +638,13 @@ void RequirementHandler::initAvailableCapabilitiesForOpenCL(
addAvailableCaps({Capability::ImageBasic, Capability::LiteralSampler,
Capability::Image1D, Capability::SampledBuffer,
Capability::ImageBuffer});
- if (ST.isAtLeastOpenCLVer(20))
+ if (ST.isAtLeastOpenCLVer(VersionTuple(2, 0)))
addAvailableCaps({Capability::ImageReadWrite});
}
- if (ST.isAtLeastSPIRVVer(11) && ST.isAtLeastOpenCLVer(22))
+ if (ST.isAtLeastSPIRVVer(VersionTuple(1, 1)) &&
+ ST.isAtLeastOpenCLVer(VersionTuple(2, 2)))
addAvailableCaps({Capability::SubgroupDispatch, Capability::PipeStorage});
- if (ST.isAtLeastSPIRVVer(13))
+ if (ST.isAtLeastSPIRVVer(VersionTuple(1, 3)))
addAvailableCaps({Capability::GroupNonUniform,
Capability::GroupNonUniformVote,
Capability::GroupNonUniformArithmetic,
@@ -645,7 +652,7 @@ void RequirementHandler::initAvailableCapabilitiesForOpenCL(
Capability::GroupNonUniformClustered,
Capability::GroupNonUniformShuffle,
Capability::GroupNonUniformShuffleRelative});
- if (ST.isAtLeastSPIRVVer(14))
+ if (ST.isAtLeastSPIRVVer(VersionTuple(1, 4)))
addAvailableCaps({Capability::DenormPreserve, Capability::DenormFlushToZero,
Capability::SignedZeroInfNanPreserve,
Capability::RoundingModeRTE,
@@ -1162,7 +1169,8 @@ static void collectReqs(const Module &M, SPIRV::ModuleAnalysisInfo &MAI,
auto Node = M.getNamedMetadata("spirv.ExecutionMode");
if (Node) {
// SPV_KHR_float_controls is not available until v1.4
- bool RequireFloatControls = false, VerLower14 = !ST.isAtLeastSPIRVVer(14);
+ bool RequireFloatControls = false,
+ VerLower14 = !ST.isAtLeastSPIRVVer(VersionTuple(1, 4));
for (unsigned i = 0; i < Node->getNumOperands(); i++) {
MDNode *MDN = cast<MDNode>(Node->getOperand(i));
const MDOperand &MDOp = MDN->getOperand(1);
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
index 6e86eed30c5dc1..79226d6d93efb2 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h
@@ -45,13 +45,13 @@ struct Requirements {
const bool IsSatisfiable;
const std::optional<Capability::Capability> Cap;
const ExtensionList Exts;
- const unsigned MinVer; // 0 if no min version is required.
- const unsigned MaxVer; // 0 if no max version is required.
+ const VersionTuple MinVer; // 0 if no min version is required.
+ const VersionTuple MaxVer; // 0 if no max version is required.
Requirements(bool IsSatisfiable = false,
std::optional<Capability::Capability> Cap = {},
- ExtensionList Exts = {}, unsigned MinVer = 0,
- unsigned MaxVer = 0)
+ ExtensionList Exts = {}, VersionTuple MinVer = VersionTuple(),
+ VersionTuple MaxVer = VersionTuple())
: IsSatisfiable(IsSatisfiable), Cap(Cap), Exts(Exts), MinVer(MinVer),
MaxVer(MaxVer) {}
Requirements(Capability::Capability Cap) : Requirements(true, {Cap}) {}
@@ -69,8 +69,8 @@ struct RequirementHandler {
DenseSet<unsigned> AvailableCaps;
SmallSet<Extension::Extension, 4> AllExtensions;
- unsigned MinVersion; // 0 if no min version is defined.
- unsigned MaxVersion; // 0 if no max version is defined.
+ VersionTuple MinVersion; // 0 if no min version is defined.
+ VersionTuple MaxVersion; // 0 if no max version is defined.
// Add capabilities to AllCaps, recursing through their implicitly declared
// capabilities too.
void recursiveAddCapabilities(const CapabilityList &ToPrune);
@@ -79,17 +79,15 @@ struct RequirementHandler {
void initAvailableCapabilitiesForVulkan(const SPIRVSubtarget &ST);
public:
- RequirementHandler() : MinVersion(0), MaxVersion(0) {}
+ RequirementHandler() {}
void clear() {
MinimalCaps.clear();
AllCaps.clear();
AvailableCaps.clear();
AllExtensions.clear();
- MinVersion = 0;
- MaxVersion = 0;
+ MinVersion = VersionTuple();
+ MaxVersion = VersionTuple();
}
- unsigned getMinVersion() const { return MinVersion; }
- unsigned getMaxVersion() const { return MaxVersion; }
const CapabilityList &getMinimalCapabilities() const { return MinimalCaps; }
const SmallSet<Extension::Extension, 4> &getExtensions() const {
return AllExtensions;
diff --git a/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp b/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
index f3864b56e1e97d..7aa0c566c75f3f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
@@ -39,18 +39,43 @@ static cl::opt<std::set<SPIRV::Extension::Extension>, false,
cl::desc("Specify list of enabled SPIR-V extensions"));
// Compare version numbers, but allow 0 to mean unspecified.
-static bool isAtLeastVer(uint32_t Target, uint32_t VerToCompareTo) {
- return Target == 0 || Target >= VerToCompareTo;
+static bool isAtLeastVer(VersionTuple Target, VersionTuple VerToCompareTo) {
+ return Target.empty() || Target >= VerToCompareTo;
}
SPIRVSubtarget::SPIRVSubtarget(const Triple &TT, const std::string &CPU,
const std::string &FS,
const SPIRVTargetMachine &TM)
: SPIRVGenSubtargetInfo(TT, CPU, /*TuneCPU=*/CPU, FS),
- PointerSize(TM.getPointerSizeInBits(/* AS= */ 0)), SPIRVVersion(0),
- OpenCLVersion(0), InstrInfo(),
+ PointerSize(TM.getPointerSizeInBits(/* AS= */ 0)), InstrInfo(),
FrameLowering(initSubtargetDependencies(CPU, FS)), TLInfo(TM, *this),
TargetTriple(TT) {
+ switch (TT.getSubArch()) {
+ case Triple::SPIRVSubArch_v10:
+ SPIRVVersion = VersionTuple(1, 0);
+ break;
+ case Triple::SPIRVSubArch_v11:
+ SPIRVVersion = VersionTuple(1, 1);
+ break;
+ case Triple::SPIRVSubArch_v12:
+ SPIRVVersion = VersionTuple(1, 2);
+ break;
+ case Triple::SPIRVSubArch_v13:
+ SPIRVVersion = VersionTuple(1, 3);
+ break;
+ case Triple::SPIRVSubArch_v14:
+ default:
+ SPIRVVersion = VersionTuple(1, 4);
+ break;
+ case Triple::SPIRVSubArch_v15:
+ SPIRVVersion = VersionTuple(1, 5);
+ break;
+ case Triple::SPIRVSubArch_v16:
+ SPIRVVersion = VersionTuple(1, 6);
+ break;
+ }
+ OpenCLVersion = VersionTuple(2, 2);
+
// The order of initialization is important.
initAvailableExtensions();
initAvailableExtInstSets();
@@ -66,10 +91,6 @@ SPIRVSubtarget::SPIRVSubtarget(const Triple &TT, const std::string &CPU,
SPIRVSubtarget &SPIRVSubtarget::initSubtargetDependencies(StringRef CPU,
StringRef FS) {
ParseSubtargetFeatures(CPU, /*TuneCPU=*/CPU, FS);
- if (SPIRVVersion == 0)
- SPIRVVersion = 14;
- if (OpenCLVersion == 0)
- OpenCLVersion = 22;
return *this;
}
@@ -82,11 +103,11 @@ bool SPIRVSubtarget::canUseExtInstSet(
return AvailableExtInstSets.contains(E);
}
-bool SPIRVSubtarget::isAtLeastSPIRVVer(uint32_t VerToCompareTo) const {
+bool SPIRVSubtarget::isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const {
return isAtLeastVer(SPIRVVersion, VerToCompareTo);
}
-bool SPIRVSubtarget::isAtLeastOpenCLVer(uint32_t VerToCompareTo) const {
+bool SPIRVSubtarget::isAtLeastOpenCLVer(VersionTuple VerToCompareTo) const {
if (!isOpenCLEnv())
return false;
return isAtLeastVer(OpenCLVersion, VerToCompareTo);
@@ -95,7 +116,7 @@ bool SPIRVSubtarget::isAtLeastOpenCLVer(uint32_t VerToCompareTo) const {
// If the SPIR-V version is >= 1.4 we can call OpPtrEqual and OpPtrNotEqual.
// In SPIR-V Translator compatibility mode this feature is not available.
bool SPIRVSubtarget::canDirectlyComparePointers() const {
- return !SPVTranslatorCompat && isAtLeastVer(SPIRVVersion, 14);
+ return !SPVTranslatorCompat && isAtLeastVer(SPIRVVersion, VersionTuple(1, 4));
}
void SPIRVSubtarget::initAvailableExtensions() {
diff --git a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h b/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
index 3b486226a93931..3e4044084266c8 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
+++ b/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
@@ -37,8 +37,8 @@ class SPIRVTargetMachine;
class SPIRVSubtarget : public SPIRVGenSubtargetInfo {
private:
const unsigned PointerSize;
- uint32_t SPIRVVersion;
- uint32_t OpenCLVersion;
+ VersionTuple SPIRVVersion;
+ VersionTuple OpenCLVersion;
SmallSet<SPIRV::Extension::Extension, 4> AvailableExtensions;
SmallSet<SPIRV::InstructionSet::InstructionSet, 4> AvailableExtInstSets;
@@ -81,9 +81,9 @@ class SPIRVSubtarget : public SPIRVGenSubtargetInfo {
TargetTriple.getArch() == Triple::spirv64;
}
bool isVulkanEnv() const { return TargetTriple.getArch() == Triple::spirv; }
- uint32_t getSPIRVVersion() const { return SPIRVVersion; };
- bool isAtLeastSPIRVVer(uint32_t VerToCompareTo) const;
- bool isAtLeastOpenCLVer(uint32_t VerToCompareTo) const;
+ VersionTuple getSPIRVVersion() const { return SPIRVVersion; };
+ bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const;
+ bool isAtLeastOpenCLVer(VersionTuple VerToCompareTo) const;
// TODO: implement command line args or other ways to determine this.
bool hasOpenCLFullProfile() const { return true; }
bool hasOpenCLImageSupport() const { return true; }
diff --git a/llvm/lib/TargetParser/Triple.cpp b/llvm/lib/TargetParser/Triple.cpp
index 77fdf31d4865c0..2c5aee3dfb2f3e 100644
--- a/llvm/lib/TargetParser/Triple.cpp
+++ b/llvm/lib/TargetParser/Triple.cpp
@@ -559,9 +559,11 @@ static Triple::ArchType parseArch(StringRef ArchName) {
.Case("spir64", Triple::spir64)
.Cases("spirv", "spirv1.5", "spirv1.6", Triple::spirv)
.Cases("spirv32", "spirv32v1.0", "spirv32v1.1", "spirv32v1.2",
- "spirv32v1.3", "spirv32v1.4", "spirv32v1.5", Triple::spirv32)
+ "spirv32v1.3", "spirv32v1.4", "spirv32v1.5",
+ "spirv32v1.6", Triple::spirv32)
.Cases("spirv64", "spirv64v1.0", "spirv64v1.1", "spirv64v1.2",
- "spirv64v1.3", "spirv64v1.4", "spirv64v1.5", Triple::spirv64)
+ "spirv64v1.3", "spirv64v1.4", "spirv64v1.5",
+ "spirv64v1.6", Triple::spirv64)
.StartsWith("kalimba", Triple::kalimba)
.Case("lanai", Triple::lanai)
.Case("renderscript32", Triple::renderscript32)
diff --git a/llvm/test/CodeGen/SPIRV/ComparePointers.ll b/llvm/test/CodeGen/SPIRV/ComparePointers.ll
index 6777fc38024b32..408b95579502e0 100644
--- a/llvm/test/CodeGen/SPIRV/ComparePointers.ll
+++ b/llvm/test/CodeGen/SPIRV/ComparePointers.ll
@@ -1,5 +1,5 @@
-; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -mtriple=spirv64v1.3-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.3-unknown-unknown %s -o - -filetype=obj | spirv-val %}
;; kernel void test(int global *in, int global *in2) {
;; if (!in)
diff --git a/llvm/test/CodeGen/SPIRV/empty-opencl32.ll b/llvm/test/CodeGen/SPIRV/empty-opencl32.ll
index 8e826ec35f3781..5b007c7e8adc19 100644
--- a/llvm/test/CodeGen/SPIRV/empty-opencl32.ll
+++ b/llvm/test/CodeGen/SPIRV/empty-opencl32.ll
@@ -1,8 +1,6 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-;; FIXME: ensure Magic Number, version number, generator's magic number, "bound" and "schema" are at least present
-
;; Ensure the required Capabilities are listed.
; CHECK-DAG: OpCapability Kernel
; CHECK-DAG: OpCapability Addresses
diff --git a/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll b/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
index 721e825a1c98e2..d3131e56068570 100644
--- a/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
+++ b/llvm/test/CodeGen/SPIRV/exec_mode_float_control_khr.ll
@@ -1,5 +1,5 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefixes=SPV
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s --mattr=+spirv1.3 --spirv-ext=+SPV_KHR_float_controls -o - | FileCheck %s --check-prefixes=SPVEXT
+; RUN: llc -O0 -mtriple=spirv32v1.3-unknown-unknown %s -o - | FileCheck %s --check-prefixes=SPV
+; RUN: llc -O0 -mtriple=spirv32v1.3-unknown-unknown %s --spirv-ext=+SPV_KHR_float_controls -o - | FileCheck %s --check-prefixes=SPVEXT
define dso_local dllexport spir_kernel void @k_float_controls_0(i32 %ibuf, i32 %obuf) local_unnamed_addr {
entry:
diff --git a/llvm/test/CodeGen/SPIRV/physical-layout/generator-magic-number.ll b/llvm/test/CodeGen/SPIRV/physical-layout/generator-magic-number.ll
new file mode 100644
index 00000000000000..afffd9e69b4544
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/physical-layout/generator-magic-number.ll
@@ -0,0 +1,4 @@
+; REQUIRES: spirv-tools
+; RUN: llc -O0 -mtriple=spirv-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s
+
+; CHECK: Generator: {{.*}}{{43|LLVM SPIR-V Backend}}{{.*}}
diff --git a/llvm/test/CodeGen/SPIRV/physical-layout/spirv-version.ll b/llvm/test/CodeGen/SPIRV/physical-layout/spirv-version.ll
new file mode 100644
index 00000000000000..686c1e97257adc
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/physical-layout/spirv-version.ll
@@ -0,0 +1,16 @@
+; REQUIRES: spirv-tools
+; RUN: llc -O0 -mtriple=spirv64v1.0-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV10
+; RUN: llc -O0 -mtriple=spirv64v1.1-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV11
+; RUN: llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV12
+; RUN: llc -O0 -mtriple=spirv64v1.3-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV13
+; RUN: llc -O0 -mtriple=spirv64v1.4-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV14
+; RUN: llc -O0 -mtriple=spirv64v1.5-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV15
+; RUN: llc -O0 -mtriple=spirv64v1.6-unknown-unknown %s -o - --filetype=obj | spirv-dis | FileCheck %s --check-prefix=CHECK-SPIRV16
+
+; CHECK-SPIRV10: Version: 1.0
+; CHECK-SPIRV11: Version: 1.1
+; CHECK-SPIRV12: Version: 1.2
+; CHECK-SPIRV13: Version: 1.3
+; CHECK-SPIRV14: Version: 1.4
+; CHECK-SPIRV15: Version: 1.5
+; CHECK-SPIRV16: Version: 1.6
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll b/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
index e0c47798cc6d09..cb5bce1375b63d 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll
@@ -1,5 +1,5 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -mtriple=spirv32v1.3-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32v1.3-unknown-unknown %s -o - -filetype=obj | spirv-val %}
;; __kernel void testAtomicCompareExchangeExplicit_cl20(
;; volatile global atomic_int* object,
More information about the llvm-commits
mailing list