diff --git a/CMakeLists.txt b/CMakeLists.txt
index 786ea78a9..e8eae2484 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -144,7 +144,9 @@ if (NOT SPIRV_TOOLS_FOUND)
   endif()
 endif()
 
-if (NOT SPIRV_TOOLS_FOUND)
+option(LLVM_SPIRV_ENABLE_LIBSPIRV_DIS "Enable --spirv-tools-dis support.")
+
+if (NOT SPIRV_TOOLS_FOUND AND LLVM_SPIRV_ENABLE_LIBSPIRV_DIS)
   message(STATUS "SPIRV-Tools not found; project will be built without "
           "--spirv-tools-dis support.")
 endif()
diff --git a/README.md b/README.md
index 6b76f2cd5..ba39dc7c1 100644
--- a/README.md
+++ b/README.md
@@ -109,6 +109,7 @@ Building clang from sources takes time and resources and it can be avoided:
 ### Build with SPIRV-Tools
 
 The translator can use [SPIRV-Tools](https://github.com/KhronosGroup/SPIRV-Tools) to generate assembly with widely adopted syntax.
+This feature can be enabled by passing `-DLLVM_SPIRV_ENABLE_LIBSPIRV_DIS=ON` option.
 If SPIRV-Tools have been installed prior to the build it will be detected and
 used automatically. However it is also possible to enable use of SPIRV-Tools
 from a custom location using the following instructions:
diff --git a/include/LLVMSPIRVExtensions.inc b/include/LLVMSPIRVExtensions.inc
index f8edd290f..1a256b912 100644
--- a/include/LLVMSPIRVExtensions.inc
+++ b/include/LLVMSPIRVExtensions.inc
@@ -46,6 +46,7 @@ EXT(SPV_INTEL_fpga_cluster_attributes)
 EXT(SPV_INTEL_loop_fuse)
 EXT(SPV_INTEL_long_constant_composite) // TODO: rename to
                                        // SPV_INTEL_long_composites later
+EXT(SPV_EXT_optnone)
 EXT(SPV_INTEL_optnone)
 EXT(SPV_INTEL_fpga_dsp_control)
 EXT(SPV_INTEL_memory_access_aliasing)
diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h
index 2698067aa..062609ff5 100644
--- a/lib/SPIRV/SPIRVInternal.h
+++ b/lib/SPIRV/SPIRVInternal.h
@@ -251,7 +251,7 @@ inline void
 SPIRVMap<Attribute::AttrKind, SPIRVFunctionControlMaskKind>::init() {
   add(Attribute::AlwaysInline, FunctionControlInlineMask);
   add(Attribute::NoInline, FunctionControlDontInlineMask);
-  add(Attribute::OptimizeNone, internal::FunctionControlOptNoneINTELMask);
+  add(Attribute::OptimizeNone, FunctionControlOptNoneEXTMask);
 }
 typedef SPIRVMap<Attribute::AttrKind, SPIRVFunctionControlMaskKind>
     SPIRSPIRVFuncCtlMaskMap;
diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp
index 548bd469f..37275ac6d 100644
--- a/lib/SPIRV/SPIRVReader.cpp
+++ b/lib/SPIRV/SPIRVReader.cpp
@@ -1099,8 +1099,43 @@ Value *SPIRVToLLVM::transConvertInst(SPIRVValue *BV, Function *F,
   case OpBitcast:
     if (Src->getType()->isPointerTy() && Dst->isPointerTy() &&
         Src->getType()->getPointerAddressSpace() != Dst->getPointerAddressSpace() &&
-        M->getTargetTriple() == "amdgcn-amd-amdhsa")
+        M->getTargetTriple() == "amdgcn-amd-amdhsa") {
       CO = Instruction::AddrSpaceCast;
+    } else {
+      // OpBitcast need to be handled as a special-case when the source is a
+      // pointer and the destination is not a pointer, and where the source is not
+      // a pointer and the destination is a pointer. This is supported by the
+      // SPIR-V bitcast, but not by the LLVM bitcast.
+      CO = Instruction::BitCast;
+      if (Src->getType()->isPointerTy() && !Dst->isPointerTy()) {
+        if (auto *DstVecTy = dyn_cast<FixedVectorType>(Dst)) {
+          unsigned TotalBitWidth =
+              DstVecTy->getElementType()->getIntegerBitWidth() *
+              DstVecTy->getNumElements();
+          auto *IntTy = Type::getIntNTy(BB->getContext(), TotalBitWidth);
+          if (BB) {
+            Src = CastInst::CreatePointerCast(Src, IntTy, "", BB);
+          } else {
+            Src = ConstantExpr::getPointerCast(dyn_cast<Constant>(Src), IntTy);
+          }
+        } else {
+          CO = Instruction::PtrToInt;
+        }
+      } else if (!Src->getType()->isPointerTy() && Dst->isPointerTy()) {
+        if (auto *SrcVecTy = dyn_cast<FixedVectorType>(Src->getType())) {
+          unsigned TotalBitWidth =
+              SrcVecTy->getElementType()->getIntegerBitWidth() *
+              SrcVecTy->getNumElements();
+          auto *IntTy = Type::getIntNTy(BB->getContext(), TotalBitWidth);
+          if (BB) {
+            Src = CastInst::Create(Instruction::BitCast, Src, IntTy, "", BB);
+          } else {
+            Src = ConstantExpr::getBitCast(dyn_cast<Constant>(Src), IntTy);
+          }
+        }
+        CO = Instruction::IntToPtr;
+      }
+    }
     break;
   default:
     CO = static_cast<CastInst::CastOps>(OpCodeMap::rmap(BC->getOpCode()));
@@ -3838,6 +3873,7 @@ Instruction *SPIRVToLLVM::transSPIRVBuiltinFromInst(SPIRVInstruction *BI,
   case internal::OpJointMatrixLoadINTEL:
   case OpCooperativeMatrixLoadKHR:
   case internal::OpCooperativeMatrixLoadCheckedINTEL:
+  case internal::OpCooperativeMatrixLoadOffsetINTEL:
   case internal::OpTaskSequenceCreateINTEL:
   case internal::OpConvertHandleToImageINTEL:
   case internal::OpConvertHandleToSampledImageINTEL:
@@ -5545,36 +5581,53 @@ void SPIRVToLLVM::transAuxDataInst(SPIRVExtInst *BC) {
     return;
   auto Args = BC->getArguments();
   // Args 0 and 1 are common between attributes and metadata.
-  // 0 is the function, 1 is the name of the attribute/metadata as a string
-  auto *SpvFcn = BC->getModule()->getValue(Args[0]);
-  auto *F = static_cast<Function *>(getTranslatedValue(SpvFcn));
-  assert(F && "Function should already have been translated!");
+  // 0 is the global object, 1 is the name of the attribute/metadata as a string
+  auto *Arg0 = BC->getModule()->getValue(Args[0]);
+  auto *GO = cast<GlobalObject>(getTranslatedValue(Arg0));
+  auto *F = dyn_cast<Function>(GO);
+  auto *GV = dyn_cast<GlobalVariable>(GO);
+  assert((F || GV) && "Value should already have been translated!");
   auto AttrOrMDName = BC->getModule()->get<SPIRVString>(Args[1])->getStr();
   switch (BC->getExtOp()) {
-  case NonSemanticAuxData::FunctionAttribute: {
+  case NonSemanticAuxData::FunctionAttribute:
+  case NonSemanticAuxData::GlobalVariableAttribute: {
     assert(Args.size() < 4 && "Unexpected FunctionAttribute Args");
     // If this attr was specially handled and added elsewhere, skip it.
     Attribute::AttrKind AsKind = Attribute::getAttrKindFromName(AttrOrMDName);
-    if (AsKind != Attribute::None && F->hasFnAttribute(AsKind))
-      return;
-    if (AsKind == Attribute::None && F->hasFnAttribute(AttrOrMDName))
-      return;
+    if (AsKind != Attribute::None)
+      if ((F && F->hasFnAttribute(AsKind)) || (GV && GV->hasAttribute(AsKind)))
+        return;
+    if (AsKind == Attribute::None)
+      if ((F && F->hasFnAttribute(AttrOrMDName)) ||
+          (GV && GV->hasAttribute(AttrOrMDName)))
+        return;
     // For attributes, arg 2 is the attribute value as a string, which may not
     // exist.
     if (Args.size() == 3) {
       auto AttrValue = BC->getModule()->get<SPIRVString>(Args[2])->getStr();
-      F->addFnAttr(AttrOrMDName, AttrValue);
-    } else {
-      if (AsKind != Attribute::None)
-        F->addFnAttr(AsKind);
+      if (F)
+        F->addFnAttr(AttrOrMDName, AttrValue);
       else
-        F->addFnAttr(AttrOrMDName);
+        GV->addAttribute(AttrOrMDName, AttrValue);
+    } else {
+      if (AsKind != Attribute::None) {
+        if (F)
+          F->addFnAttr(AsKind);
+        else
+          GV->addAttribute(AsKind);
+      } else {
+        if (F)
+          F->addFnAttr(AttrOrMDName);
+        else
+          GV->addAttribute(AttrOrMDName);
+      }
     }
     break;
   }
-  case NonSemanticAuxData::FunctionMetadata: {
+  case NonSemanticAuxData::FunctionMetadata:
+  case NonSemanticAuxData::GlobalVariableMetadata: {
     // If this metadata was specially handled and added elsewhere, skip it.
-    if (F->hasMetadata(AttrOrMDName))
+    if (GO->hasMetadata(AttrOrMDName))
       return;
     SmallVector<Metadata *> MetadataArgs;
     // Process the metadata values.
@@ -5584,14 +5637,14 @@ void SPIRVToLLVM::transAuxDataInst(SPIRVExtInst *BC) {
       if (Arg->getOpCode() == OpString) {
         auto *ArgAsStr = static_cast<SPIRVString *>(Arg);
         MetadataArgs.push_back(
-            MDString::get(F->getContext(), ArgAsStr->getStr()));
+            MDString::get(GO->getContext(), ArgAsStr->getStr()));
       } else {
         auto *ArgAsVal = static_cast<SPIRVValue *>(Arg);
-        auto *TranslatedMD = transValue(ArgAsVal, F, nullptr);
+        auto *TranslatedMD = transValue(ArgAsVal, nullptr, nullptr);
         MetadataArgs.push_back(ValueAsMetadata::get(TranslatedMD));
       }
     }
-    F->setMetadata(AttrOrMDName, MDNode::get(*Context, MetadataArgs));
+    GO->setMetadata(AttrOrMDName, MDNode::get(*Context, MetadataArgs));
     break;
   }
   default:
diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp
index 53dbd2355..123cdd446 100644
--- a/lib/SPIRV/SPIRVWriter.cpp
+++ b/lib/SPIRV/SPIRVWriter.cpp
@@ -1239,23 +1239,27 @@ void LLVMToSPIRVBase::transFunctionMetadataAsUserSemanticDecoration(
   }
 }
 
-void LLVMToSPIRVBase::transAuxDataInst(SPIRVFunction *BF, Function *F) {
-  auto *BM = BF->getModule();
+void LLVMToSPIRVBase::transAuxDataInst(SPIRVValue *BV, Value *V) {
+  auto *GO = cast<GlobalObject>(V);
+  auto *F = dyn_cast<Function>(GO);
+  auto *GV = dyn_cast<GlobalVariable>(GO);
+  assert((F || GV) && "Invalid value type");
+  auto *BM = BV->getModule();
   if (!BM->preserveAuxData())
     return;
   if (!BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_6))
     BM->addExtension(SPIRV::ExtensionID::SPV_KHR_non_semantic_info);
   else
     BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_6);
-  const auto &FnAttrs = F->getAttributes().getFnAttrs();
-  for (const auto &Attr : FnAttrs) {
+  const auto &Attrs = F ? F->getAttributes().getFnAttrs() : GV->getAttributes();
+  for (const auto &Attr : Attrs) {
     std::vector<SPIRVWord> Ops;
-    Ops.push_back(BF->getId());
+    Ops.push_back(BV->getId());
     if (Attr.isStringAttribute()) {
       // Format for String attributes is:
-      // NonSemanticAuxDataFunctionAttribute Fcn AttrName AttrValue
+      // NonSemanticAuxData*Attribute ValueName AttrName AttrValue
       // or, if no value:
-      // NonSemanticAuxDataFunctionAttribute Fcn AttrName
+      // NonSemanticAuxData*Attribute ValueName AttrName
       //
       // AttrName and AttrValue are always Strings
       StringRef AttrKind = Attr.getKindAsString();
@@ -1268,19 +1272,20 @@ void LLVMToSPIRVBase::transAuxDataInst(SPIRVFunction *BF, Function *F) {
       }
     } else {
       // Format for other types is:
-      // NonSemanticAuxDataFunctionAttribute Fcn AttrStr
+      // NonSemanticAuxData*Attribute ValueName AttrStr
       // AttrStr is always a String.
       std::string AttrStr = Attr.getAsString();
       auto *AttrSpvString = BM->getString(AttrStr);
       Ops.push_back(AttrSpvString->getId());
     }
-    BM->addAuxData(NonSemanticAuxData::FunctionAttribute,
-                   transType(Type::getVoidTy(F->getContext())), Ops);
+    BM->addAuxData(F ? NonSemanticAuxData::FunctionAttribute
+                     : NonSemanticAuxData::GlobalVariableAttribute,
+                   transType(Type::getVoidTy(V->getContext())), Ops);
   }
   SmallVector<std::pair<unsigned, MDNode *>> AllMD;
   SmallVector<StringRef> MDNames;
-  F->getContext().getMDKindNames(MDNames);
-  F->getAllMetadata(AllMD);
+  V->getContext().getMDKindNames(MDNames);
+  GO->getAllMetadata(AllMD);
   for (const auto &MD : AllMD) {
     std::string MDName = MDNames[MD.first].str();
 
@@ -1293,11 +1298,11 @@ void LLVMToSPIRVBase::transAuxDataInst(SPIRVFunction *BF, Function *F) {
       continue;
 
     // Format for metadata is:
-    // NonSemanticAuxDataFunctionMetadata Fcn MDName MDVals...
+    // NonSemanticAuxData*Metadata ValueName MDName MDVals...
     // MDName is always a String, MDVals have different types as explained
     // below. Also note this instruction has a variable number of operands
     std::vector<SPIRVWord> Ops;
-    Ops.push_back(BF->getId());
+    Ops.push_back(BV->getId());
     Ops.push_back(BM->getString(MDName)->getId());
     for (unsigned int OpIdx = 0; OpIdx < MD.second->getNumOperands(); OpIdx++) {
       const auto &CurOp = MD.second->getOperand(OpIdx);
@@ -1313,8 +1318,9 @@ void LLVMToSPIRVBase::transAuxDataInst(SPIRVFunction *BF, Function *F) {
         assert(false && "Unsupported metadata type");
       }
     }
-    BM->addAuxData(NonSemanticAuxData::FunctionMetadata,
-                   transType(Type::getVoidTy(F->getContext())), Ops);
+    BM->addAuxData(F ? NonSemanticAuxData::FunctionMetadata
+                     : NonSemanticAuxData::GlobalVariableMetadata,
+                   transType(Type::getVoidTy(V->getContext())), Ops);
   }
 }
 
@@ -2045,6 +2051,7 @@ LLVMToSPIRVBase::transValueWithoutDecoration(Value *V, SPIRVBasicBlock *BB,
     if (ST && ST->hasName() && isSPIRVConstantName(ST->getName())) {
       auto *BV = transConstant(Init);
       assert(BV);
+      transAuxDataInst(BV, V);
       return mapValue(V, BV);
     }
     if (isa_and_nonnull<ConstantExpr>(Init)) {
@@ -2148,6 +2155,8 @@ LLVMToSPIRVBase::transValueWithoutDecoration(Value *V, SPIRVBasicBlock *BB,
             GV->getAttribute(kVCMetadata::VCSingleElementVector), BVar);
     }
 
+    transAuxDataInst(BVar, V);
+
     mapValue(V, BVar);
     spv::BuiltIn Builtin = spv::BuiltInPosition;
     if (!GV->hasName() || !getSPIRVBuiltin(GV->getName().str(), Builtin))
@@ -4356,7 +4365,7 @@ SPIRVValue *LLVMToSPIRVBase::transIntrinsicInst(IntrinsicInst *II,
           static_cast<std::vector<SPIRVValue *>::size_type>(VecSize);
       auto *ElemOne = BM->addConstant(ElemTy, 1);
       auto *ElemZero = BM->addConstant(ElemTy, 0);
-      auto *ElemMinusOne = BM->addConstant(ElemTy, MinusOneValue);
+      auto *ElemMinusOne = BM->addConstant(ElemTy, std::move(MinusOneValue));
       std::vector<SPIRVValue *> ElemsOne(ElemCount, ElemOne);
       std::vector<SPIRVValue *> ElemsZero(ElemCount, ElemZero);
       std::vector<SPIRVValue *> ElemsMinusOne(ElemCount, ElemMinusOne);
@@ -5575,10 +5584,15 @@ SPIRVWord LLVMToSPIRVBase::transFunctionControlMask(Function *F) {
       [&](Attribute::AttrKind Attr, SPIRVFunctionControlMaskKind Mask) {
         if (F->hasFnAttribute(Attr)) {
           if (Attr == Attribute::OptimizeNone) {
-            if (!BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_optnone))
+            if (BM->isAllowedToUseExtension(ExtensionID::SPV_EXT_optnone)) {
+              BM->addExtension(ExtensionID::SPV_EXT_optnone);
+              BM->addCapability(CapabilityOptNoneEXT);
+            } else if (BM->isAllowedToUseExtension(
+                           ExtensionID::SPV_INTEL_optnone)) {
+              BM->addExtension(ExtensionID::SPV_INTEL_optnone);
+              BM->addCapability(CapabilityOptNoneINTEL);
+            } else
               return;
-            BM->addExtension(ExtensionID::SPV_INTEL_optnone);
-            BM->addCapability(internal::CapabilityOptNoneINTEL);
           }
           FCM |= Mask;
         }
@@ -7071,6 +7085,7 @@ bool runSpirvBackend(Module *M, std::string &Result, std::string &ErrMsg,
       SPIRV::ExtensionID::SPV_INTEL_cache_controls,
       SPIRV::ExtensionID::SPV_INTEL_global_variable_fpga_decorations,
       SPIRV::ExtensionID::SPV_INTEL_global_variable_host_access,
+      SPIRV::ExtensionID::SPV_EXT_optnone,
       SPIRV::ExtensionID::SPV_INTEL_optnone,
       SPIRV::ExtensionID::SPV_INTEL_usm_storage_classes,
       SPIRV::ExtensionID::SPV_INTEL_subgroups,
diff --git a/lib/SPIRV/SPIRVWriter.h b/lib/SPIRV/SPIRVWriter.h
index 244d8dc40..46399f3cf 100644
--- a/lib/SPIRV/SPIRVWriter.h
+++ b/lib/SPIRV/SPIRVWriter.h
@@ -138,7 +138,7 @@ class LLVMToSPIRVBase : protected BuiltinCallHelper {
   void transFunctionMetadataAsExecutionMode(SPIRVFunction *BF, Function *F);
   void transFunctionMetadataAsUserSemanticDecoration(SPIRVFunction *BF,
                                                      Function *F);
-  void transAuxDataInst(SPIRVFunction *BF, Function *F);
+  void transAuxDataInst(SPIRVValue *BV, Value *V);
 
   bool transGlobalVariables();
 
diff --git a/lib/SPIRV/libSPIRV/NonSemantic.AuxData.h b/lib/SPIRV/libSPIRV/NonSemantic.AuxData.h
index 240734afd..aa01871ca 100644
--- a/lib/SPIRV/libSPIRV/NonSemantic.AuxData.h
+++ b/lib/SPIRV/libSPIRV/NonSemantic.AuxData.h
@@ -28,6 +28,7 @@ namespace NonSemanticAuxData {
 enum Instruction {
   FunctionMetadata = 0,
   FunctionAttribute = 1,
-  PreserveCount = 2
+  GlobalVariableMetadata = 2,
+  GlobalVariableAttribute = 3
 };
 } // namespace NonSemanticAuxData
diff --git a/lib/SPIRV/libSPIRV/SPIRVEnum.h b/lib/SPIRV/libSPIRV/SPIRVEnum.h
index 4c318be39..ce2084157 100644
--- a/lib/SPIRV/libSPIRV/SPIRVEnum.h
+++ b/lib/SPIRV/libSPIRV/SPIRVEnum.h
@@ -221,6 +221,8 @@ template <> inline void SPIRVMap<SPIRVCapabilityKind, SPIRVCapVec>::init() {
                {CapabilityCooperativeMatrixKHR});
   ADD_VEC_INIT(internal::CapabilityCooperativeMatrixCheckedInstructionsINTEL,
                {CapabilityCooperativeMatrixKHR});
+  ADD_VEC_INIT(internal::CapabilityCooperativeMatrixOffsetInstructionsINTEL,
+               {CapabilityCooperativeMatrixKHR});
 }
 
 template <> inline void SPIRVMap<SPIRVExecutionModelKind, SPIRVCapVec>::init() {
diff --git a/lib/SPIRV/libSPIRV/SPIRVExtInst.h b/lib/SPIRV/libSPIRV/SPIRVExtInst.h
index 693d8dab6..bbce66fc3 100644
--- a/lib/SPIRV/libSPIRV/SPIRVExtInst.h
+++ b/lib/SPIRV/libSPIRV/SPIRVExtInst.h
@@ -278,6 +278,10 @@ inline void SPIRVMap<NonSemanticAuxDataOpKind, std::string>::init() {
       "NonSemanticAuxDataFunctionMetadata");
   add(NonSemanticAuxData::FunctionAttribute,
       "NonSemanticAuxDataFunctionAttribute");
+  add(NonSemanticAuxData::GlobalVariableMetadata,
+      "NonSemanticAuxDataGlobalVariableMetadata");
+  add(NonSemanticAuxData::GlobalVariableAttribute,
+      "NonSemanticAuxDataGlobalVariableAttribute");
 }
 SPIRV_DEF_NAMEMAP(NonSemanticAuxDataOpKind, NonSemanticAuxDataOpMap)
 
diff --git a/lib/SPIRV/libSPIRV/SPIRVInstruction.h b/lib/SPIRV/libSPIRV/SPIRVInstruction.h
index fe10b0a07..ba4b20dd4 100644
--- a/lib/SPIRV/libSPIRV/SPIRVInstruction.h
+++ b/lib/SPIRV/libSPIRV/SPIRVInstruction.h
@@ -3740,6 +3740,26 @@ _SPIRV_OP(CooperativeMatrixStoreChecked, false, 8, true, 8)
 _SPIRV_OP(CooperativeMatrixConstructChecked, true, 8)
 #undef _SPIRV_OP
 
+class SPIRVCooperativeMatrixOffsetInstructionsINTELInstBase
+    : public SPIRVInstTemplateBase {
+protected:
+  std::optional<ExtensionID> getRequiredExtension() const override {
+    return ExtensionID::SPV_INTEL_joint_matrix;
+  }
+  SPIRVCapVec getRequiredCapability() const override {
+    return getVec(internal::CapabilityCooperativeMatrixOffsetInstructionsINTEL);
+  }
+};
+
+#define _SPIRV_OP(x, ...)                                                      \
+  typedef SPIRVInstTemplate<                                                   \
+      SPIRVCooperativeMatrixOffsetInstructionsINTELInstBase,                   \
+      internal::Op##x##INTEL, __VA_ARGS__>                                     \
+      SPIRV##x##INTEL;
+_SPIRV_OP(CooperativeMatrixLoadOffset, true, 8, true, 5)
+_SPIRV_OP(CooperativeMatrixStoreOffset, false, 7, true, 6)
+#undef _SPIRV_OP
+
 class SPIRVCooperativeMatrixInvocationInstructionsINTELInstBase
     : public SPIRVInstTemplateBase {
 protected:
diff --git a/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h b/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h
index 34b73691c..de46aa731 100644
--- a/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h
+++ b/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h
@@ -294,7 +294,7 @@ inline bool isValidFunctionControlMask(SPIRVWord Mask) {
   ValidMask |= FunctionControlDontInlineMask;
   ValidMask |= FunctionControlPureMask;
   ValidMask |= FunctionControlConstMask;
-  ValidMask |= internal::FunctionControlOptNoneINTELMask;
+  ValidMask |= FunctionControlOptNoneEXTMask;
 
   return (Mask & ~ValidMask) == 0;
 }
diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
index ec442aacf..8ec0f2b06 100644
--- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
+++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h
@@ -623,7 +623,7 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
   add(CapabilityAtomicFloat32AddEXT, "AtomicFloat32AddEXT");
   add(CapabilityAtomicFloat64AddEXT, "AtomicFloat64AddEXT");
   add(CapabilityLongCompositesINTEL, "LongCompositesINTEL");
-  add(CapabilityOptNoneINTEL, "OptNoneINTEL");
+  add(CapabilityOptNoneEXT, "OptNoneEXT");
   add(CapabilityAtomicFloat16AddEXT, "AtomicFloat16AddEXT");
   add(CapabilityDebugInfoModuleINTEL, "DebugInfoModuleINTEL");
   add(CapabilitySplitBarrierINTEL, "SplitBarrierINTEL");
@@ -642,7 +642,6 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
   add(CapabilityRegisterLimitsINTEL, "RegisterLimitsINTEL");
   // From spirv_internal.hpp
   add(internal::CapabilityFastCompositeINTEL, "FastCompositeINTEL");
-  add(internal::CapabilityOptNoneINTEL, "OptNoneINTEL");
   add(internal::CapabilityTokenTypeINTEL, "TokenTypeINTEL");
   add(internal::CapabilityFPArithmeticFenceINTEL, "FPArithmeticFenceINTEL");
   add(internal::CapabilityBfloat16ConversionINTEL, "Bfloat16ConversionINTEL");
@@ -671,6 +670,8 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
       "CooperativeMatrixInvocationInstructionsINTEL");
   add(internal::CapabilityCooperativeMatrixCheckedInstructionsINTEL,
       "CooperativeMatrixCheckedInstructionsINTEL");
+  add(internal::CapabilityCooperativeMatrixOffsetInstructionsINTEL,
+      "CooperativeMatrixOffsetInstructionsINTEL");
   add(internal::CapabilitySubgroupRequirementsINTEL,
       "SubgroupRequirementsINTEL");
   add(internal::CapabilityTaskSequenceINTEL, "TaskSequenceINTEL");
diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCodeEnumInternal.h b/lib/SPIRV/libSPIRV/SPIRVOpCodeEnumInternal.h
index f29008037..5317be2f8 100644
--- a/lib/SPIRV/libSPIRV/SPIRVOpCodeEnumInternal.h
+++ b/lib/SPIRV/libSPIRV/SPIRVOpCodeEnumInternal.h
@@ -24,6 +24,10 @@ _SPIRV_OP_INTERNAL(CooperativeMatrixStoreCheckedINTEL,
                    internal::OpCooperativeMatrixStoreCheckedINTEL)
 _SPIRV_OP_INTERNAL(CooperativeMatrixConstructCheckedINTEL,
                    internal::OpCooperativeMatrixConstructCheckedINTEL)
+_SPIRV_OP_INTERNAL(CooperativeMatrixLoadOffsetINTEL,
+                   internal::OpCooperativeMatrixLoadOffsetINTEL)
+_SPIRV_OP_INTERNAL(CooperativeMatrixStoreOffsetINTEL,
+                   internal::OpCooperativeMatrixStoreOffsetINTEL)
 _SPIRV_OP_INTERNAL(CooperativeMatrixApplyFunctionINTEL,
                    internal::OpCooperativeMatrixApplyFunctionINTEL)
 _SPIRV_OP_INTERNAL(ComplexFMulINTEL, internal::ComplexFMulINTEL)
diff --git a/lib/SPIRV/libSPIRV/spirv_internal.hpp b/lib/SPIRV/libSPIRV/spirv_internal.hpp
index cdec3a959..d796e7e0d 100644
--- a/lib/SPIRV/libSPIRV/spirv_internal.hpp
+++ b/lib/SPIRV/libSPIRV/spirv_internal.hpp
@@ -77,6 +77,8 @@ enum InternalOp {
   IOpCooperativeMatrixLoadCheckedINTEL = 6193,
   IOpCooperativeMatrixStoreCheckedINTEL = 6194,
   IOpCooperativeMatrixConstructCheckedINTEL = 6195,
+  IOpCooperativeMatrixLoadOffsetINTEL = 6239,
+  IOpCooperativeMatrixStoreOffsetINTEL = 6240,
   IOpJointMatrixWorkItemLengthINTEL = 6410,
   IOpTypeTaskSequenceINTEL = 6199,
   IOpComplexFMulINTEL = 6415,
@@ -105,7 +107,6 @@ enum InternalDecoration {
 
 enum InternalCapability {
   ICapFastCompositeINTEL = 6093,
-  ICapOptNoneINTEL = 6094,
   ICapTokenTypeINTEL = 6112,
   ICapBfloat16ConversionINTEL = 6115,
   ICapabilityJointMatrixINTEL = 6118,
@@ -114,6 +115,7 @@ enum InternalCapability {
   ICapGlobalVariableDecorationsINTEL = 6146,
   ICapabilityTaskSequenceINTEL = 6162,
   ICapabilityCooperativeMatrixCheckedInstructionsINTEL = 6192,
+  ICapabilityCooperativeMatrixOffsetInstructionsINTEL = 6238,
   ICapabilityCooperativeMatrixPrefetchINTEL = 6411,
   ICapabilityComplexFloatMulDivINTEL = 6414,
   ICapabilityTensorFloat32RoundingINTEL = 6425,
@@ -129,8 +131,6 @@ enum InternalCapability {
   ICapabilityBindlessImagesINTEL = 6528
 };
 
-enum InternalFunctionControlMask { IFunctionControlOptNoneINTELMask = 0x10000 };
-
 enum InternalExecutionMode {
   IExecModeFastCompositeKernelINTEL = 6088,
   IExecModeNamedSubgroupSizeINTEL = 6446,
@@ -187,6 +187,10 @@ _SPIRV_OP(Op, CooperativeMatrixLoadCheckedINTEL)
 _SPIRV_OP(Op, CooperativeMatrixStoreCheckedINTEL)
 _SPIRV_OP(Op, CooperativeMatrixConstructCheckedINTEL)
 
+_SPIRV_OP(Capability, CooperativeMatrixOffsetInstructionsINTEL)
+_SPIRV_OP(Op, CooperativeMatrixLoadOffsetINTEL)
+_SPIRV_OP(Op, CooperativeMatrixStoreOffsetINTEL)
+
 _SPIRV_OP(Capability, CooperativeMatrixInvocationInstructionsINTEL)
 _SPIRV_OP(Op, CooperativeMatrixApplyFunctionINTEL)
 
@@ -284,8 +288,6 @@ constexpr Decoration DecorationArgumentAttributeINTEL =
 
 constexpr Capability CapabilityFastCompositeINTEL =
     static_cast<Capability>(ICapFastCompositeINTEL);
-constexpr Capability CapabilityOptNoneINTEL =
-    static_cast<Capability>(ICapOptNoneINTEL);
 constexpr Capability CapabilityTokenTypeINTEL =
     static_cast<Capability>(ICapTokenTypeINTEL);
 constexpr Capability CapabilityFPArithmeticFenceINTEL =
@@ -295,9 +297,6 @@ constexpr Capability CapabilityBfloat16ConversionINTEL =
 constexpr Capability CapabilityGlobalVariableDecorationsINTEL =
     static_cast<Capability>(ICapGlobalVariableDecorationsINTEL);
 
-constexpr FunctionControlMask FunctionControlOptNoneINTELMask =
-    static_cast<FunctionControlMask>(IFunctionControlOptNoneINTELMask);
-
 constexpr ExecutionMode ExecutionModeFastCompositeKernelINTEL =
     static_cast<ExecutionMode>(IExecModeFastCompositeKernelINTEL);
 
diff --git a/spirv-headers-tag.conf b/spirv-headers-tag.conf
index dbca53dc3..066641d4b 100644
--- a/spirv-headers-tag.conf
+++ b/spirv-headers-tag.conf
@@ -1 +1 @@
-efb6b4099ddb8fa60f62956dee592c4b94ec6a49
+3f17b2af6784bfa2c5aa5dbb8e0e74a607dd8b3b
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 0e5f7bf40..747cd1b30 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -31,6 +31,9 @@ if(SPIRV_TOOLS_FOUND AND NOT SPIRV-Tools-tools_FOUND)
   endif()
 
   set(SPIRV_TOOLS_BINDIR "${SPIRV_TOOLS_PREFIX}/bin")
+  if (LLVM_SPIRV_ENABLE_LIBSPIRV_DIS)
+    set(SPIRV_ENABLE_LIBSPIRV_DIS ON)
+  endif()
 elseif(SPIRV-Tools-tools_FOUND)
   # we found SPIRV-Tools via cmake targets
 
diff --git a/test/DebugInfo/X86/header.ll b/test/DebugInfo/X86/header.ll
index 9ec6647fd..d98e95750 100644
--- a/test/DebugInfo/X86/header.ll
+++ b/test/DebugInfo/X86/header.ll
@@ -16,8 +16,8 @@ target triple = "spir64-unknown-unknown"
 
 ; Test that we don't pollute the start of the file with debug sections
 
-; CHECK:       .text
-; CHECK-NEXT: .file	"<stdin>"
+; CHECK: .file	"<stdin>"
+; CHECK-NEXT: .text
 ; CHECK-NEXT: .globl	f
 ; CHECK-NEXT: .p2align	4
 ; CHECK-NEXT: .type	f,@function
diff --git a/test/OpBitcast_ptr_scalar.spvasm b/test/OpBitcast_ptr_scalar.spvasm
new file mode 100644
index 000000000..4e5b417a3
--- /dev/null
+++ b/test/OpBitcast_ptr_scalar.spvasm
@@ -0,0 +1,31 @@
+; Check support of OpBitcast with pointer operands
+; Converts to scalar integers, which is supported by all SPIR-V versions
+
+; REQUIRES: spirv-as
+; RUN: spirv-as --target-env spv1.0 -o %t.spv %s
+; RUN: spirv-val %t.spv
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc
+; RUN: FileCheck < %t.rev.ll %s --check-prefix=CHECK-LLVM
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "test"
+       %uint = OpTypeInt 32 0
+      %ulong = OpTypeInt 64 0
+       %void = OpTypeVoid
+   %pptr_int = OpTypePointer Function %uint
+ %kernel_sig = OpTypeFunction %void
+     %kernel = OpFunction %void None %kernel_sig
+      %entry = OpLabel
+     %srcptr = OpVariable %pptr_int Function
+     %dstint = OpBitcast %ulong %srcptr
+     %dstptr = OpBitcast %pptr_int %dstint
+               OpReturn
+               OpFunctionEnd
+
+
+; CHECK-LLVM: [[SRCPTR:%[a-z0-9.]+]] = alloca i32, align 4
+; CHECK-LLVM: [[DSTINT:%[a-z0-9.]+]] = ptrtoint ptr [[SRCPTR]] to i64
+; CHECK-LLVM: [[DSTPTR:%[a-z0-9.]+]] = inttoptr i64 [[DSTINT]] to ptr
diff --git a/test/OpBitcast_ptr_vector.spvasm b/test/OpBitcast_ptr_vector.spvasm
new file mode 100644
index 000000000..852b41ce7
--- /dev/null
+++ b/test/OpBitcast_ptr_vector.spvasm
@@ -0,0 +1,33 @@
+; Check support of OpBitcast with pointer operands
+; Converts to vectors of integers, which is supported by SPIR-V 1.5
+
+; REQUIRES: spirv-as
+; RUN: spirv-as --target-env spv1.5 -o %t.spv %s
+; RUN: spirv-val %t.spv
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc
+; RUN: FileCheck < %t.rev.ll %s --check-prefix=CHECK-LLVM
+               OpCapability Addresses
+               OpCapability Kernel
+               OpCapability Int64
+               OpMemoryModel Physical64 OpenCL
+               OpEntryPoint Kernel %kernel "test"
+       %uint = OpTypeInt 32 0
+      %uint2 = OpTypeVector %uint 2
+       %void = OpTypeVoid
+   %pptr_int = OpTypePointer Function %uint
+ %kernel_sig = OpTypeFunction %void
+     %kernel = OpFunction %void None %kernel_sig
+      %entry = OpLabel
+     %srcptr = OpVariable %pptr_int Function
+    %dstint2 = OpBitcast %uint2 %srcptr
+     %dstptr = OpBitcast %pptr_int %dstint2
+               OpReturn
+               OpFunctionEnd
+
+
+; CHECK-LLVM: [[SRCPTR:%[a-z0-9.]+]] = alloca i32, align 4
+; CHECK-LLVM: [[TMPLONG0:%[a-z0-9.]+]] = ptrtoint ptr [[SRCPTR]] to i64
+; CHECK-LLVM: [[DSTINT2:%[a-z0-9.]+]] = bitcast i64 [[TMPLONG0]] to <2 x i32>
+; CHECK-LLVM: [[TMPLONG1:%[a-z0-9.]+]] = bitcast <2 x i32> [[DSTINT2]] to i64
+; CHECK-LLVM: [[DSTPTR:%[a-z0-9.]+]] = inttoptr i64 [[TMPLONG1]] to ptr
diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_load_store_offset.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_load_store_offset.ll
new file mode 100644
index 000000000..261965c4e
--- /dev/null
+++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_load_store_offset.ll
@@ -0,0 +1,149 @@
+; This is an adapted copy of test/extensions/KHR/SPV_KHR_cooperative_matrix/cooperative_matrix.ll
+
+; RUN: llvm-as < %s -o %t.bc
+; RUN: amd-llvm-spirv %t.bc --spirv-ext=+SPV_KHR_cooperative_matrix,+SPV_INTEL_joint_matrix -o %t.spv
+; RUN: amd-llvm-spirv %t.spv -to-text -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
+
+; CHECK-SPIRV-DAG: Capability CooperativeMatrixKHR
+; CHECK-SPIRV-DAG: Capability CooperativeMatrixOffsetInstructionsINTEL
+; CHECK-SPIRV-DAG: Extension "SPV_INTEL_joint_matrix"
+; CHECK-SPIRV-DAG: TypeInt [[#Int16Ty:]] 16 0
+; CHECK-SPIRV-DAG: TypeInt [[#Int32Ty:]] 32 0
+; CHECK-SPIRV-DAG: TypeInt [[#Int64Ty:]] 64 0
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const0:]] 0
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const1:]] 1
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const2:]] 2
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const3:]] 3
+; CHECK-SPIRV-DAG: Constant [[#Int32Ty]] [[#Const16:]] 16
+; CHECK-SPIRV-DAG: Constant [[#Int64Ty]] [[#Const128:]] 128 0
+; CHECK-SPIRV-DAG: Constant [[#Int64Ty:]] [[#Const256:]] 256 0
+; CHECK-SPIRV-DAG: TypeFloat [[#Float32Ty:]] 32
+; CHECK-SPIRV-DAG: TypeCooperativeMatrixKHR [[#MatTy1:]] [[#Float32Ty]] [[#Const3]] [[#Const1]] [[#Const16]] [[#Const2]]
+; CHECK-SPIRV-DAG: TypeCooperativeMatrixKHR [[#MatTy2:]] [[#Int16Ty]] [[#Const3]] [[#Const1]] [[#Const16]] [[#Const0:]]
+; CHECK-SPIRV-DAG: TypeCooperativeMatrixKHR [[#MatTy3:]] [[#Int16Ty]] [[#Const3]] [[#Const16]] [[#Const16]] [[#Const1]]
+; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy1]] [[#]] [[#Ptr1:]] [[#]] [[#Index1:]] [[#Const0]] [[#Const128]] 0
+; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy2]] [[#Load2:]] [[#]] [[#Index2:]] [[#]] [[#Const0]] [[#Const128]] 0
+; CHECK-SPIRV: CooperativeMatrixLoadOffsetINTEL [[#MatTy3]] [[#Load3:]] [[#]] [[#]] [[#]] [[#Const2:]] [[#Const256:]] 0
+; CHECK-SPIRV: CooperativeMatrixMulAddKHR [[#MatTy1]] [[#]] [[#Load2]] [[#Load3]] [[#Result:]] 64
+; CHECK-SPIRV: CooperativeMatrixStoreOffsetINTEL [[#Ptr1]] [[#Index2]] [[#Index1]] [[#Result]] [[#Const0]] [[#Const128]] 0
+
+; CHECK-LLVM: call spir_func target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) @_Z94__spirv_CooperativeMatrixLoadOffsetINTEL_RPU3AS144__spirv_CooperativeMatrixKHR__float_3_1_16_2PU3AS1fiiili(ptr addrspace(1) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0, i64 128, i32 0)
+; CHECK-LLVM: call spir_func target("spirv.CooperativeMatrixKHR", i16, 3, 1, 16, 0) @"_Z94__spirv_CooperativeMatrixLoadOffsetINTEL_RPU3AS144__spirv_CooperativeMatrixKHR__short_3_1_16_0PU3AS138class.sycl::_V1::ext::oneapi::bfloat16iiili"(ptr addrspace(1) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 0, i64 128, i32 0)
+; CHECK-LLVM: call spir_func target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) @"_Z95__spirv_CooperativeMatrixLoadOffsetINTEL_RPU3AS145__spirv_CooperativeMatrixKHR__short_3_16_16_1PU3AS138class.sycl::_V1::ext::oneapi::bfloat16iiili"(ptr addrspace(1) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, i32 2, i64 256, i32 0)
+; CHECK-LLVM: call spir_func target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) @_Z34__spirv_CooperativeMatrixMulAddKHRPU3AS144__spirv_CooperativeMatrixKHR__short_3_1_16_0PU3AS145__spirv_CooperativeMatrixKHR__short_3_16_16_1PU3AS144__spirv_CooperativeMatrixKHR__float_3_1_16_2i(target("spirv.CooperativeMatrixKHR", i16, 3, 1, 16, 0) %{{.*}}, target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) %{{.*}}, target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) %{{.*}}, i32 64)
+; CHECK-LLVM: call spir_func void @_Z41__spirv_CooperativeMatrixStoreOffsetINTELPU3AS1fiiPU3AS144__spirv_CooperativeMatrixKHR__float_3_1_16_2ili(ptr addrspace(1) %{{.*}}, i32 %{{.*}}, i32 %{{.*}}, target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) %{{.*}}, i32 0, i64 128, i32 0)
+
+; ModuleID = 'joint_matrix_all_sizes.cpp'
+source_filename = "joint_matrix_all_sizes.cpp"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
+%"class.sycl::_V1::detail::array" = type { [2 x i64] }
+%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
+%"class.sycl::_V1::ext::oneapi::bfloat16" = type { i16 }
+
+$_ZTSZZ15matrix_multiply = comdat any
+
+@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@__spirv_BuiltInLocalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+
+; Function Attrs: convergent mustprogress norecurse nounwind
+define weak_odr dso_local spir_kernel void @_ZTSZZ15matrix_multiply(ptr addrspace(1) noundef align 4 %_arg_accC, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_accC2, ptr noundef byval(%"class.sycl::_V1::id") align 8 %_arg_accC3, i64 noundef %_arg_sg_size, ptr addrspace(1) noundef readonly align 2 %_arg_accA, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_accA5, ptr noundef byval(%"class.sycl::_V1::id") align 8 %_arg_accA6, ptr addrspace(1) noundef readonly align 2 %_arg_accB, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_accB8, ptr noundef byval(%"class.sycl::_V1::id") align 8 %_arg_accB9) comdat {
+entry:
+  %agg.tmp11.sroa.0.sroa.2.0._arg_accC2.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accC2, i64 8
+  %agg.tmp11.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp11.sroa.0.sroa.2.0._arg_accC2.ascast.sroa_idx, align 8
+  %agg.tmp12.sroa.0.sroa.0.0.copyload = load i64, ptr %_arg_accC3, align 8
+  %agg.tmp12.sroa.0.sroa.2.0._arg_accC3.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accC3, i64 8
+  %agg.tmp12.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp12.sroa.0.sroa.2.0._arg_accC3.ascast.sroa_idx, align 8
+  %mul.i6.i.i.i.i = mul i64 %agg.tmp12.sroa.0.sroa.0.0.copyload, %agg.tmp11.sroa.0.sroa.2.0.copyload
+  %0 = getelementptr float, ptr addrspace(1) %_arg_accC, i64 %mul.i6.i.i.i.i
+  %add.ptr.i = getelementptr float, ptr addrspace(1) %0, i64 %agg.tmp12.sroa.0.sroa.2.0.copyload
+  %agg.tmp15.sroa.0.sroa.2.0._arg_accA5.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accA5, i64 8
+  %agg.tmp15.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp15.sroa.0.sroa.2.0._arg_accA5.ascast.sroa_idx, align 8
+  %agg.tmp16.sroa.0.sroa.0.0.copyload = load i64, ptr %_arg_accA6, align 8
+  %agg.tmp16.sroa.0.sroa.2.0._arg_accA6.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accA6, i64 8
+  %agg.tmp16.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp16.sroa.0.sroa.2.0._arg_accA6.ascast.sroa_idx, align 8
+  %mul.i6.i.i.i.i91 = mul i64 %agg.tmp16.sroa.0.sroa.0.0.copyload, %agg.tmp15.sroa.0.sroa.2.0.copyload
+  %1 = getelementptr %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %_arg_accA, i64 %mul.i6.i.i.i.i91
+  %add.ptr.i92 = getelementptr %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %1, i64 %agg.tmp16.sroa.0.sroa.2.0.copyload
+  %agg.tmp19.sroa.0.sroa.2.0._arg_accB8.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accB8, i64 8
+  %agg.tmp19.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp19.sroa.0.sroa.2.0._arg_accB8.ascast.sroa_idx, align 8
+  %agg.tmp20.sroa.0.sroa.0.0.copyload = load i64, ptr %_arg_accB9, align 8
+  %agg.tmp20.sroa.0.sroa.2.0._arg_accB9.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accB9, i64 8
+  %agg.tmp20.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp20.sroa.0.sroa.2.0._arg_accB9.ascast.sroa_idx, align 8
+  %mul.i6.i.i.i.i107 = mul i64 %agg.tmp20.sroa.0.sroa.0.0.copyload, %agg.tmp19.sroa.0.sroa.2.0.copyload
+  %2 = getelementptr %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %_arg_accB, i64 %mul.i6.i.i.i.i107
+  %add.ptr.i108 = getelementptr %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %2, i64 %agg.tmp20.sroa.0.sroa.2.0.copyload
+  %3 = load i64, ptr addrspace(1) getelementptr inbounds (i8, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, i64 8), align 8
+  %cmp.i28 = icmp ult i64 %3, 2147483648
+  tail call void @llvm.assume(i1 %cmp.i28)
+  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32
+  %cmp.i24 = icmp ult i64 %4, 2147483648
+  tail call void @llvm.assume(i1 %cmp.i24)
+  %5 = load i64, ptr addrspace(1) getelementptr inbounds (i8, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, i64 8), align 8
+  %cmp.i35 = icmp ult i64 %5, 2147483648
+  tail call void @llvm.assume(i1 %cmp.i35)
+  %sub.i = sub nsw i64 %3, %5
+  %6 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32
+  %cmp.i31 = icmp ult i64 %6, 2147483648
+  tail call void @llvm.assume(i1 %cmp.i31)
+  %sub5.i = sub nsw i64 %4, %6
+  %add.i7.i.i.i.i.i = add i64 %mul.i6.i.i.i.i, %agg.tmp12.sroa.0.sroa.2.0.copyload
+  %idx.neg.i.i = sub i64 0, %add.i7.i.i.i.i.i
+  %add.ptr.i.i = getelementptr inbounds float, ptr addrspace(1) %add.ptr.i, i64 %idx.neg.i.i
+  %div.i = udiv i64 %sub5.i, %_arg_sg_size
+  %conv.i = trunc i64 %sub.i to i32
+  %div.i.tr = trunc i64 %div.i to i32
+  %conv2.i = shl i32 %div.i.tr, 4
+  %call4.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) @_Z40__spirv_CooperativeMatrixLoadOffsetINTELIU3AS1ffLm1ELm16ELN5__spv9MatrixUseE2ELNS1_12MatrixLayoutE3ELNS1_5Scope4FlagE3EEPNS1_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEEPT_iiS3_mi(ptr addrspace(1) noundef %add.ptr.i.i, i32 noundef %conv.i, i32 noundef %conv2.i, i32 noundef 0, i64 noundef 128, i32 noundef 0)
+  %add.i7.i.i.i.i.i118 = add i64 %mul.i6.i.i.i.i91, %agg.tmp16.sroa.0.sroa.2.0.copyload
+  %idx.neg.i.i119 = sub i64 0, %add.i7.i.i.i.i.i118
+  %add.ptr.i.i120 = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %add.ptr.i92, i64 %idx.neg.i.i119
+  %add.i7.i.i.i.i.i126 = add i64 %mul.i6.i.i.i.i107, %agg.tmp20.sroa.0.sroa.2.0.copyload
+  %idx.neg.i.i127 = sub i64 0, %add.i7.i.i.i.i.i126
+  %add.ptr.i.i128 = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %add.ptr.i108, i64 %idx.neg.i.i127
+  %conv2.i60 = shl i32 %div.i.tr, 5
+  br label %for.cond.i
+
+for.cond.i:                                       ; preds = %for.body.i, %entry
+  %sub_c.i.sroa.0.0 = phi target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) [ %call4.i, %entry ], [ %call.i63, %for.body.i ]
+  %k.0.i = phi i32 [ 0, %entry ], [ %add.i, %for.body.i ]
+  %cmp.i = icmp samesign ult i32 %k.0.i, 8
+  br i1 %cmp.i, label %for.body.i, label %_ZZZ15matrix_multiplyIfN4sycl3_V13ext6oneapi8bfloat16ELm16ELm128ELm128ELi2ELm1ELm16ELm16E4multIS4_Lm1ELm16ELm16EEEvR10big_matrixIT_XT1_EXT2_EERS7_IT0_XT1_EXT3_EERS7_ISB_XdvT3_T4_EXmlT2_T4_EEENKUlRNS1_7handlerEE_clESH_ENKUlNS1_7nd_itemILi2EEEE_clESK_.exit
+
+for.body.i:                                       ; preds = %for.cond.i
+  %7 = shl nuw nsw i32 %k.0.i, 4
+  %call3.i50 = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 1, 16, 0) @_Z40__spirv_CooperativeMatrixLoadOffsetINTELIU3AS1KN4sycl3_V13ext6oneapi8bfloat16ES4_Lm1ELm16ELN5__spv9MatrixUseE0ELNS6_12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEEPT_iiS8_mi(ptr addrspace(1) noundef %add.ptr.i.i120, i32 noundef %conv.i, i32 noundef %7, i32 noundef 0, i64 noundef 128, i32 noundef 0)
+  %8 = shl nuw nsw i32 %k.0.i, 3
+  %call3.i61 = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) @_Z40__spirv_CooperativeMatrixLoadOffsetINTELIU3AS1KN4sycl3_V13ext6oneapi8bfloat16ES4_Lm16ELm16ELN5__spv9MatrixUseE1ELNS6_12MatrixLayoutE2ELNS6_5Scope4FlagE3EEPNS6_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEEPT_iiS8_mi(ptr addrspace(1) noundef %add.ptr.i.i128, i32 noundef %8, i32 noundef %conv2.i60, i32 noundef 2, i64 noundef 256, i32 noundef 0)
+  %call.i63 = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) @_Z34__spirv_CooperativeMatrixMulAddKHRIN4sycl3_V13ext6oneapi8bfloat16ES4_fLm1ELm16ELm16ELN5__spv9MatrixUseE0ELS6_1ELS6_2ELNS5_12MatrixLayoutE0ELS7_0ELS7_0ELNS5_5Scope4FlagE3EEPNS5_28__spirv_CooperativeMatrixKHRIT1_XT11_EXT2_EXT4_EXT7_EEEPNSA_IT_XT11_EXT2_EXT3_EXT5_EEEPNSA_IT0_XT11_EXT3_EXT4_EXT6_EEESD_m(target("spirv.CooperativeMatrixKHR", i16, 3, 1, 16, 0) noundef %call3.i50, target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) noundef %call3.i61, target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) noundef %sub_c.i.sroa.0.0, i64 noundef 64)
+  %add.i = add nuw nsw i32 %k.0.i, 1
+  br label %for.cond.i
+
+_ZZZ15matrix_multiplyIfN4sycl3_V13ext6oneapi8bfloat16ELm16ELm128ELm128ELi2ELm1ELm16ELm16E4multIS4_Lm1ELm16ELm16EEEvR10big_matrixIT_XT1_EXT2_EERS7_IT0_XT1_EXT3_EERS7_ISB_XdvT3_T4_EXmlT2_T4_EEENKUlRNS1_7handlerEE_clESH_ENKUlNS1_7nd_itemILi2EEEE_clESK_.exit: ; preds = %for.cond.i
+  tail call spir_func void @_Z41__spirv_CooperativeMatrixStoreOffsetINTELIU3AS1ffLm1ELm16ELN5__spv9MatrixUseE2ELNS1_12MatrixLayoutE3ELNS1_5Scope4FlagE3EEvPT_iiPNS1_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEES3_mi(ptr addrspace(1) noundef %add.ptr.i.i, i32 noundef %conv.i, i32 noundef %conv2.i, target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) noundef %sub_c.i.sroa.0.0, i32 noundef 0, i64 noundef 128, i32 noundef 0)
+  ret void
+}
+
+; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write)
+declare void @llvm.assume(i1 noundef)
+
+; Function Attrs: convergent nounwind
+declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) @_Z40__spirv_CooperativeMatrixLoadOffsetINTELIU3AS1ffLm1ELm16ELN5__spv9MatrixUseE2ELNS1_12MatrixLayoutE3ELNS1_5Scope4FlagE3EEPNS1_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEEPT_iiS3_mi(ptr addrspace(1) noundef, i32 noundef, i32 noundef, i32 noundef, i64 noundef, i32 noundef)
+
+; Function Attrs: convergent nounwind
+declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 1, 16, 0) @_Z40__spirv_CooperativeMatrixLoadOffsetINTELIU3AS1KN4sycl3_V13ext6oneapi8bfloat16ES4_Lm1ELm16ELN5__spv9MatrixUseE0ELNS6_12MatrixLayoutE0ELNS6_5Scope4FlagE3EEPNS6_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEEPT_iiS8_mi(ptr addrspace(1) noundef, i32 noundef, i32 noundef, i32 noundef, i64 noundef, i32 noundef)
+
+; Function Attrs: convergent nounwind
+declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) @_Z40__spirv_CooperativeMatrixLoadOffsetINTELIU3AS1KN4sycl3_V13ext6oneapi8bfloat16ES4_Lm16ELm16ELN5__spv9MatrixUseE1ELNS6_12MatrixLayoutE2ELNS6_5Scope4FlagE3EEPNS6_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEEPT_iiS8_mi(ptr addrspace(1) noundef, i32 noundef, i32 noundef, i32 noundef, i64 noundef, i32 noundef)
+
+; Function Attrs: convergent nounwind
+declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) @_Z34__spirv_CooperativeMatrixMulAddKHRIN4sycl3_V13ext6oneapi8bfloat16ES4_fLm1ELm16ELm16ELN5__spv9MatrixUseE0ELS6_1ELS6_2ELNS5_12MatrixLayoutE0ELS7_0ELS7_0ELNS5_5Scope4FlagE3EEPNS5_28__spirv_CooperativeMatrixKHRIT1_XT11_EXT2_EXT4_EXT7_EEEPNSA_IT_XT11_EXT2_EXT3_EXT5_EEEPNSA_IT0_XT11_EXT3_EXT4_EXT6_EEESD_m(target("spirv.CooperativeMatrixKHR", i16, 3, 1, 16, 0) noundef, target("spirv.CooperativeMatrixKHR", i16, 3, 16, 16, 1) noundef, target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) noundef, i64 noundef)
+
+; Function Attrs: convergent nounwind
+declare dso_local spir_func void @_Z41__spirv_CooperativeMatrixStoreOffsetINTELIU3AS1ffLm1ELm16ELN5__spv9MatrixUseE2ELNS1_12MatrixLayoutE3ELNS1_5Scope4FlagE3EEvPT_iiPNS1_28__spirv_CooperativeMatrixKHRIT0_XT5_EXT1_EXT2_EXT3_EEES3_mi(ptr addrspace(1) noundef, i32 noundef, i32 noundef, target("spirv.CooperativeMatrixKHR", float, 3, 1, 16, 2) noundef, i32 noundef, i64 noundef, i32 noundef)
diff --git a/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-gv-attributes.ll b/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-gv-attributes.ll
new file mode 100644
index 000000000..fb9a63e70
--- /dev/null
+++ b/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-gv-attributes.ll
@@ -0,0 +1,56 @@
+; RUN: llvm-as < %s -o %t.bc
+; RUN: not amd-llvm-spirv %t.bc -spirv-text --spirv-preserve-auxdata --spirv-max-version=1.5 --spirv-ext=-SPV_KHR_non_semantic_info,+SPV_INTEL_global_variable_decorations -o - 2>&1 | FileCheck %s --check-prefix=CHECK-SPIRV-EXT-DISABLED
+; RUN: amd-llvm-spirv %t.bc -o %t.spv --spirv-preserve-auxdata --spirv-max-version=1.5 --spirv-ext=+SPV_INTEL_global_variable_decorations
+; RUN: amd-llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-EXT
+; RUN: amd-llvm-spirv -r --spirv-preserve-auxdata %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.without.bc
+; RUN: llvm-dis %t.rev.without.bc -o - | FileCheck %s --implicit-check-not="{{foo|bar|baz}}"
+
+; RUN: amd-llvm-spirv %t.bc -spirv-text --spirv-preserve-auxdata --spirv-ext=+SPV_KHR_non_semantic_info,+SPV_INTEL_global_variable_decorations -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-NOEXT
+; RUN: amd-llvm-spirv %t.bc -o %t.spv --spirv-preserve-auxdata --spirv-ext=+SPV_INTEL_global_variable_decorations
+; RUN: amd-llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-NOEXT
+; RUN: amd-llvm-spirv -r --spirv-preserve-auxdata %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.without.bc
+; RUN: llvm-dis %t.rev.without.bc -o - | FileCheck %s --implicit-check-not="{{foo|bar|baz}}"
+
+; Check SPIR-V versions in a format magic number + version
+; CHECK-SPIRV-EXT: 119734787 65536
+; CHECK-SPIRV-EXT: Extension "SPV_KHR_non_semantic_info"
+; CHECK-SPIRV-NOEXT: 119734787 67072
+
+; CHECK-SPIRV: ExtInstImport [[#Import:]] "NonSemantic.AuxData"
+
+; CHECK-SPIRV: String [[#Attr0LHS:]] "sycl-device-global-size"
+; CHECK-SPIRV: String [[#Attr0RHS:]] "32"
+; CHECK-SPIRV: String [[#Attr1:]] "sycl-device-image-scope"
+; CHECK-SPIRV: String [[#Attr2LHS:]] "sycl-host-access"
+; CHECK-SPIRV: String [[#Attr2RHS:]] "0"
+; CHECK-SPIRV: String [[#Attr3LHS:]] "sycl-unique-id"
+; CHECK-SPIRV: String [[#Attr3RHS:]] "_Z20__AsanKernelMetadata"
+
+; CHECK-SPIRV: Name [[#GVName:]] "__AsanKernelMetadata"
+
+; CHECK-SPIRV: TypeVoid [[#VoidT:]]
+
+; CHECK-SPIRV: ExtInst [[#VoidT]] [[#Attr0Inst:]] [[#Import]] NonSemanticAuxDataGlobalVariableAttribute [[#GVName]] [[#Attr0LHS]] [[#Attr0RHS]] {{$}}
+; CHECK-SPIRV: ExtInst [[#VoidT]] [[#Attr1Inst:]] [[#Import]] NonSemanticAuxDataGlobalVariableAttribute [[#GVName]] [[#Attr1]] {{$}}
+; CHECK-SPIRV: ExtInst [[#VoidT]] [[#Attr1Inst:]] [[#Import]] NonSemanticAuxDataGlobalVariableAttribute [[#GVName]] [[#Attr2LHS]] [[#Attr2RHS]] {{$}}
+; CHECK-SPIRV: ExtInst [[#VoidT]] [[#Attr1Inst:]] [[#Import]] NonSemanticAuxDataGlobalVariableAttribute [[#GVName]] [[#Attr3LHS]] [[#Attr3RHS]] {{$}}
+
+target triple = "spir64-unknown-unknown"
+
+; CHECK-LLVM: @__AsanKernelMetadata = addrspace(1) global [1 x %structtype] [%structtype { i64 0, i64 92 }] #[[#GVIRAttr:]]
+%structtype = type { i64, i64 }
+
+@__AsanKernelMetadata = addrspace(1) global [1 x %structtype] [%structtype { i64 ptrtoint (ptr addrspace(2) null to i64), i64 92 }], !spirv.Decorations !0 #0
+
+; CHECK-LLVM: attributes #[[#GVIRAttr]] = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
+attributes #0 = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
+
+!0 = !{!1}
+!1 = !{i32 6147, i32 0, !"_Z20__AsanKernelMetadata"}
+
+; CHECK-SPIRV-EXT-DISABLED: RequiresExtension: Feature requires the following SPIR-V extension:
+; CHECK-SPIRV-EXT-DISABLED-NEXT: SPV_KHR_non_semantic_info
diff --git a/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-gv-metadata.ll b/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-gv-metadata.ll
new file mode 100644
index 000000000..cc63f04a3
--- /dev/null
+++ b/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-gv-metadata.ll
@@ -0,0 +1,41 @@
+; RUN: llvm-as < %s -o %t.bc
+; RUN: amd-llvm-spirv %t.bc -o %t.spv --spirv-preserve-auxdata --spirv-max-version=1.5
+; RUN: amd-llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-EXT
+; RUN: amd-llvm-spirv -r --spirv-preserve-auxdata %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.without.bc
+; RUN: llvm-dis %t.rev.without.bc -o - | FileCheck %s --implicit-check-not="{{foo|bar|baz}}"
+
+; RUN: amd-llvm-spirv %t.bc -o %t.spv --spirv-preserve-auxdata
+; RUN: amd-llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-NOEXT
+; RUN: amd-llvm-spirv -r --spirv-preserve-auxdata %t.spv -o %t.rev.bc
+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
+; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.without.bc
+; RUN: llvm-dis %t.rev.without.bc -o - | FileCheck %s --implicit-check-not="{{foo|bar|baz}}"
+
+; Check SPIR-V versions in a format magic number + version
+; CHECK-SPIRV-EXT: 119734787 65536
+; CHECK-SPIRV-EXT: Extension "SPV_KHR_non_semantic_info"
+; CHECK-SPIRV-NOEXT: 119734787 67072
+
+; CHECK-SPIRV: ExtInstImport [[#Import:]] "NonSemantic.AuxData"
+
+; CHECK-SPIRV: String [[#MDName:]] "absolute_symbol"
+
+; CHECK-SPIRV: Name [[#GVName:]] "a"
+
+; CHECK-SPIRV: TypeInt [[#Int32T:]] 64 0
+; CHECK-SPIRV: Constant [[#Int32T]] [[#MDValue0:]] 0
+; CHECK-SPIRV: Constant [[#Int32T]] [[#MDValue1:]] 16
+
+; CHECK-SPIRV: TypeVoid [[#VoidT:]]
+
+; CHECK-SPIRV: ExtInst [[#VoidT]] [[#ValInst:]] [[#Import]] NonSemanticAuxDataGlobalVariableMetadata [[#GVName]] [[#MDName]] [[#MDValue0]] [[#MDValue1]] {{$}}
+
+target triple = "spir64-unknown-unknown"
+
+; CHECK-LLVM: @a = external addrspace(1) global i8, !absolute_symbol ![[#LLVMVal:]]
+@a = external addrspace(1) global i8, !absolute_symbol !0
+
+; CHECK-LLVM: ![[#LLVMVal]] = !{i64 0, i64 16}
+!0 = !{i64 0, i64 16}
diff --git a/test/lit.cfg.py b/test/lit.cfg.py
index 7d161425d..486dc4983 100644
--- a/test/lit.cfg.py
+++ b/test/lit.cfg.py
@@ -21,7 +21,7 @@
 # excludes: A list of directories  and fles to exclude from the testsuite.
 config.excludes = ['CMakeLists.txt']
 
-if config.spirv_tools_found:
+if config.libspirv_dis:
     config.available_features.add('libspirv_dis')
 
 if not config.spirv_skip_debug_info_tests:
diff --git a/test/lit.site.cfg.py.in b/test/lit.site.cfg.py.in
index a47fb197c..f7a23ce58 100644
--- a/test/lit.site.cfg.py.in
+++ b/test/lit.site.cfg.py.in
@@ -19,6 +19,7 @@ config.host_arch = "@HOST_ARCH@"
 config.python_executable = "@PYTHON_EXECUTABLE@"
 config.test_run_dir = "@CMAKE_CURRENT_BINARY_DIR@"
 config.spirv_tools_found = "@SPIRV_TOOLS_FOUND@"
+config.libspirv_dis = "@SPIRV_ENABLE_LIBSPIRV_DIS@"
 config.spirv_tools_have_spirv_as = @SPIRV_TOOLS_SPIRV_AS_FOUND@
 config.spirv_tools_have_spirv_dis = @SPIRV_TOOLS_SPIRV_DIS_FOUND@
 config.spirv_tools_have_spirv_link = @SPIRV_TOOLS_SPIRV_LINK_FOUND@
diff --git a/test/optnone.ll b/test/optnone.ll
index bb90775d9..abcaa62d2 100644
--- a/test/optnone.ll
+++ b/test/optnone.ll
@@ -1,32 +1,44 @@
 ; RUN: llvm-as %s -o %t.bc
-; RUN: amd-llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_optnone -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: amd-llvm-spirv %t.bc --spirv-ext=+SPV_EXT_optnone -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-EXT
+; RUN: amd-llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_optnone -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-INTEL
+; RUN: amd-llvm-spirv %t.bc --spirv-ext=+SPV_EXT_optnone,+SPV_INTEL_optnone -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-BOTH
 
-; RUN: amd-llvm-spirv --spirv-ext=+SPV_INTEL_optnone %t.bc -o %t.spv
+
+; RUN: amd-llvm-spirv --spirv-ext=+SPV_EXT_optnone %t.bc -o %t.spv
 ; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.bc
 ; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
 
 ; Check that optnone is correctly ignored when extension is not enabled
-; RUN: amd-llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-NO-EXT
+; RUN: amd-llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-NONE
 
 ; RUN: amd-llvm-spirv %t.bc -o %t.spv
 ; RUN: amd-llvm-spirv -r %t.spv -o %t.rev.bc
-; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-NO-EXT
+; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-NONE
+
+; Note: The capability is unconditionally printed with the EXT suffix.
+; CHECK-SPIRV-EXT: Capability OptNoneEXT
+; CHECK-SPIRV-INTEL: Capability OptNoneEXT
+; CHECK-SPIRV-BOTH: Capability OptNoneEXT
 
-; CHECK-SPIRV: Capability OptNoneINTEL
-; CHECK-SPIRV: Extension "SPV_INTEL_optnone"
+; CHECK-SPIRV-EXT: Extension "SPV_EXT_optnone"
+; CHECK-SPIRV-INTEL: Extension "SPV_INTEL_optnone"
+; Note: When both extensions are enabled, prefer the EXT extension.
+; CHECK-SPIRV-BOTH: Extension "SPV_EXT_optnone"
 
 ; Per SPIR-V spec:
 ; FunctionControlDontInlineMask = 0x2 (2)
-; Per SPIR-V spec extension SPV_INTEL_optnone:
-; FunctionControlOptNoneINTELMask = 0x10000 (65536)
-; CHECK-SPIRV: Function {{[0-9]+}} {{[0-9]+}} 65538
-; CHECK-SPIRV-NO-EXT: Function {{[0-9]+}} {{[0-9]+}} 2
+; Per SPIR-V spec extension spec:
+; FunctionControlOptNoneMask = 0x10000 (65536)
+; CHECK-SPIRV-EXT: Function {{[0-9]+}} {{[0-9]+}} 65538
+; CHECK-SPIRV-INTEL: Function {{[0-9]+}} {{[0-9]+}} 65538
+; CHECK-SPIRV-BOTH: Function {{[0-9]+}} {{[0-9]+}} 65538
+; CHECK-SPIRV-NONE: Function {{[0-9]+}} {{[0-9]+}} 2
 
 ; CHECK-LLVM: define spir_func void @_Z3foov() #[[ATTR:[0-9]+]]
 ; CHECK-LLVM: #[[ATTR]] = { {{.*}}noinline{{.*}}optnone{{.*}} }
 
-; CHECK-LLVM-NO-EXT: define spir_func void @_Z3foov() #[[ATTR:[0-9]+]]
-; CHECK-LLVM-NO-EXT-NOT: #[[ATTR]] = { {{.*}}noinline{{.*}}optnone{{.*}} }
+; CHECK-LLVM-NONE: define spir_func void @_Z3foov() #[[ATTR:[0-9]+]]
+; CHECK-LLVM-NONE-NOT: #[[ATTR]] = { {{.*}}noinline{{.*}}optnone{{.*}} }
 
 target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir-unknown-unknown"
diff --git a/test/transcoding/OpImageSampleExplicitLod_arg.cl b/test/transcoding/OpImageSampleExplicitLod_arg.cl
index ddc3b2f9a..c65861721 100644
--- a/test/transcoding/OpImageSampleExplicitLod_arg.cl
+++ b/test/transcoding/OpImageSampleExplicitLod_arg.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -O1 -triple spir-unknown-unknown -cl-std=CL2.0 %s -finclude-default-header -emit-llvm-bc -o %t.bc
+// RUN: %clang_cc1 -O1 -triple spir-unknown-unknown -cl-std=CL2.0 %s -fdeclare-opencl-builtins -finclude-default-header -emit-llvm-bc -o %t.bc
 // RUN: amd-llvm-spirv %t.bc -spirv-text -o %t.txt
 // RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV
 // RUN: amd-llvm-spirv %t.bc -o %t.spv
diff --git a/test/transcoding/OpImageWrite.cl b/test/transcoding/OpImageWrite.cl
index 6f7e545b5..17abcc948 100644
--- a/test/transcoding/OpImageWrite.cl
+++ b/test/transcoding/OpImageWrite.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -O1 -triple spir-unknown-unknown -cl-std=CL2.0 %s -finclude-default-header -emit-llvm-bc -o %t.bc
+// RUN: %clang_cc1 -O1 -triple spir-unknown-unknown -cl-std=CL2.0 %s -fdeclare-opencl-builtins -finclude-default-header -emit-llvm-bc -o %t.bc
 // RUN: amd-llvm-spirv %t.bc -spirv-text -o %t.txt
 // RUN: FileCheck < %t.txt %s --check-prefix=CHECK-SPIRV
 // RUN: amd-llvm-spirv %t.bc -o %t.spv
diff --git a/tools/llvm-spirv/CMakeLists.txt b/tools/llvm-spirv/CMakeLists.txt
index b42df6d20..f3e6596af 100644
--- a/tools/llvm-spirv/CMakeLists.txt
+++ b/tools/llvm-spirv/CMakeLists.txt
@@ -32,8 +32,8 @@ target_include_directories(amd-llvm-spirv
     ${LLVM_SPIRV_INCLUDE_DIRS}
 )
 
-if(SPIRV_TOOLS_FOUND)
-  target_compile_definitions(amd-llvm-spirv PRIVATE LLVM_SPIRV_HAVE_SPIRV_TOOLS=1)
-  target_include_directories(amd-llvm-spirv PRIVATE ${SPIRV_TOOLS_INCLUDE_DIRS})
-  target_link_libraries(amd-llvm-spirv PRIVATE ${SPIRV_TOOLS_LDFLAGS})
-endif(SPIRV_TOOLS_FOUND)
+if(SPIRV_TOOLS_FOUND AND LLVM_SPIRV_ENABLE_LIBSPIRV_DIS)
+  target_compile_definitions(llvm-spirv PRIVATE LLVM_SPIRV_HAVE_SPIRV_TOOLS=1)
+  target_include_directories(llvm-spirv PRIVATE ${SPIRV_TOOLS_INCLUDE_DIRS})
+  target_link_libraries(llvm-spirv PRIVATE ${SPIRV_TOOLS_LDFLAGS})
+endif()