[llvm] [SPIR-V] Emit SPIR-V generator magic number and version (PR #87951)

Michal Paszkowski via llvm-commits llvm-commits at lists.llvm.org
Sun Apr 7 16:19:58 PDT 2024


https://github.com/michalpaszkowski created https://github.com/llvm/llvm-project/pull/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

>From 2764a43be37bcabfa904336b5340f55bee356bd8 Mon Sep 17 00:00:00 2001
From: "Levytskyy, Vyacheslav" <vyacheslav.levytskyy at intel.com>
Date: Fri, 5 Apr 2024 12:30:00 -0700
Subject: [PATCH 1/3] implement switch; improve validation of forward calls

---
 llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp     |   5 +-
 llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp |  35 ++-
 llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h   |   7 +-
 llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp   |  17 +-
 llvm/lib/Target/SPIRV/SPIRVISelLowering.h     |   4 +
 llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp   | 223 +++++-------------
 .../SPIRV/branching/OpSwitchUnreachable.ll    |   5 +-
 .../SPIRV/branching/switch-range-check.ll     | 118 +++++++++
 8 files changed, 232 insertions(+), 182 deletions(-)
 create mode 100644 llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll

diff --git a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
index 4eee8062f28248..ffaa7ada9a8060 100644
--- a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
@@ -43,6 +43,8 @@ using namespace llvm;
 
 namespace {
 class SPIRVAsmPrinter : public AsmPrinter {
+  unsigned NLabels = 0;
+
 public:
   explicit SPIRVAsmPrinter(TargetMachine &TM,
                            std::unique_ptr<MCStreamer> Streamer)
@@ -112,7 +114,7 @@ void SPIRVAsmPrinter::emitEndOfAsmFile(Module &M) {
   // TODO: calculate Bound more carefully from maximum used register number,
   // accounting for generated OpLabels and other related instructions if
   // needed.
-  unsigned Bound = 2 * (ST->getBound() + 1);
+  unsigned Bound = 2 * (ST->getBound() + 1) + NLabels;
   bool FlagToRestore = OutStreamer->getUseAssemblerInfoForParsing();
   OutStreamer->setUseAssemblerInfoForParsing(true);
   if (MCAssembler *Asm = OutStreamer->getAssemblerPtr())
@@ -158,6 +160,7 @@ void SPIRVAsmPrinter::emitOpLabel(const MachineBasicBlock &MBB) {
   LabelInst.setOpcode(SPIRV::OpLabel);
   LabelInst.addOperand(MCOperand::createReg(MAI->getOrCreateMBBRegister(MBB)));
   outputMCInst(LabelInst);
+  ++NLabels;
 }
 
 void SPIRVAsmPrinter::emitBasicBlockStart(const MachineBasicBlock &MBB) {
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index b341fcb41d0312..e8ce5a35b457d5 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -460,15 +460,36 @@ void SPIRVEmitIntrinsics::preprocessCompositeConstants(IRBuilder<> &B) {
 }
 
 Instruction *SPIRVEmitIntrinsics::visitSwitchInst(SwitchInst &I) {
-  IRBuilder<> B(I.getParent());
+  BasicBlock *ParentBB = I.getParent();
+  IRBuilder<> B(ParentBB);
+  B.SetInsertPoint(&I);
   SmallVector<Value *, 4> Args;
-  for (auto &Op : I.operands())
-    if (Op.get()->getType()->isSized())
+  SmallVector<BasicBlock *> BBCases;
+  for (auto &Op : I.operands()) {
+    if (Op.get()->getType()->isSized()) {
       Args.push_back(Op);
-  B.SetInsertPoint(&I);
-  B.CreateIntrinsic(Intrinsic::spv_switch, {I.getOperand(0)->getType()},
-                    {Args});
-  return &I;
+    } else if (BasicBlock *BB = dyn_cast<BasicBlock>(Op.get())) {
+      BBCases.push_back(BB);
+      Args.push_back(BlockAddress::get(BB->getParent(), BB));
+    } else {
+      report_fatal_error("Unexpected switch operand");
+    }
+  }
+  CallInst *NewI = B.CreateIntrinsic(Intrinsic::spv_switch,
+                                     {I.getOperand(0)->getType()}, {Args});
+  // remove switch to avoid its unneeded and undesirable unwrap into branches
+  // and conditions
+  I.replaceAllUsesWith(NewI);
+  I.eraseFromParent();
+  // insert artificial and temporary instruction to preserve valid CFG,
+  // it will be removed after IR translation pass
+  B.SetInsertPoint(ParentBB);
+  IndirectBrInst *BrI = B.CreateIndirectBr(
+      Constant::getNullValue(PointerType::getUnqual(ParentBB->getContext())),
+      BBCases.size());
+  for (BasicBlock *BBCase : BBCases)
+    BrI->addDestination(BBCase);
+  return BrI;
 }
 
 Instruction *SPIRVEmitIntrinsics::visitGetElementPtrInst(GetElementPtrInst &I) {
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
index ac799374adce8c..37f575e884ef48 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h
@@ -284,7 +284,12 @@ class SPIRVGlobalRegistry {
   // Return the VReg holding the result of the given OpTypeXXX instruction.
   Register getSPIRVTypeID(const SPIRVType *SpirvType) const;
 
-  void setCurrentFunc(MachineFunction &MF) { CurMF = &MF; }
+  // Return previous value of the current machine function
+  MachineFunction *setCurrentFunc(MachineFunction &MF) {
+    MachineFunction *Ret = CurMF;
+    CurMF = &MF;
+    return Ret;
+  }
 
   // Whether the given VReg has an OpTypeXXX instruction mapped to it with the
   // given opcode (e.g. OpTypeFloat).
diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
index d450078d793fb7..8db54c74f23690 100644
--- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.cpp
@@ -160,12 +160,15 @@ void validateFunCallMachineDef(const SPIRVSubtarget &STI,
             : nullptr;
     if (DefElemType) {
       const Type *DefElemTy = GR.getTypeForSPIRVType(DefElemType);
-      // Switch GR context to the call site instead of the (default) definition
-      // side
-      GR.setCurrentFunc(*FunCall.getParent()->getParent());
+      // validatePtrTypes() works in the context if the call site
+      // When we process historical records about forward calls
+      // we need to switch context to the (forward) call site and
+      // then restore it back to the current machine function.
+      MachineFunction *CurMF =
+          GR.setCurrentFunc(*FunCall.getParent()->getParent());
       validatePtrTypes(STI, CallMRI, GR, FunCall, OpIdx, DefElemType,
                        DefElemTy);
-      GR.setCurrentFunc(*FunDef->getParent()->getParent());
+      GR.setCurrentFunc(*CurMF);
     }
   }
 }
@@ -215,6 +218,11 @@ void validateAccessChain(const SPIRVSubtarget &STI, MachineRegisterInfo *MRI,
 // TODO: the logic of inserting additional bitcast's is to be moved
 // to pre-IRTranslation passes eventually
 void SPIRVTargetLowering::finalizeLowering(MachineFunction &MF) const {
+  // finalizeLowering() is called twice (see GlobalISel/InstructionSelect.cpp)
+  // We'd like to avoid the needless second processing pass.
+  if (ProcessedMF.find(&MF) != ProcessedMF.end())
+    return;
+
   MachineRegisterInfo *MRI = &MF.getRegInfo();
   SPIRVGlobalRegistry &GR = *STI.getSPIRVGlobalRegistry();
   GR.setCurrentFunc(MF);
@@ -302,5 +310,6 @@ void SPIRVTargetLowering::finalizeLowering(MachineFunction &MF) const {
       }
     }
   }
+  ProcessedMF.insert(&MF);
   TargetLowering::finalizeLowering(MF);
 }
diff --git a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
index b01571bfc1eeb5..8c1de7d97d1a3c 100644
--- a/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
+++ b/llvm/lib/Target/SPIRV/SPIRVISelLowering.h
@@ -16,6 +16,7 @@
 
 #include "SPIRVGlobalRegistry.h"
 #include "llvm/CodeGen/TargetLowering.h"
+#include <set>
 
 namespace llvm {
 class SPIRVSubtarget;
@@ -23,6 +24,9 @@ class SPIRVSubtarget;
 class SPIRVTargetLowering : public TargetLowering {
   const SPIRVSubtarget &STI;
 
+  // Record of already processed machine functions
+  mutable std::set<const MachineFunction *> ProcessedMF;
+
 public:
   explicit SPIRVTargetLowering(const TargetMachine &TM,
                                const SPIRVSubtarget &ST)
diff --git a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
index b133f0ae85de20..7e155a36aadbc4 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp
@@ -438,186 +438,75 @@ static void processInstrsWithTypeFolding(MachineFunction &MF,
   }
 }
 
+// Find basic blocks of the switch and replace registers in spv_switch() by its
+// MBB equivalent.
 static void processSwitches(MachineFunction &MF, SPIRVGlobalRegistry *GR,
                             MachineIRBuilder MIB) {
-  // Before IRTranslator pass, calls to spv_switch intrinsic are inserted before
-  // each switch instruction. IRTranslator lowers switches to G_ICMP + G_BRCOND
-  // + G_BR triples. A switch with two cases may be transformed to this MIR
-  // sequence:
-  //
-  //   intrinsic(@llvm.spv.switch), %CmpReg, %Const0, %Const1
-  //   %Dst0 = G_ICMP intpred(eq), %CmpReg, %Const0
-  //   G_BRCOND %Dst0, %bb.2
-  //   G_BR %bb.5
-  // bb.5.entry:
-  //   %Dst1 = G_ICMP intpred(eq), %CmpReg, %Const1
-  //   G_BRCOND %Dst1, %bb.3
-  //   G_BR %bb.4
-  // bb.2.sw.bb:
-  //   ...
-  // bb.3.sw.bb1:
-  //   ...
-  // bb.4.sw.epilog:
-  //   ...
-  //
-  // Sometimes (in case of range-compare switches), additional G_SUBs
-  // instructions are inserted before G_ICMPs. Those need to be additionally
-  // processed.
-  //
-  // This function modifies spv_switch call's operands to include destination
-  // MBBs (default and for each constant value).
-  //
-  // At the end, the function removes redundant [G_SUB] + G_ICMP + G_BRCOND +
-  // G_BR sequences.
-
-  MachineRegisterInfo &MRI = MF.getRegInfo();
-
-  // Collect spv_switches and G_ICMPs across all MBBs in MF.
-  std::vector<MachineInstr *> RelevantInsts;
-
-  // Collect redundant MIs from [G_SUB] + G_ICMP + G_BRCOND + G_BR sequences.
-  // After updating spv_switches, the instructions can be removed.
-  std::vector<MachineInstr *> PostUpdateArtifacts;
-
-  // Temporary set of compare registers. G_SUBs and G_ICMPs relating to
-  // spv_switch use these registers.
-  DenseSet<Register> CompareRegs;
+  DenseMap<const BasicBlock *, MachineBasicBlock *> BB2MBB;
+  SmallVector<std::pair<MachineInstr *, SmallVector<MachineInstr *, 8>>>
+      Switches;
   for (MachineBasicBlock &MBB : MF) {
+    MachineRegisterInfo &MRI = MF.getRegInfo();
+    BB2MBB[MBB.getBasicBlock()] = &MBB;
     for (MachineInstr &MI : MBB) {
+      if (!isSpvIntrinsic(MI, Intrinsic::spv_switch))
+        continue;
       // Calls to spv_switch intrinsics representing IR switches.
-      if (isSpvIntrinsic(MI, Intrinsic::spv_switch)) {
-        assert(MI.getOperand(1).isReg());
-        CompareRegs.insert(MI.getOperand(1).getReg());
-        RelevantInsts.push_back(&MI);
-      }
-
-      // G_SUBs coming from range-compare switch lowering. G_SUBs are found
-      // after spv_switch but before G_ICMP.
-      if (MI.getOpcode() == TargetOpcode::G_SUB && MI.getOperand(1).isReg() &&
-          CompareRegs.contains(MI.getOperand(1).getReg())) {
-        assert(MI.getOperand(0).isReg() && MI.getOperand(1).isReg());
-        Register Dst = MI.getOperand(0).getReg();
-        CompareRegs.insert(Dst);
-        PostUpdateArtifacts.push_back(&MI);
-      }
-
-      // G_ICMPs relating to switches.
-      if (MI.getOpcode() == TargetOpcode::G_ICMP && MI.getOperand(2).isReg() &&
-          CompareRegs.contains(MI.getOperand(2).getReg())) {
-        Register Dst = MI.getOperand(0).getReg();
-        RelevantInsts.push_back(&MI);
-        PostUpdateArtifacts.push_back(&MI);
-        MachineInstr *CBr = MRI.use_begin(Dst)->getParent();
-        assert(CBr->getOpcode() == SPIRV::G_BRCOND);
-        PostUpdateArtifacts.push_back(CBr);
-        MachineInstr *Br = CBr->getNextNode();
-        assert(Br->getOpcode() == SPIRV::G_BR);
-        PostUpdateArtifacts.push_back(Br);
+      SmallVector<MachineInstr *, 8> NewOps;
+      for (unsigned i = 2; i < MI.getNumOperands(); ++i) {
+        Register Reg = MI.getOperand(i).getReg();
+        if (i % 2 == 1) {
+          MachineInstr *ConstInstr = getDefInstrMaybeConstant(Reg, &MRI);
+          NewOps.push_back(ConstInstr);
+        } else {
+          MachineInstr *BuildMBB = MRI.getVRegDef(Reg);
+          assert(BuildMBB &&
+                 BuildMBB->getOpcode() == TargetOpcode::G_BLOCK_ADDR &&
+                 BuildMBB->getOperand(1).isBlockAddress() &&
+                 BuildMBB->getOperand(1).getBlockAddress());
+          NewOps.push_back(BuildMBB);
+        }
       }
+      Switches.push_back(std::make_pair(&MI, NewOps));
     }
   }
 
-  // Update each spv_switch with destination MBBs.
-  for (auto i = RelevantInsts.begin(); i != RelevantInsts.end(); i++) {
-    if (!isSpvIntrinsic(**i, Intrinsic::spv_switch))
-      continue;
-
-    // Currently considered spv_switch.
-    MachineInstr *Switch = *i;
-    // Set the first successor as default MBB to support empty switches.
-    MachineBasicBlock *DefaultMBB = *Switch->getParent()->succ_begin();
-    // Container for mapping values to MMBs.
-    SmallDenseMap<uint64_t, MachineBasicBlock *> ValuesToMBBs;
-
-    // Walk all G_ICMPs to collect ValuesToMBBs. Start at currently considered
-    // spv_switch (i) and break at any spv_switch with the same compare
-    // register (indicating we are back at the same scope).
-    Register CompareReg = Switch->getOperand(1).getReg();
-    for (auto j = i + 1; j != RelevantInsts.end(); j++) {
-      if (isSpvIntrinsic(**j, Intrinsic::spv_switch) &&
-          (*j)->getOperand(1).getReg() == CompareReg)
-        break;
-
-      if (!((*j)->getOpcode() == TargetOpcode::G_ICMP &&
-            (*j)->getOperand(2).getReg() == CompareReg))
-        continue;
-
-      MachineInstr *ICMP = *j;
-      Register Dst = ICMP->getOperand(0).getReg();
-      MachineOperand &PredOp = ICMP->getOperand(1);
-      const auto CC = static_cast<CmpInst::Predicate>(PredOp.getPredicate());
-      (void)CC;
-      assert((CC == CmpInst::ICMP_EQ || CC == CmpInst::ICMP_ULE) &&
-             MRI.hasOneUse(Dst) && MRI.hasOneDef(CompareReg));
-      uint64_t Value = getIConstVal(ICMP->getOperand(3).getReg(), &MRI);
-      MachineInstr *CBr = MRI.use_begin(Dst)->getParent();
-      assert(CBr->getOpcode() == SPIRV::G_BRCOND && CBr->getOperand(1).isMBB());
-      MachineBasicBlock *MBB = CBr->getOperand(1).getMBB();
-
-      // Map switch case Value to target MBB.
-      ValuesToMBBs[Value] = MBB;
-
-      // Add target MBB as successor to the switch's MBB.
-      Switch->getParent()->addSuccessor(MBB);
-
-      // The next MI is always G_BR to either the next case or the default.
-      MachineInstr *NextMI = CBr->getNextNode();
-      assert(NextMI->getOpcode() == SPIRV::G_BR &&
-             NextMI->getOperand(0).isMBB());
-      MachineBasicBlock *NextMBB = NextMI->getOperand(0).getMBB();
-      // Default MBB does not begin with G_ICMP using spv_switch compare
-      // register.
-      if (NextMBB->front().getOpcode() != SPIRV::G_ICMP ||
-          (NextMBB->front().getOperand(2).isReg() &&
-           NextMBB->front().getOperand(2).getReg() != CompareReg)) {
-        // Set default MBB and add it as successor to the switch's MBB.
-        DefaultMBB = NextMBB;
-        Switch->getParent()->addSuccessor(DefaultMBB);
+  SmallPtrSet<MachineInstr *, 8> ToEraseMI;
+  for (auto &SwIt : Switches) {
+    MachineInstr &MI = *SwIt.first;
+    SmallVector<MachineInstr *, 8> &Ins = SwIt.second;
+    SmallVector<MachineOperand, 8> NewOps;
+    for (unsigned i = 0; i < Ins.size(); ++i) {
+      if (Ins[i]->getOpcode() == TargetOpcode::G_BLOCK_ADDR) {
+        BasicBlock *CaseBB =
+            Ins[i]->getOperand(1).getBlockAddress()->getBasicBlock();
+        auto It = BB2MBB.find(CaseBB);
+        if (It == BB2MBB.end())
+          report_fatal_error("cannot find a machine basic block by a basic "
+                             "block in a switch statement");
+        NewOps.push_back(MachineOperand::CreateMBB(It->second));
+        MI.getParent()->addSuccessor(It->second);
+        ToEraseMI.insert(Ins[i]);
+      } else {
+        NewOps.push_back(
+            MachineOperand::CreateCImm(Ins[i]->getOperand(1).getCImm()));
       }
     }
-
-    // Modify considered spv_switch operands using collected Values and
-    // MBBs.
-    SmallVector<const ConstantInt *, 3> Values;
-    SmallVector<MachineBasicBlock *, 3> MBBs;
-    for (unsigned k = 2; k < Switch->getNumExplicitOperands(); k++) {
-      Register CReg = Switch->getOperand(k).getReg();
-      uint64_t Val = getIConstVal(CReg, &MRI);
-      MachineInstr *ConstInstr = getDefInstrMaybeConstant(CReg, &MRI);
-      if (!ValuesToMBBs[Val])
-        continue;
-
-      Values.push_back(ConstInstr->getOperand(1).getCImm());
-      MBBs.push_back(ValuesToMBBs[Val]);
-    }
-
-    for (unsigned k = Switch->getNumExplicitOperands() - 1; k > 1; k--)
-      Switch->removeOperand(k);
-
-    Switch->addOperand(MachineOperand::CreateMBB(DefaultMBB));
-    for (unsigned k = 0; k < Values.size(); k++) {
-      Switch->addOperand(MachineOperand::CreateCImm(Values[k]));
-      Switch->addOperand(MachineOperand::CreateMBB(MBBs[k]));
-    }
-  }
-
-  for (MachineInstr *MI : PostUpdateArtifacts) {
-    MachineBasicBlock *ParentMBB = MI->getParent();
-    MI->eraseFromParent();
-    // If G_ICMP + G_BRCOND + G_BR were the only MIs in MBB, erase this MBB. It
-    // can be safely assumed, there are no breaks or phis directing into this
-    // MBB. However, we need to remove this MBB from the CFG graph. MBBs must be
-    // erased top-down.
-    if (ParentMBB->empty()) {
-      while (!ParentMBB->pred_empty())
-        (*ParentMBB->pred_begin())->removeSuccessor(ParentMBB);
-
-      while (!ParentMBB->succ_empty())
-        ParentMBB->removeSuccessor(ParentMBB->succ_begin());
-
-      ParentMBB->eraseFromParent();
+    for (unsigned i = MI.getNumOperands() - 1; i > 1; --i)
+      MI.removeOperand(i);
+    for (auto &MO : NewOps)
+      MI.addOperand(MO);
+    if (MachineInstr *Next = MI.getNextNode()) {
+      if (isSpvIntrinsic(*Next, Intrinsic::spv_track_constant)) {
+        ToEraseMI.insert(Next);
+        Next = MI.getNextNode();
+      }
+      if (Next && Next->getOpcode() == TargetOpcode::G_BRINDIRECT)
+        ToEraseMI.insert(Next);
     }
   }
+  for (MachineInstr *BlockAddrI : ToEraseMI)
+    BlockAddrI->eraseFromParent();
 }
 
 static bool isImplicitFallthrough(MachineBasicBlock &MBB) {
diff --git a/llvm/test/CodeGen/SPIRV/branching/OpSwitchUnreachable.ll b/llvm/test/CodeGen/SPIRV/branching/OpSwitchUnreachable.ll
index e73efbeade70dc..6eb36e5756ecf6 100644
--- a/llvm/test/CodeGen/SPIRV/branching/OpSwitchUnreachable.ll
+++ b/llvm/test/CodeGen/SPIRV/branching/OpSwitchUnreachable.ll
@@ -1,8 +1,9 @@
 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
 define void @test_switch_with_unreachable_block(i1 %a) {
   %value = zext i1 %a to i32
-; CHECK-SPIRV:      OpSwitch %[[#]] %[[#REACHABLE:]]
+; CHECK-SPIRV:      OpSwitch %[[#]] %[[#UNREACHABLE:]] 0 %[[#REACHABLE:]] 1 %[[#REACHABLE:]]
   switch i32 %value, label %unreachable [
     i32 0, label %reachable
     i32 1, label %reachable
@@ -13,7 +14,7 @@ reachable:
 ; CHECK-SPIRV-NEXT: OpReturn
   ret void
 
-; CHECK-SPIRV:      %[[#]] = OpLabel
+; CHECK-SPIRV:      %[[#UNREACHABLE]] = OpLabel
 ; CHECK-SPIRV-NEXT: OpUnreachable
 unreachable:
   unreachable
diff --git a/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll b/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll
new file mode 100644
index 00000000000000..8ec384a0b07d2d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll
@@ -0,0 +1,118 @@
+; 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 %}
+
+; CHECK: %[[#Var:]] = OpPhi
+; CHECK: OpSwitch %[[#Var]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]] [[#]] %[[#]]
+; CHECK-COUNT-11: OpBranch
+; CHECK-NOT: OpBranch
+
+define spir_func void @foo(i64 noundef %addr, i64 noundef %as) {
+entry:
+  %0 = inttoptr i64 %as to ptr addrspace(4)
+  %1 = load i8, ptr addrspace(4) %0
+  %cmp = icmp sgt i8 %1, 0
+  br i1 %cmp, label %if.then, label %if.end
+
+if.then:                                          ; preds = %entry
+  %add.ptr = getelementptr inbounds i8, ptr addrspace(4) %0, i64 1
+  %2 = load i8, ptr addrspace(4) %add.ptr
+  br label %if.end
+
+if.end:                                           ; preds = %if.then, %entry
+  %shadow_value.0.in = phi i8 [ %2, %if.then ], [ %1, %entry ]
+  switch i8 %shadow_value.0.in, label %sw.default [
+    i8 -127, label %sw.epilog
+    i8 -126, label %sw.bb3
+    i8 -125, label %sw.bb4
+    i8 -111, label %sw.bb5
+    i8 -110, label %sw.bb6
+    i8 -109, label %sw.bb7
+    i8 -15, label %sw.bb8
+    i8 -14, label %sw.bb8
+    i8 -13, label %sw.bb8
+    i8 -124, label %sw.bb9
+    i8 -95, label %sw.bb10
+    i8 -123, label %sw.bb11
+  ]
+
+sw.bb3:                                           ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb4:                                           ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb5:                                           ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb6:                                           ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb7:                                           ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb8:                                           ; preds = %if.end, %if.end, %if.end
+  br label %sw.epilog
+
+sw.bb9:                                           ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb10:                                          ; preds = %if.end
+  br label %sw.epilog
+
+sw.bb11:                                          ; preds = %if.end
+  br label %sw.epilog
+
+sw.default:                                       ; preds = %if.end
+  br label %sw.epilog
+
+sw.epilog:                                        ; preds = %sw.default, %sw.bb11, %sw.bb10, %sw.bb9, %sw.bb8, %sw.bb7, %sw.bb6, %sw.bb5, %sw.bb4, %sw.bb3, %if.end
+  br label %exit
+
+if.then.i:                                        ; preds = %sw.epilog
+  br label %exit
+
+for.cond.i:                                       ; preds = %for.inc.i, %if.then.i
+  br label %exit
+
+for.inc.i:                                        ; preds = %for.cond.i
+  br label %exit
+
+if.end.i:                                         ; preds = %for.cond.i, %if.then.i
+  br label %exit
+
+if.end18.thread.i:                                ; preds = %if.end.i
+  br label %5
+
+for.cond8.i:                                      ; preds = %for.inc14.i, %if.end.i
+  br label %exit
+
+for.inc14.i:                                      ; preds = %for.cond8.i
+  br label %exit
+
+if.end18.i:                                       ; preds = %for.cond8.i
+  br label %5
+
+5:                                                ; preds = %if.end18.i, %if.end18.thread.i
+  br label %for.cond25.i
+
+for.cond25.i:                                     ; preds = %for.body29.i, %5
+  br label %exit
+
+for.cond.cleanup27.i:                             ; preds = %for.cond25.i
+  br label %for.cond41.i
+
+for.body29.i:                                     ; preds = %for.cond25.i
+  br label %for.cond25.i
+
+for.cond41.i:                                     ; preds = %for.body45.i, %for.cond.cleanup27.i
+  br label %exit
+
+for.cond.cleanup43.i:                             ; preds = %for.cond41.i
+  br label %exit
+
+for.body45.i:                                     ; preds = %for.cond41.i
+  br label %for.cond41.i
+
+exit:
+  ret void
+}

>From e2cf06bb04d62adc4646a13ddeca58950449bc48 Mon Sep 17 00:00:00 2001
From: "Levytskyy, Vyacheslav" <vyacheslav.levytskyy at intel.com>
Date: Sat, 6 Apr 2024 00:14:01 -0700
Subject: [PATCH 2/3] fix a test, edit a comment

---
 llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp     |  5 +-
 .../SPIRV/branching/switch-range-check.ll     | 85 +++++--------------
 2 files changed, 22 insertions(+), 68 deletions(-)

diff --git a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
index ffaa7ada9a8060..1de4616fd5b774 100644
--- a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp
@@ -111,9 +111,8 @@ void SPIRVAsmPrinter::emitEndOfAsmFile(Module &M) {
   uint32_t DecSPIRVVersion = ST->getSPIRVVersion();
   uint32_t Major = DecSPIRVVersion / 10;
   uint32_t Minor = DecSPIRVVersion - Major * 10;
-  // TODO: calculate Bound more carefully from maximum used register number,
-  // accounting for generated OpLabels and other related instructions if
-  // needed.
+  // Bound is an approximation that accounts for the maximum used register
+  // number and number of generated OpLabels
   unsigned Bound = 2 * (ST->getBound() + 1) + NLabels;
   bool FlagToRestore = OutStreamer->getUseAssemblerInfoForParsing();
   OutStreamer->setUseAssemblerInfoForParsing(true);
diff --git a/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll b/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll
index 8ec384a0b07d2d..85a4d4db089cb4 100644
--- a/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll
+++ b/llvm/test/CodeGen/SPIRV/branching/switch-range-check.ll
@@ -8,19 +8,19 @@
 
 define spir_func void @foo(i64 noundef %addr, i64 noundef %as) {
 entry:
-  %0 = inttoptr i64 %as to ptr addrspace(4)
-  %1 = load i8, ptr addrspace(4) %0
-  %cmp = icmp sgt i8 %1, 0
+  %src = inttoptr i64 %as to ptr addrspace(4)
+  %val = load i8, ptr addrspace(4) %src
+  %cmp = icmp sgt i8 %val, 0
   br i1 %cmp, label %if.then, label %if.end
 
-if.then:                                          ; preds = %entry
-  %add.ptr = getelementptr inbounds i8, ptr addrspace(4) %0, i64 1
-  %2 = load i8, ptr addrspace(4) %add.ptr
+if.then:
+  %add.ptr = getelementptr inbounds i8, ptr addrspace(4) %src, i64 1
+  %cond = load i8, ptr addrspace(4) %add.ptr
   br label %if.end
 
-if.end:                                           ; preds = %if.then, %entry
-  %shadow_value.0.in = phi i8 [ %2, %if.then ], [ %1, %entry ]
-  switch i8 %shadow_value.0.in, label %sw.default [
+if.end:
+  %swval = phi i8 [ %cond, %if.then ], [ %val, %entry ]
+  switch i8 %swval, label %sw.default [
     i8 -127, label %sw.epilog
     i8 -126, label %sw.bb3
     i8 -125, label %sw.bb4
@@ -35,84 +35,39 @@ if.end:                                           ; preds = %if.then, %entry
     i8 -123, label %sw.bb11
   ]
 
-sw.bb3:                                           ; preds = %if.end
+sw.bb3:
   br label %sw.epilog
 
-sw.bb4:                                           ; preds = %if.end
+sw.bb4:
   br label %sw.epilog
 
-sw.bb5:                                           ; preds = %if.end
+sw.bb5:
   br label %sw.epilog
 
-sw.bb6:                                           ; preds = %if.end
+sw.bb6:
   br label %sw.epilog
 
-sw.bb7:                                           ; preds = %if.end
+sw.bb7:
   br label %sw.epilog
 
-sw.bb8:                                           ; preds = %if.end, %if.end, %if.end
+sw.bb8:
   br label %sw.epilog
 
-sw.bb9:                                           ; preds = %if.end
+sw.bb9:
   br label %sw.epilog
 
-sw.bb10:                                          ; preds = %if.end
+sw.bb10:
   br label %sw.epilog
 
-sw.bb11:                                          ; preds = %if.end
+sw.bb11:
   br label %sw.epilog
 
-sw.default:                                       ; preds = %if.end
+sw.default:
   br label %sw.epilog
 
-sw.epilog:                                        ; preds = %sw.default, %sw.bb11, %sw.bb10, %sw.bb9, %sw.bb8, %sw.bb7, %sw.bb6, %sw.bb5, %sw.bb4, %sw.bb3, %if.end
+sw.epilog:
   br label %exit
 
-if.then.i:                                        ; preds = %sw.epilog
-  br label %exit
-
-for.cond.i:                                       ; preds = %for.inc.i, %if.then.i
-  br label %exit
-
-for.inc.i:                                        ; preds = %for.cond.i
-  br label %exit
-
-if.end.i:                                         ; preds = %for.cond.i, %if.then.i
-  br label %exit
-
-if.end18.thread.i:                                ; preds = %if.end.i
-  br label %5
-
-for.cond8.i:                                      ; preds = %for.inc14.i, %if.end.i
-  br label %exit
-
-for.inc14.i:                                      ; preds = %for.cond8.i
-  br label %exit
-
-if.end18.i:                                       ; preds = %for.cond8.i
-  br label %5
-
-5:                                                ; preds = %if.end18.i, %if.end18.thread.i
-  br label %for.cond25.i
-
-for.cond25.i:                                     ; preds = %for.body29.i, %5
-  br label %exit
-
-for.cond.cleanup27.i:                             ; preds = %for.cond25.i
-  br label %for.cond41.i
-
-for.body29.i:                                     ; preds = %for.cond25.i
-  br label %for.cond25.i
-
-for.cond41.i:                                     ; preds = %for.body45.i, %for.cond.cleanup27.i
-  br label %exit
-
-for.cond.cleanup43.i:                             ; preds = %for.cond41.i
-  br label %exit
-
-for.body45.i:                                     ; preds = %for.cond41.i
-  br label %for.cond41.i
-
 exit:
   ret void
 }

>From 04e76baa6eb57f7f5be150f8259750fedf368562 Mon Sep 17 00:00:00 2001
From: Michal Paszkowski <michal at paszkowski.org>
Date: Sun, 7 Apr 2024 16:12:01 -0700
Subject: [PATCH 3/3] [SPIR-V] Emit SPIR-V generator magic number and version

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
---
 llvm/lib/MC/SPIRVObjectWriter.cpp             |  6 +--
 llvm/lib/Target/SPIRV/SPIRV.td                | 13 -----
 llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp     | 10 ++--
 llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 48 ++++++++++---------
 llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h   | 20 ++++----
 llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp      | 43 ++++++++++++-----
 llvm/lib/Target/SPIRV/SPIRVSubtarget.h        | 12 ++---
 llvm/lib/TargetParser/Triple.cpp              | 14 +++---
 llvm/test/CodeGen/SPIRV/ComparePointers.ll    |  4 +-
 llvm/test/CodeGen/SPIRV/empty-opencl32.ll     |  2 -
 .../SPIRV/exec_mode_float_control_khr.ll      |  4 +-
 .../physical-layout/generator-magic-number.ll |  4 ++
 .../SPIRV/physical-layout/spirv-version.ll    | 16 +++++++
 .../AtomicCompareExchangeExplicit_cl20.ll     |  4 +-
 14 files changed, 115 insertions(+), 85 deletions(-)
 create mode 100644 llvm/test/CodeGen/SPIRV/physical-layout/generator-magic-number.ll
 create mode 100644 llvm/test/CodeGen/SPIRV/physical-layout/spirv-version.ll

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/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 40c3e5f9c6bdab..d6c1b81b56048a 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -76,18 +76,21 @@ 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;
+  
+  unsigned ReqMinVerDec = getSymbolicOperandMinVersion(Category, i);
+  VersionTuple ReqMinVer = VersionTuple(ReqMinVerDec / 10, ReqMinVerDec % 10);
+  unsigned ReqMaxVerDec = getSymbolicOperandMaxVersion(Category, i);
+  VersionTuple ReqMaxVer = VersionTuple(ReqMaxVerDec / 10, ReqMaxVerDec % 10);
+  VersionTuple SPIRVVersion = ST.getSPIRVVersion();
+  bool MinVerOK = ReqMinVer.empty() || 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 +121,9 @@ 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) {
@@ -499,25 +502,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;
   }
 }
@@ -528,7 +531,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 "
@@ -536,14 +539,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"
@@ -621,12 +624,12 @@ 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,
@@ -634,7 +637,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,
@@ -1151,7 +1154,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..bd1530f059b960 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..bfed006308750f 100644
--- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
+++ b/llvm/lib/Target/SPIRV/SPIRVSubtarget.h
@@ -37,9 +37,9 @@ 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;
   std::unique_ptr<SPIRVGlobalRegistry> GR;
@@ -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 624679ff507a7f..9e14863b4f13c3 100644
--- a/llvm/lib/TargetParser/Triple.cpp
+++ b/llvm/lib/TargetParser/Triple.cpp
@@ -527,13 +527,13 @@ static Triple::ArchType parseArch(StringRef ArchName) {
     .Case("m68k", Triple::m68k)
     .Case("msp430", Triple::msp430)
     .Cases("mips", "mipseb", "mipsallegrex", "mipsisa32r6",
-           "mipsr6", Triple::mips)
+            "mipsr6", Triple::mips)
     .Cases("mipsel", "mipsallegrexel", "mipsisa32r6el", "mipsr6el",
-           Triple::mipsel)
+            Triple::mipsel)
     .Cases("mips64", "mips64eb", "mipsn32", "mipsisa64r6",
-           "mips64r6", "mipsn32r6", Triple::mips64)
+            "mips64r6", "mipsn32r6", Triple::mips64)
     .Cases("mips64el", "mipsn32el", "mipsisa64r6el", "mips64r6el",
-           "mipsn32r6el", Triple::mips64el)
+            "mipsn32r6el", Triple::mips64el)
     .Case("r600", Triple::r600)
     .Case("amdgcn", Triple::amdgcn)
     .Case("riscv32", Triple::riscv32)
@@ -558,9 +558,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