diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 47022b3f89a8b..36facbd05aa67 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -21,6 +21,7 @@
 #include "SPIRVUtils.h"
 #include "llvm/ADT/APInt.h"
 #include "llvm/IR/Constants.h"
+#include "llvm/IR/Function.h"
 #include "llvm/IR/IntrinsicInst.h"
 #include "llvm/IR/Intrinsics.h"
 #include "llvm/IR/IntrinsicsSPIRV.h"
@@ -223,14 +224,37 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeVoid(MachineIRBuilder &MIRBuilder) {
 }
 
 void SPIRVGlobalRegistry::invalidateMachineInstr(MachineInstr *MI) {
-  // TODO:
-  // - review other data structure wrt. possible issues related to removal
-  //   of a machine instruction during instruction selection.
+  // Other maps that may hold MachineInstr*:
+  // - VRegToTypeMap: Clearing would require a linear search. If we are deleting
+  // type, then no registers remaining in the code should have this type. Should
+  // be safe to leave as is.
+  // - FunctionToInstr & FunctionToInstrRev: At this point, we should not be
+  // deleting functions. No need to update.
+  // - AliasInstMDMap: Would require a linear search, and the Intel Alias
+  // instruction are not instructions instruction selection will be able to
+  // remove.
+
+  const SPIRVSubtarget &ST = MI->getMF()->getSubtarget<SPIRVSubtarget>();
+  const SPIRVInstrInfo *TII = ST.getInstrInfo();
+  assert(!TII->isAliasingInstr(*MI) &&
+         "Cannot invalidate aliasing instructions.");
+  assert(MI->getOpcode() != SPIRV::OpFunction &&
+         "Cannot invalidate OpFunction.");
+
+  if (MI->getOpcode() == SPIRV::OpFunctionCall) {
+    if (const auto *F = dyn_cast<Function>(MI->getOperand(2).getGlobal())) {
+      auto It = ForwardCalls.find(F);
+      if (It != ForwardCalls.end()) {
+        It->second.erase(MI);
+        if (It->second.empty())
+          ForwardCalls.erase(It);
+      }
+    }
+  }
+
   const MachineFunction *MF = MI->getMF();
   auto It = LastInsertedTypeMap.find(MF);
-  if (It == LastInsertedTypeMap.end())
-    return;
-  if (It->second == MI)
+  if (It != LastInsertedTypeMap.end() && It->second == MI)
     LastInsertedTypeMap.erase(MF);
   // remove from the duplicate tracker to avoid incorrect reuse
   erase(MI);
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index fc87288a4a212..d6e371c178392 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -94,6 +94,9 @@ class SPIRVInstructionSelector : public InstructionSelector {
 
 private:
   void resetVRegsType(MachineFunction &MF);
+  // New helper function for dead instruction removal
+  void removeDeadInstruction(MachineInstr &MI) const;
+  void removeOpNamesForDeadMI(MachineInstr &MI) const;
 
   // tblgen-erated 'select' implementation, used as the initial selector for
   // the patterns that don't require complex C++.
@@ -506,22 +509,195 @@ static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
   return false;
 }
 
+// TODO(168736): We should make this either a flag in tabelgen
+// or reduce our dependence on the global registry, so we can remove this
+// function. It can easily be missed when new intrinsics are added.
+static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
+  switch (ID) {
+  // Intrinsics that do not have side effects.
+  // This is not an exhaustive list and may need to be updated.
+  case Intrinsic::spv_all:
+  case Intrinsic::spv_alloca:
+  case Intrinsic::spv_any:
+  case Intrinsic::spv_bitcast:
+  case Intrinsic::spv_const_composite:
+  case Intrinsic::spv_cross:
+  case Intrinsic::spv_degrees:
+  case Intrinsic::spv_distance:
+  case Intrinsic::spv_extractelt:
+  case Intrinsic::spv_extractv:
+  case Intrinsic::spv_faceforward:
+  case Intrinsic::spv_fdot:
+  case Intrinsic::spv_firstbitlow:
+  case Intrinsic::spv_firstbitshigh:
+  case Intrinsic::spv_firstbituhigh:
+  case Intrinsic::spv_frac:
+  case Intrinsic::spv_gep:
+  case Intrinsic::spv_global_offset:
+  case Intrinsic::spv_global_size:
+  case Intrinsic::spv_group_id:
+  case Intrinsic::spv_insertelt:
+  case Intrinsic::spv_insertv:
+  case Intrinsic::spv_isinf:
+  case Intrinsic::spv_isnan:
+  case Intrinsic::spv_lerp:
+  case Intrinsic::spv_length:
+  case Intrinsic::spv_normalize:
+  case Intrinsic::spv_num_subgroups:
+  case Intrinsic::spv_num_workgroups:
+  case Intrinsic::spv_ptrcast:
+  case Intrinsic::spv_radians:
+  case Intrinsic::spv_reflect:
+  case Intrinsic::spv_refract:
+  case Intrinsic::spv_resource_getpointer:
+  case Intrinsic::spv_resource_handlefrombinding:
+  case Intrinsic::spv_resource_handlefromimplicitbinding:
+  case Intrinsic::spv_resource_nonuniformindex:
+  case Intrinsic::spv_rsqrt:
+  case Intrinsic::spv_saturate:
+  case Intrinsic::spv_sdot:
+  case Intrinsic::spv_sign:
+  case Intrinsic::spv_smoothstep:
+  case Intrinsic::spv_step:
+  case Intrinsic::spv_subgroup_id:
+  case Intrinsic::spv_subgroup_local_invocation_id:
+  case Intrinsic::spv_subgroup_max_size:
+  case Intrinsic::spv_subgroup_size:
+  case Intrinsic::spv_thread_id:
+  case Intrinsic::spv_thread_id_in_group:
+  case Intrinsic::spv_udot:
+  case Intrinsic::spv_undef:
+  case Intrinsic::spv_value_md:
+  case Intrinsic::spv_workgroup_size:
+    return false;
+  default:
+    return true;
+  }
+}
+
+// TODO(168736): We should make this either a flag in tabelgen
+// or reduce our dependence on the global registry, so we can remove this
+// function. It can easily be missed when new intrinsics are added.
+static bool isOpcodeWithNoSideEffects(unsigned Opcode) {
+  switch (Opcode) {
+  case SPIRV::OpTypeVoid:
+  case SPIRV::OpTypeBool:
+  case SPIRV::OpTypeInt:
+  case SPIRV::OpTypeFloat:
+  case SPIRV::OpTypeVector:
+  case SPIRV::OpTypeMatrix:
+  case SPIRV::OpTypeImage:
+  case SPIRV::OpTypeSampler:
+  case SPIRV::OpTypeSampledImage:
+  case SPIRV::OpTypeArray:
+  case SPIRV::OpTypeRuntimeArray:
+  case SPIRV::OpTypeStruct:
+  case SPIRV::OpTypeOpaque:
+  case SPIRV::OpTypePointer:
+  case SPIRV::OpTypeFunction:
+  case SPIRV::OpTypeEvent:
+  case SPIRV::OpTypeDeviceEvent:
+  case SPIRV::OpTypeReserveId:
+  case SPIRV::OpTypeQueue:
+  case SPIRV::OpTypePipe:
+  case SPIRV::OpTypeForwardPointer:
+  case SPIRV::OpTypePipeStorage:
+  case SPIRV::OpTypeNamedBarrier:
+  case SPIRV::OpTypeAccelerationStructureNV:
+  case SPIRV::OpTypeCooperativeMatrixNV:
+  case SPIRV::OpTypeCooperativeMatrixKHR:
+    return true;
+  default:
+    return false;
+  }
+}
+
 bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
+  // If there are no definitions, then assume there is some other
+  // side-effect that makes this instruction live.
+  if (MI.getNumDefs() == 0)
+    return false;
+
   for (const auto &MO : MI.all_defs()) {
     Register Reg = MO.getReg();
-    if (Reg.isPhysical() || !MRI.use_nodbg_empty(Reg))
+    if (Reg.isPhysical()) {
+      LLVM_DEBUG(dbgs() << "Not dead: def of physical register " << Reg);
       return false;
+    }
+    for (const auto &UseMI : MRI.use_nodbg_instructions(Reg)) {
+      if (UseMI.getOpcode() != SPIRV::OpName) {
+        LLVM_DEBUG(dbgs() << "Not dead: def " << MO << " has use in " << UseMI);
+        return false;
+      }
+    }
   }
+
   if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
-      MI.isLifetimeMarker())
+      MI.isLifetimeMarker()) {
+    LLVM_DEBUG(
+        dbgs()
+        << "Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
     return false;
-  if (MI.isPHI())
+  }
+  if (MI.isPHI()) {
+    LLVM_DEBUG(dbgs() << "Dead: Phi instruction with no uses.\n");
     return true;
+  }
+
+  if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
+      MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
+    const auto &Intr = cast<GIntrinsic>(MI);
+    if (!intrinsicHasSideEffects(Intr.getIntrinsicID())) {
+      LLVM_DEBUG(dbgs() << "Dead: Intrinsic with no real side effects.\n");
+      return true;
+    }
+  }
+
   if (MI.mayStore() || MI.isCall() ||
       (MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
-      MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo())
+      MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo()) {
+    LLVM_DEBUG(dbgs() << "Not dead: instruction has side effects.\n");
     return false;
-  return true;
+  }
+
+  if (isPreISelGenericOpcode(MI.getOpcode())) {
+    // TODO: Is there a generic way to check if the opcode has side effects?
+    LLVM_DEBUG(dbgs() << "Dead: Generic opcode with no uses.\n");
+    return true;
+  }
+
+  if (isOpcodeWithNoSideEffects(MI.getOpcode())) {
+    LLVM_DEBUG(dbgs() << "Dead: known opcode with no side effects\n");
+    return true;
+  }
+
+  return false;
+}
+
+void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const {
+  // Delete the OpName that uses the result if there is one.
+  for (const auto &MO : MI.all_defs()) {
+    Register Reg = MO.getReg();
+    if (Reg.isPhysical())
+      continue;
+    SmallVector<MachineInstr *, 4> UselessOpNames;
+    for (MachineInstr &UseMI : MRI->use_nodbg_instructions(Reg)) {
+      assert(UseMI.getOpcode() == SPIRV::OpName &&
+             "There is still a use of the dead function.");
+      UselessOpNames.push_back(&UseMI);
+    }
+    for (MachineInstr *OpNameMI : UselessOpNames) {
+      GR.invalidateMachineInstr(OpNameMI);
+      OpNameMI->eraseFromParent();
+    }
+  }
+}
+
+void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &MI) const {
+  salvageDebugInfo(*MRI, MI);
+  GR.invalidateMachineInstr(&MI);
+  removeOpNamesForDeadMI(MI);
+  MI.eraseFromParent();
 }
 
 bool SPIRVInstructionSelector::select(MachineInstr &I) {
@@ -530,6 +706,13 @@ bool SPIRVInstructionSelector::select(MachineInstr &I) {
   assert(I.getParent() && "Instruction should be in a basic block!");
   assert(I.getParent()->getParent() && "Instruction should be in a function!");
 
+  LLVM_DEBUG(dbgs() << "Checking if instruction is dead: " << I;);
+  if (isDead(I, *MRI)) {
+    LLVM_DEBUG(dbgs() << "Instruction is dead.\n");
+    removeDeadInstruction(I);
+    return true;
+  }
+
   Register Opcode = I.getOpcode();
   // If it's not a GMIR instruction, we've selected it already.
   if (!isPreISelGenericOpcode(Opcode)) {
@@ -581,9 +764,7 @@ bool SPIRVInstructionSelector::select(MachineInstr &I) {
     // if the instruction has been already made dead by folding it away
     // erase it
     LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");
-    salvageDebugInfo(*MRI, I);
-    GR.invalidateMachineInstr(&I);
-    I.eraseFromParent();
+    removeDeadInstruction(I);
     return true;
   }
 
diff --git a/llvm/test/CodeGen/SPIRV/OpVariable_order.ll b/llvm/test/CodeGen/SPIRV/OpVariable_order.ll
index 1e94be0886307..a43a4d66d04bb 100644
--- a/llvm/test/CodeGen/SPIRV/OpVariable_order.ll
+++ b/llvm/test/CodeGen/SPIRV/OpVariable_order.ll
@@ -13,7 +13,9 @@
 define void @main() {
 entry:
   %0 = alloca <2 x i32>, align 4
+  store <2 x i32> zeroinitializer, ptr %0, align 4
   %1 = getelementptr <2 x i32>, ptr %0, i32 0, i32 0
   %2 = alloca float, align 4
+  store float 0.0, ptr %2, align 4
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll b/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll
index 9e91854de1172..b0bad1819a25d 100644
--- a/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll
+++ b/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll
@@ -29,9 +29,12 @@
 %Struct7 = type [2 x %Struct]
 %Nested = type { %Struct7 }
 
+@G = global %Struct zeroinitializer
+
 define spir_kernel void @foo(ptr addrspace(4) %arg1, ptr addrspace(4) %arg2) {
 entry:
   %var = alloca %Struct
+  store %Struct zeroinitializer, ptr %var
   %r1 = call %Struct @_Z29__spirv_SpecConstantComposite_1(float 1.0)
   store %Struct %r1, ptr addrspace(4) %arg1
   %r2 = call %Struct7 @_Z29__spirv_SpecConstantComposite_2(%Struct %r1, %Struct %r1)
diff --git a/llvm/test/CodeGen/SPIRV/basic_float_types.ll b/llvm/test/CodeGen/SPIRV/basic_float_types.ll
index a0ba97e1d1f14..6cdc67bbf24ee 100644
--- a/llvm/test/CodeGen/SPIRV/basic_float_types.ll
+++ b/llvm/test/CodeGen/SPIRV/basic_float_types.ll
@@ -2,6 +2,9 @@
 ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_bfloat16 %s -o - | FileCheck %s
 ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-unknown --spirv-ext=+SPV_KHR_bfloat16 %s -o - -filetype=obj | spirv-val %}
 
+// TODO: Open bug bfloat16 cannot be stored to.
+XFAIL: *
+
 define void @main() {
 entry:
 
@@ -49,50 +52,66 @@ entry:
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_half]] Function
   %half_Val = alloca half, align 2
+  store half 0.0, ptr %half_Val, align 2
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_bfloat]] Function
   %bfloat_Val = alloca bfloat, align 2
+  store bfloat 0.0, ptr %bfloat_Val, align 2
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_float]] Function
   %float_Val = alloca float, align 4
+  store float 0.0, ptr %float_Val, align 4
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_double]] Function
   %double_Val = alloca double, align 8
+  store double 0.0, ptr %double_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2half]] Function
   %half2_Val = alloca <2 x half>, align 4
+  store <2 x half> zeroinitializer, ptr %half2_Val, align 4
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3half]] Function
   %half3_Val = alloca <3 x half>, align 8
+  store <3 x half> zeroinitializer, ptr %half3_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4half]] Function
   %half4_Val = alloca <4 x half>, align 8
+  store <4 x half> zeroinitializer, ptr %half4_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2bfloat]] Function
   %bfloat2_Val = alloca <2 x bfloat>, align 4
+  store <2 x bfloat> zeroinitializer, ptr %bfloat2_Val, align 4
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3bfloat]] Function
   %bfloat3_Val = alloca <3 x bfloat>, align 8
+  store <3 x bfloat> zeroinitializer, ptr %bfloat3_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4bfloat]] Function
   %bfloat4_Val = alloca <4 x bfloat>, align 8
+  store <4 x bfloat> zeroinitializer, ptr %bfloat4_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2float]] Function
   %float2_Val = alloca <2 x float>, align 8
+  store <2 x float> zeroinitializer, ptr %float2_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3float]] Function
   %float3_Val = alloca <3 x float>, align 16
+  store <3 x float> zeroinitializer, ptr %float3_Val, align 16
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4float]] Function
   %float4_Val = alloca <4 x float>, align 16
+  store <4 x float> zeroinitializer, ptr %float4_Val, align 16
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2double]] Function
   %double2_Val = alloca <2 x double>, align 16
+  store <2 x double> zeroinitializer, ptr %double2_Val, align 16
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3double]] Function
   %double3_Val = alloca <3 x double>, align 32
+  store <3 x double> zeroinitializer, ptr %double3_Val, align 32
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4double]] Function
   %double4_Val = alloca <4 x double>, align 32
+  store <4 x double> zeroinitializer, ptr %double4_Val, align 32
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/basic_int_types.ll b/llvm/test/CodeGen/SPIRV/basic_int_types.ll
index 5aa7aaf6fbd01..1ed241eed4019 100644
--- a/llvm/test/CodeGen/SPIRV/basic_int_types.ll
+++ b/llvm/test/CodeGen/SPIRV/basic_int_types.ll
@@ -37,39 +37,51 @@ entry:
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_short]] Function
   %int16_t_Val = alloca i16, align 2
+  store i16 0, ptr %int16_t_Val, align 2
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_int]] Function
   %int_Val = alloca i32, align 4
+  store i32 0, ptr %int_Val, align 4
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_long]] Function
   %int64_t_Val = alloca i64, align 8
+  store i64 0, ptr %int64_t_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2short]] Function
   %int16_t2_Val = alloca <2 x i16>, align 4
+  store <2 x i16> zeroinitializer, ptr %int16_t2_Val, align 4
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3short]] Function
   %int16_t3_Val = alloca <3 x i16>, align 8
+  store <3 x i16> zeroinitializer, ptr %int16_t3_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4short]] Function
   %int16_t4_Val = alloca <4 x i16>, align 8
+  store <4 x i16> zeroinitializer, ptr %int16_t4_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2int]] Function
   %int2_Val = alloca <2 x i32>, align 8
+  store <2 x i32> zeroinitializer, ptr %int2_Val, align 8
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3int]] Function
   %int3_Val = alloca <3 x i32>, align 16
+  store <3 x i32> zeroinitializer, ptr %int3_Val, align 16
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4int]] Function
   %int4_Val = alloca <4 x i32>, align 16
+  store <4 x i32> zeroinitializer, ptr %int4_Val, align 16
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2long]] Function
   %int64_t2_Val = alloca <2 x i64>, align 16
+  store <2 x i64> zeroinitializer, ptr %int64_t2_Val, align 16
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3long]] Function
   %int64_t3_Val = alloca <3 x i64>, align 32
+  store <3 x i64> zeroinitializer, ptr %int64_t3_Val, align 32
 
 ; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4long]] Function
   %int64_t4_Val = alloca <4 x i64>, align 32
+  store <4 x i64> zeroinitializer, ptr %int64_t4_Val, align 32
 
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll b/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll
index 56b5f48715533..f3c8f9967211a 100644
--- a/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll
+++ b/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll
@@ -6,39 +6,51 @@ define void @main() {
 entry:
 ; CHECK: %int16_t_Val = OpVariable %_ptr_Function_ushort Function
   %int16_t_Val = alloca i16, align 2
+  store i16 0, i16* %int16_t_Val, align 2
 
 ; CHECK: %int_Val = OpVariable %_ptr_Function_uint Function
   %int_Val = alloca i32, align 4
+  store i32 0, i32* %int_Val, align 4
 
 ; CHECK: %int64_t_Val = OpVariable %_ptr_Function_ulong Function
   %int64_t_Val = alloca i64, align 8
+  store i64 0, i64* %int64_t_Val, align 8
 
 ; CHECK: %int16_t2_Val = OpVariable %_ptr_Function_v2ushort Function
   %int16_t2_Val = alloca <2 x i16>, align 4
+  store <2 x i16> zeroinitializer, <2 x i16>* %int16_t2_Val, align 4
 
 ; CHECK: %int16_t3_Val = OpVariable %_ptr_Function_v3ushort Function
   %int16_t3_Val = alloca <3 x i16>, align 8
+  store <3 x i16> zeroinitializer, <3 x i16>* %int16_t3_Val, align 8
 
 ; CHECK: %int16_t4_Val = OpVariable %_ptr_Function_v4ushort Function
   %int16_t4_Val = alloca <4 x i16>, align 8
+  store <4 x i16> zeroinitializer, <4 x i16>* %int16_t4_Val, align 8
 
 ; CHECK: %int2_Val = OpVariable %_ptr_Function_v2uint Function
   %int2_Val = alloca <2 x i32>, align 8
+  store <2 x i32> zeroinitializer, <2 x i32>* %int2_Val, align 8
 
 ; CHECK: %int3_Val = OpVariable %_ptr_Function_v3uint Function
   %int3_Val = alloca <3 x i32>, align 16
+  store <3 x i32> zeroinitializer, <3 x i32>* %int3_Val, align 16
 
 ; CHECK: %int4_Val = OpVariable %_ptr_Function_v4uint Function
   %int4_Val = alloca <4 x i32>, align 16
+  store <4 x i32> zeroinitializer, <4 x i32>* %int4_Val, align 16
 
 ; CHECK: %int64_t2_Val = OpVariable %_ptr_Function_v2ulong Function
   %int64_t2_Val = alloca <2 x i64>, align 16
+  store <2 x i64> zeroinitializer, <2 x i64>* %int64_t2_Val, align 16
 
 ; CHECK: %int64_t3_Val = OpVariable %_ptr_Function_v3ulong Function
   %int64_t3_Val = alloca <3 x i64>, align 32
+  store <3 x i64> zeroinitializer, <3 x i64>* %int64_t3_Val, align 32
 
 ; CHECK: %int64_t4_Val = OpVariable %_ptr_Function_v4ulong Function
   %int64_t4_Val = alloca <4 x i64>, align 32
+  store <4 x i64> zeroinitializer, <4 x i64>* %int64_t4_Val, align 32
 
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
index 39a755e736081..bca90f4ebd151 100644
--- a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
+++ b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
@@ -33,6 +33,28 @@ target triple = "spirv32-unknown-unknown"
 ; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input
 ; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input
 
+@G_spv_num_workgroups_0 = global i32 0
+@G_spv_num_workgroups_1 = global i32 0
+@G_spv_num_workgroups_2 = global i32 0
+@G_spv_workgroup_size_0 = global i32 0
+@G_spv_workgroup_size_1 = global i32 0
+@G_spv_workgroup_size_2 = global i32 0
+@G_spv_group_id_0 = global i32 0
+@G_spv_group_id_1 = global i32 0
+@G_spv_group_id_2 = global i32 0
+@G_spv_thread_id_in_group_0 = global i32 0
+@G_spv_thread_id_in_group_1 = global i32 0
+@G_spv_thread_id_in_group_2 = global i32 0
+@G_spv_thread_id_0 = global i32 0
+@G_spv_thread_id_1 = global i32 0
+@G_spv_thread_id_2 = global i32 0
+@G_spv_global_size_0 = global i32 0
+@G_spv_global_size_1 = global i32 0
+@G_spv_global_size_2 = global i32 0
+@G_spv_global_offset_0 = global i32 0
+@G_spv_global_offset_1 = global i32 0
+@G_spv_global_offset_2 = global i32 0
+
 ; Function Attrs: convergent noinline norecurse nounwind optnone
 define spir_func void @test_id_and_range() {
 entry:
@@ -44,66 +66,87 @@ entry:
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.num.workgroups = call i32 @llvm.spv.num.workgroups.i32(i32 0)
+  store i32 %spv.num.workgroups, i32* @G_spv_num_workgroups_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.num.workgroups1 = call i32 @llvm.spv.num.workgroups.i32(i32 1)
+  store i32 %spv.num.workgroups1, i32* @G_spv_num_workgroups_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.num.workgroups2 = call i32 @llvm.spv.num.workgroups.i32(i32 2)
+  store i32 %spv.num.workgroups2, i32* @G_spv_num_workgroups_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.workgroup.size = call i32 @llvm.spv.workgroup.size.i32(i32 0)
+  store i32 %spv.workgroup.size, i32* @G_spv_workgroup_size_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.workgroup.size3 = call i32 @llvm.spv.workgroup.size.i32(i32 1)
+  store i32 %spv.workgroup.size3, i32* @G_spv_workgroup_size_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.workgroup.size4 = call i32 @llvm.spv.workgroup.size.i32(i32 2)
+  store i32 %spv.workgroup.size4, i32* @G_spv_workgroup_size_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.group.id = call i32 @llvm.spv.group.id.i32(i32 0)
+  store i32 %spv.group.id, i32* @G_spv_group_id_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.group.id5 = call i32 @llvm.spv.group.id.i32(i32 1)
+  store i32 %spv.group.id5, i32* @G_spv_group_id_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.group.id6 = call i32 @llvm.spv.group.id.i32(i32 2)
+  store i32 %spv.group.id6, i32* @G_spv_group_id_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.thread.id.in.group = call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
+  store i32 %spv.thread.id.in.group, i32* @G_spv_thread_id_in_group_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.thread.id.in.group7 = call i32 @llvm.spv.thread.id.in.group.i32(i32 1)
+  store i32 %spv.thread.id.in.group7, i32* @G_spv_thread_id_in_group_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.thread.id.in.group8 = call i32 @llvm.spv.thread.id.in.group.i32(i32 2)
+  store i32 %spv.thread.id.in.group8, i32* @G_spv_thread_id_in_group_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.thread.id = call i32 @llvm.spv.thread.id.i32(i32 0)
+  store i32 %spv.thread.id, i32* @G_spv_thread_id_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.thread.id9 = call i32 @llvm.spv.thread.id.i32(i32 1)
+  store i32 %spv.thread.id9, i32* @G_spv_thread_id_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.thread.id10 = call i32 @llvm.spv.thread.id.i32(i32 2)
+  store i32 %spv.thread.id10, i32* @G_spv_thread_id_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.num.workgroups11 = call i32 @llvm.spv.global.size.i32(i32 0)
+  store i32 %spv.num.workgroups11, i32* @G_spv_global_size_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.num.workgroups12 = call i32 @llvm.spv.global.size.i32(i32 1)
+  store i32 %spv.num.workgroups12, i32* @G_spv_global_size_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.num.workgroups13 = call i32 @llvm.spv.global.size.i32(i32 2)
+  store i32 %spv.num.workgroups13, i32* @G_spv_global_size_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
   %spv.global.offset = call i32 @llvm.spv.global.offset.i32(i32 0)
+  store i32 %spv.global.offset, i32* @G_spv_global_offset_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
   %spv.global.offset14 = call i32 @llvm.spv.global.offset.i32(i32 1)
+  store i32 %spv.global.offset14, i32* @G_spv_global_offset_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
 ; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
   %spv.global.offset15 = call i32 @llvm.spv.global.offset.i32(i32 2)
+  store i32 %spv.global.offset15, i32* @G_spv_global_offset_2
 ; CHECK: OpLoad %5 [[SubgroupSize]]
   %0 = call i32 @llvm.spv.subgroup.size()
   store i32 %0, ptr %ssize, align 4
diff --git a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
index dcdf8992ce1c4..26c2d866d14c7 100644
--- a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
+++ b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
@@ -34,6 +34,28 @@ target triple = "spirv64-unknown-unknown"
 ; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input
 ; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input
 
+@G_spv_num_workgroups_0 = global i64 0
+@G_spv_num_workgroups_1 = global i64 0
+@G_spv_num_workgroups_2 = global i64 0
+@G_spv_workgroup_size_0 = global i64 0
+@G_spv_workgroup_size_1 = global i64 0
+@G_spv_workgroup_size_2 = global i64 0
+@G_spv_group_id_0 = global i64 0
+@G_spv_group_id_1 = global i64 0
+@G_spv_group_id_2 = global i64 0
+@G_spv_thread_id_in_group_0 = global i64 0
+@G_spv_thread_id_in_group_1 = global i64 0
+@G_spv_thread_id_in_group_2 = global i64 0
+@G_spv_thread_id_0 = global i64 0
+@G_spv_thread_id_1 = global i64 0
+@G_spv_thread_id_2 = global i64 0
+@G_spv_global_size_0 = global i64 0
+@G_spv_global_size_1 = global i64 0
+@G_spv_global_size_2 = global i64 0
+@G_spv_global_offset_0 = global i64 0
+@G_spv_global_offset_1 = global i64 0
+@G_spv_global_offset_2 = global i64 0
+
 ; Function Attrs: convergent noinline norecurse nounwind optnone
 define spir_func void @test_id_and_range() {
 entry:
@@ -45,66 +67,87 @@ entry:
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.num.workgroups = call i64 @llvm.spv.num.workgroups.i64(i32 0)
+  store i64 %spv.num.workgroups, i64* @G_spv_num_workgroups_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.num.workgroups1 = call i64 @llvm.spv.num.workgroups.i64(i32 1)
+  store i64 %spv.num.workgroups1, i64* @G_spv_num_workgroups_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.num.workgroups2 = call i64 @llvm.spv.num.workgroups.i64(i32 2)
+  store i64 %spv.num.workgroups2, i64* @G_spv_num_workgroups_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.workgroup.size = call i64 @llvm.spv.workgroup.size.i64(i32 0)
+  store i64 %spv.workgroup.size, i64* @G_spv_workgroup_size_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.workgroup.size3 = call i64 @llvm.spv.workgroup.size.i64(i32 1)
+  store i64 %spv.workgroup.size3, i64* @G_spv_workgroup_size_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.workgroup.size4 = call i64 @llvm.spv.workgroup.size.i64(i32 2)
+  store i64 %spv.workgroup.size4, i64* @G_spv_workgroup_size_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.group.id = call i64 @llvm.spv.group.id.i64(i32 0)
+  store i64 %spv.group.id, i64* @G_spv_group_id_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.group.id5 = call i64 @llvm.spv.group.id.i64(i32 1)
+  store i64 %spv.group.id5, i64* @G_spv_group_id_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.group.id6 = call i64 @llvm.spv.group.id.i64(i32 2)
+  store i64 %spv.group.id6, i64* @G_spv_group_id_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.thread.id.in.group = call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
+  store i64 %spv.thread.id.in.group, i64* @G_spv_thread_id_in_group_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.thread.id.in.group7 = call i64 @llvm.spv.thread.id.in.group.i64(i32 1)
+  store i64 %spv.thread.id.in.group7, i64* @G_spv_thread_id_in_group_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.thread.id.in.group8 = call i64 @llvm.spv.thread.id.in.group.i64(i32 2)
+  store i64 %spv.thread.id.in.group8, i64* @G_spv_thread_id_in_group_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.thread.id = call i64 @llvm.spv.thread.id.i64(i32 0)
+  store i64 %spv.thread.id, i64* @G_spv_thread_id_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.thread.id9 = call i64 @llvm.spv.thread.id.i64(i32 1)
+  store i64 %spv.thread.id9, i64* @G_spv_thread_id_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.thread.id10 = call i64 @llvm.spv.thread.id.i64(i32 2)
+  store i64 %spv.thread.id10, i64* @G_spv_thread_id_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.num.workgroups11 = call i64 @llvm.spv.global.size.i64(i32 0)
+  store i64 %spv.num.workgroups11, i64* @G_spv_global_size_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.num.workgroups12 = call i64 @llvm.spv.global.size.i64(i32 1)
+  store i64 %spv.num.workgroups12, i64* @G_spv_global_size_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.num.workgroups13 = call i64 @llvm.spv.global.size.i64(i32 2)
+  store i64 %spv.num.workgroups13, i64* @G_spv_global_size_2
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
   %spv.global.offset = call i64 @llvm.spv.global.offset.i64(i32 0)
+  store i64 %spv.global.offset, i64* @G_spv_global_offset_0
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
   %spv.global.offset14 = call i64 @llvm.spv.global.offset.i64(i32 1)
+  store i64 %spv.global.offset14, i64* @G_spv_global_offset_1
 ; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
 ; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
   %spv.global.offset15 = call i64 @llvm.spv.global.offset.i64(i32 2)
+  store i64 %spv.global.offset15, i64* @G_spv_global_offset_2
 ; CHECK: OpLoad %5 [[SubgroupSize]]
   %0 = call i32 @llvm.spv.subgroup.size()
   store i32 %0, ptr %ssize, align 4
diff --git a/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll b/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
index 0c9b29de890d4..8dd9b387a6d84 100644
--- a/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
+++ b/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
@@ -81,17 +81,36 @@
 @__spirv_BuiltInSubgroupId = external addrspace(1) global i32
 @__spirv_BuiltInSubgroupLocalInvocationId = external addrspace(1) global i32
 
+@G_r1 = global i64 0
+@G_r2 = global i64 0
+@G_r3 = global i32 0
+@G_r4 = global i32 0
+@G_r5 = global i32 0
+@G_r6 = global i32 0
+@G_r7 = global i32 0
+@G_r8 = global i32 0
+@G_r9 = global i32 0
+
 define spir_kernel void @_Z1wv() {
 entry:
   %r1 = tail call spir_func i64 @get_global_linear_id()
+  store i64 %r1, i64* @G_r1
   %r2 = tail call spir_func i64 @get_local_linear_id()
+  store i64 %r2, i64* @G_r2
   %r3 = tail call spir_func i32 @get_work_dim()
+  store i32 %r3, i32* @G_r3
   %r4 = tail call spir_func i32 @get_sub_group_size()
+  store i32 %r4, i32* @G_r4
   %r5 = tail call spir_func i32 @get_max_sub_group_size()
+  store i32 %r5, i32* @G_r5
   %r6 = tail call spir_func i32 @get_num_sub_groups()
+  store i32 %r6, i32* @G_r6
   %r7 = tail call spir_func i32 @get_enqueued_num_sub_groups()
+  store i32 %r7, i32* @G_r7
   %r8 = tail call spir_func i32 @get_sub_group_id()
+  store i32 %r8, i32* @G_r8
   %r9 = tail call spir_func i32 @get_sub_group_local_id()
+  store i32 %r9, i32* @G_r9
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
index 3e0d0cc4cd8e2..d260c9f94d4ad 100644
--- a/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
@@ -126,6 +126,7 @@ define spir_func i32 @test0() !dbg !17 {
   %14 = load ptr addrspace(4), ptr %11, align 4, !dbg !65
   store ptr addrspace(4) %14, ptr %12, align 4, !dbg !64
     #dbg_declare(ptr %13, !66, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !70)
+  store [8 x i32] zeroinitializer, ptr %13, align 4
   ret i32 0, !dbg !71
 }
 
@@ -169,6 +170,7 @@ define spir_func i32 @test1() !dbg !72 {
   %14 = load ptr addrspace(4), ptr %11, align 4, !dbg !97
   store ptr addrspace(4) %14, ptr %12, align 4, !dbg !96
     #dbg_declare(ptr %13, !98, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !99)
+  store [8 x i32] zeroinitializer, ptr %13, align 4
   ret i32 0, !dbg !100
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/event-zero-const.ll b/llvm/test/CodeGen/SPIRV/event-zero-const.ll
index 523d2ad9825f3..2bf8259e78785 100644
--- a/llvm/test/CodeGen/SPIRV/event-zero-const.ll
+++ b/llvm/test/CodeGen/SPIRV/event-zero-const.ll
@@ -12,11 +12,15 @@
 ; CHECK: OpINotEqual %[[#]] %[[#]] %[[#LongNull]]
 ; CHECK: OpGroupAsyncCopy %[[#EventTy]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#EventNull]]
 
+@G_r1 = global i1 0
+@G_e1 = global target("spirv.Event") poison
 
 define weak_odr dso_local spir_kernel void @foo(i64 %_arg_i, ptr addrspace(1) %_arg_ptr, ptr addrspace(3) %_arg_local) {
 entry:
   %r1 = icmp ne i64 %_arg_i, 0
+  store i1 %r1, ptr @G_r1
   %e1 = tail call spir_func target("spirv.Event") @__spirv_GroupAsyncCopy(i32 2, ptr addrspace(3) %_arg_local, ptr addrspace(1) %_arg_ptr, i64 1, i64 1, target("spirv.Event") zeroinitializer)
+  store target("spirv.Event") %e1, ptr @G_e1
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll
index e5736b88b63a3..a9a0d3358f8cc 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll
@@ -11,15 +11,22 @@
 @G1 = addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr @bar to ptr addrspace(4))] }
 @G2 = addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr null to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr @bar to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4))] }
 
+@G_r1_foo = global ptr addrspace(4) null
+@G_r2_foo = global ptr addrspace(4) null
+@G_r1_bar = global ptr addrspace(4) null
+
 define void @foo(ptr addrspace(4) %p) {
 entry:
   %r1 = addrspacecast ptr @foo to ptr addrspace(4)
+  store ptr addrspace(4) %r1, ptr @G_r1_foo
   %r2 = addrspacecast ptr null to ptr addrspace(4)
+  store ptr addrspace(4) %r2, ptr @G_r2_foo
   ret void
 }
 
 define void @bar(ptr addrspace(4) %p) {
 entry:
   %r1 = addrspacecast ptr @bar to ptr addrspace(4)
+  store ptr addrspace(4) %r1, ptr @G_r1_bar
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll
index 22668e71fb257..92652f1faefc0 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll
@@ -12,11 +12,16 @@
 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"
 target triple = "spir64-unknown-unknown"
 
+@G1 = global bfloat 0.0
+@G2 = global <2 x bfloat> zeroinitializer
+
 define spir_kernel void @test() {
 entry:
   %addr1 = alloca bfloat
   %addr2 = alloca <2 x bfloat>
   %data1 = load bfloat, ptr %addr1
   %data2 = load <2 x bfloat>, ptr %addr2
+  store bfloat %data1, ptr @G1
+  store <2 x bfloat> %data2, ptr @G2
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll
index d3fe9e43450cd..81497f26f1aef 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll
@@ -79,6 +79,54 @@
 ; CHECK: OpDecorate %[[#maxResV]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|AllowContract|AllowReassoc|AllowTransform
 ; CHECK: OpDecorate %[[#maxCommonResV]] FPFastMathMode NotNaN|NotInf
 
+@G_addRes = global float 0.0
+@G_subRes = global float 0.0
+@G_mulRes = global float 0.0
+@G_divRes = global float 0.0
+@G_remRes = global float 0.0
+@G_negRes = global float 0.0
+@G_oeqRes = global i1 0
+@G_oneRes = global i1 0
+@G_oltRes = global i1 0
+@G_ogtRes = global i1 0
+@G_oleRes = global i1 0
+@G_ogeRes = global i1 0
+@G_ordRes = global i1 0
+@G_ueqRes = global i1 0
+@G_uneRes = global i1 0
+@G_ultRes = global i1 0
+@G_ugtRes = global i1 0
+@G_uleRes = global i1 0
+@G_ugeRes = global i1 0
+@G_unoRes = global i1 0
+@G_modRes = global float 0.0
+@G_maxRes = global float 0.0
+@G_maxCommonRes = global float 0.0
+
+@G_addResV = global <2 x float> zeroinitializer
+@G_subResV = global <2 x float> zeroinitializer
+@G_mulResV = global <2 x float> zeroinitializer
+@G_divResV = global <2 x float> zeroinitializer
+@G_remResV = global <2 x float> zeroinitializer
+@G_negResV = global <2 x float> zeroinitializer
+@G_oeqResV = global <2 x i1> zeroinitializer
+@G_oneResV = global <2 x i1> zeroinitializer
+@G_oltResV = global <2 x i1> zeroinitializer
+@G_ogtResV = global <2 x i1> zeroinitializer
+@G_oleResV = global <2 x i1> zeroinitializer
+@G_ogeResV = global <2 x i1> zeroinitializer
+@G_ordResV = global <2 x i1> zeroinitializer
+@G_ueqResV = global <2 x i1> zeroinitializer
+@G_uneResV = global <2 x i1> zeroinitializer
+@G_ultResV = global <2 x i1> zeroinitializer
+@G_ugtResV = global <2 x i1> zeroinitializer
+@G_uleResV = global <2 x i1> zeroinitializer
+@G_ugeResV = global <2 x i1> zeroinitializer
+@G_unoResV = global <2 x i1> zeroinitializer
+@G_modResV = global <2 x float> zeroinitializer
+@G_maxResV = global <2 x float> zeroinitializer
+@G_maxCommonResV = global <2 x float> zeroinitializer
+
 ; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
 declare spir_func float @_Z4fmodff(float, float)
 declare dso_local spir_func noundef nofpclass(nan inf) float @_Z16__spirv_ocl_fmaxff(float noundef nofpclass(nan inf), float noundef nofpclass(nan inf)) local_unnamed_addr #1
@@ -91,55 +139,101 @@ declare dso_local spir_func noundef nofpclass(nan inf) <2 x float> @_Z23__spirv_
 define weak_odr dso_local spir_kernel void @foo(float %1, float %2) {
 entry:
   %addRes = fadd float %1,  %2
+  store float %addRes, float* @G_addRes
   %subRes = fsub nnan float %1,  %2
+  store float %subRes, float* @G_subRes
   %mulRes = fmul ninf float %1,  %2
+  store float %mulRes, float* @G_mulRes
   %divRes = fdiv nsz float %1,  %2
+  store float %divRes, float* @G_divRes
   %remRes = frem arcp float %1,  %2
+  store float %remRes, float* @G_remRes
   %negRes = fneg fast float %1
+  store float %negRes, float* @G_negRes
   %oeqRes = fcmp nnan ninf oeq float %1,  %2
+  store i1 %oeqRes, i1* @G_oeqRes
   %oneRes = fcmp one float %1,  %2, !spirv.Decorations !3
+  store i1 %oneRes, i1* @G_oneRes
   %oltRes = fcmp nnan olt float %1,  %2, !spirv.Decorations !3
+  store i1 %oltRes, i1* @G_oltRes
   %ogtRes = fcmp ninf ogt float %1,  %2, !spirv.Decorations !3
+  store i1 %ogtRes, i1* @G_ogtRes
   %oleRes = fcmp nsz ole float %1,  %2, !spirv.Decorations !3
+  store i1 %oleRes, i1* @G_oleRes
   %ogeRes = fcmp arcp oge float %1,  %2, !spirv.Decorations !3
+  store i1 %ogeRes, i1* @G_ogeRes
   %ordRes = fcmp fast ord float %1,  %2, !spirv.Decorations !3
+  store i1 %ordRes, i1* @G_ordRes
   %ueqRes = fcmp nnan ninf ueq float %1,  %2, !spirv.Decorations !3
+  store i1 %ueqRes, i1* @G_ueqRes
   %uneRes = fcmp une float %1,  %2, !spirv.Decorations !3
+  store i1 %uneRes, i1* @G_uneRes
   %ultRes = fcmp ult float %1,  %2, !spirv.Decorations !3
+  store i1 %ultRes, i1* @G_ultRes
   %ugtRes = fcmp ugt float %1,  %2, !spirv.Decorations !3
+  store i1 %ugtRes, i1* @G_ugtRes
   %uleRes = fcmp ule float %1,  %2, !spirv.Decorations !3
+  store i1 %uleRes, i1* @G_uleRes
   %ugeRes = fcmp uge float %1,  %2, !spirv.Decorations !3
+  store i1 %ugeRes, i1* @G_ugeRes
   %unoRes = fcmp uno float %1,  %2, !spirv.Decorations !3
+  store i1 %unoRes, i1* @G_unoRes
   %modRes = call spir_func float @_Z4fmodff(float %1, float %2)
+  store float %modRes, float* @G_modRes
   %maxRes = tail call fast spir_func noundef nofpclass(nan inf) float @_Z16__spirv_ocl_fmaxff(float noundef nofpclass(nan inf) %1, float noundef nofpclass(nan inf) %2)
+  store float %maxRes, float* @G_maxRes
    %maxCommonRes = tail call spir_func noundef float @_Z23__spirv_ocl_fmax_commonff(float noundef nofpclass(nan inf) %1, float noundef nofpclass(nan inf) %2)
+  store float %maxCommonRes, float* @G_maxCommonRes
   ret void
 }
 
 define weak_odr dso_local spir_kernel void @fooV(<2 x float> %v1, <2 x float> %v2) {
   %addResV = fadd <2 x float> %v1,  %v2
+  store <2 x float> %addResV, <2 x float>* @G_addResV
   %subResV = fsub nnan <2 x float> %v1,  %v2
+  store <2 x float> %subResV, <2 x float>* @G_subResV
   %mulResV = fmul ninf <2 x float> %v1,  %v2
+  store <2 x float> %mulResV, <2 x float>* @G_mulResV
   %divResV = fdiv nsz <2 x float> %v1,  %v2
+  store <2 x float> %divResV, <2 x float>* @G_divResV
   %remResV = frem arcp <2 x float> %v1,  %v2
+  store <2 x float> %remResV, <2 x float>* @G_remResV
   %negResV = fneg fast <2 x float> %v1
+  store <2 x float> %negResV, <2 x float>* @G_negResV
   %oeqResV = fcmp nnan ninf oeq <2 x float> %v1,  %v2
+  store <2 x i1> %oeqResV, <2 x i1>* @G_oeqResV
   %oneResV = fcmp one <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %oneResV, <2 x i1>* @G_oneResV
   %oltResV = fcmp nnan olt <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %oltResV, <2 x i1>* @G_oltResV
   %ogtResV = fcmp ninf ogt <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ogtResV, <2 x i1>* @G_ogtResV
   %oleResV = fcmp nsz ole <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %oleResV, <2 x i1>* @G_oleResV
   %ogeResV = fcmp arcp oge <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ogeResV, <2 x i1>* @G_ogeResV
   %ordResV = fcmp fast ord <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ordResV, <2 x i1>* @G_ordResV
   %ueqResV = fcmp nnan ninf ueq <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ueqResV, <2 x i1>* @G_ueqResV
   %uneResV = fcmp une <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %uneResV, <2 x i1>* @G_uneResV
   %ultResV = fcmp ult <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ultResV, <2 x i1>* @G_ultResV
   %ugtResV = fcmp ugt <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ugtResV, <2 x i1>* @G_ugtResV
   %uleResV = fcmp ule <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %uleResV, <2 x i1>* @G_uleResV
   %ugeResV = fcmp uge <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %ugeResV, <2 x i1>* @G_ugeResV
   %unoResV = fcmp uno <2 x float> %v1,  %v2, !spirv.Decorations !3
+  store <2 x i1> %unoResV, <2 x i1>* @G_unoResV
   %modResV = call spir_func <2 x float> @_Z4fmodDv2_fDv2_f(<2 x float> %v1, <2 x float> %v2)
+  store <2 x float> %modResV, <2 x float>* @G_modResV
   %maxResV = tail call fast spir_func noundef nofpclass(nan inf) <2 x float> @_Z16__spirv_ocl_fmaxDv2_fDv2_f(<2 x float> noundef nofpclass(nan inf) %v1, <2 x float> noundef nofpclass(nan inf) %v2)
+  store <2 x float> %maxResV, <2 x float>* @G_maxResV
    %maxCommonResV = tail call spir_func noundef <2 x float> @_Z23__spirv_ocl_fmax_commonDv2_fDv2_f(<2 x float> noundef nofpclass(nan inf) %v1, <2 x float> noundef nofpclass(nan inf) %v2)
+  store <2 x float> %maxCommonResV, <2 x float>* @G_maxCommonResV
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll b/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll
index 4db0ba33d52c9..face4a9f5e615 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll
@@ -2,10 +2,15 @@
 ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=KHR %s -o - | FileCheck %s
 ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=khr %s -o - | FileCheck %s
 
+@G = global i32 0
+
 define i6 @foo() {
   %call = tail call i32 @llvm.bitreverse.i32(i32 42)
+  store i32 %call, ptr @G
   ret i6 2
 }
 
 ; CHECK-NOT: OpExtension "SPV_INTEL_arbitrary_precision_integers"
 ; CHECK-DAG: OpExtension "SPV_KHR_bit_instructions"
+
+declare i32 @llvm.bitreverse.i32(i32)
diff --git a/llvm/test/CodeGen/SPIRV/freeze.ll b/llvm/test/CodeGen/SPIRV/freeze.ll
index 9077d2ede72a9..4f7e7794ed03b 100644
--- a/llvm/test/CodeGen/SPIRV/freeze.ll
+++ b/llvm/test/CodeGen/SPIRV/freeze.ll
@@ -1,15 +1,15 @@
 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
 ; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
-; CHECK: OpName %[[Arg1:.*]] "arg1"
-; CHECK: OpName %[[Arg2:.*]] "arg2"
-; CHECK: OpName %[[NotAStaticPoison:.*]] "poison1"
-; CHECK: OpName %[[NotAStaticPoison]] "nil0"
-; CHECK: OpName %[[StaticPoisonIntFreeze:.*]] "nil1"
-; CHECK: OpName %[[StaticPoisonFloatFreeze:.*]] "nil2"
-; CHECK: OpName %[[Arg1]] "val1"
-; CHECK: OpName %[[Const100:.*]] "val2"
-; CHECK: OpName %[[Const100]] "val3"
+; CHECK-DAG: OpName %[[Arg1:.*]] "arg1"
+; CHECK-DAG: OpName %[[Arg2:.*]] "arg2"
+; CHECK-DAG: OpName %[[NotAStaticPoison:.*]] "poison1"
+; CHECK-DAG: OpName %[[NotAStaticPoison]] "nil0"
+; CHECK-DAG: OpName %[[StaticPoisonIntFreeze:.*]] "nil1"
+; CHECK-DAG: OpName %[[StaticPoisonFloatFreeze:.*]] "nil2"
+; CHECK-DAG: OpName %[[Arg1]] "val1"
+; CHECK-DAG: OpName %[[Const100:.*]] "val2"
+; CHECK-DAG: OpName %[[Const100]] "val3"
 ; CHECK: OpDecorate
 ; CHECK-DAG: %[[FloatTy:.*]] = OpTypeFloat 32
 ; CHECK-DAG: %[[ShortTy:.*]] = OpTypeInt 16 0
@@ -18,17 +18,37 @@
 ; CHECK-DAG: %[[Undef32:.*]] = OpUndef %[[IntTy]]
 ; CHECK-DAG: %[[UndefFloat:.*]] = OpUndef %[[FloatTy]]
 ; CHECK-DAG: %[[Const100]] = OpConstant %[[IntTy]] 100
-; CHECK: %[[Arg1]] = OpFunctionParameter %[[FloatTy]]
-; CHECK: %[[NotAStaticPoison]] = OpIAdd %[[ShortTy]] %[[Arg2]] %[[Undef16]]
 
-define spir_func void @foo(float %arg1, i16 %arg2) {
+define spir_func i16 @test_nil0(i16 %arg2) {
 entry:
+; CHECK: %[[NotAStaticPoison]] = OpIAdd %[[ShortTy]] %[[Arg2]] %[[Undef16]]
   %poison1 = add i16 %arg2, undef
   %nil0 = freeze i16 %poison1
+  ret i16 %nil0
+}
+
+define spir_func i32 @test_nil1() {
+entry:
   %nil1 = freeze i32 undef
+  ret i32 %nil1
+}
+
+define spir_func float @test_nil2() {
+entry:
   %nil2 = freeze float poison
+  ret float %nil2
+}
+
+define spir_func float @freeze_float(float %arg1) {
+entry:
+; CHECK: %[[Arg1]] = OpFunctionParameter %[[FloatTy]]
   %val1 = freeze float %arg1
+  ret float %val1
+}
+
+define spir_func i32 @foo() {
+entry:
   %val2 = freeze i32 100
   %val3 = freeze i32 %val2
-  ret void
+  ret i32 %val3
 }
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll
index a97492b8453ea..a15d628cc3614 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll
@@ -63,7 +63,7 @@ entry:
 ; CHECK: %[[#a_high:]] = OpVectorShuffle %[[#vec2_int_32]] %[[#a]] %[[#undef_v4i32]] 1 3
 ; CHECK: %[[#b_low:]] = OpVectorShuffle %[[#vec2_int_32]] %[[#b]] %[[#undef_v4i32]] 0 2
 ; CHECK: %[[#b_high:]] = OpVectorShuffle %[[#vec2_int_32]] %[[#b]] %[[#undef_v4i32]] 1 3
-; CHECK: %[[#iaddcarry:]] = OpIAddCarry %[[#struct_v2i32_v2i32]] %[[#a_low]] %[[#vec2_int_32]]
+; CHECK: %[[#iaddcarry:]] = OpIAddCarry %[[#struct_v2i32_v2i32]] %[[#a_low]] %[[#b_low]]
 ; CHECK: %[[#lowsum:]] = OpCompositeExtract %[[#vec2_int_32]] %[[#iaddcarry]] 0
 ; CHECK: %[[#carry:]] = OpCompositeExtract %[[#vec2_int_32]] %[[#iaddcarry]] 1
 ; CHECK: %[[#carry_ne0:]] = OpINotEqual %[[#vec2_bool]] %[[#carry]] %[[#const_v2i32_0_0]]
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll
index 4a15fa8b14537..75fac211f1108 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll
@@ -3,24 +3,25 @@
 
 ; CHECK: OpExtInstImport "GLSL.std.450"
 
+@i = global i32 0, align 4
+@absi = global i32 0, align 4
+@f = global float 0.0, align 4
+@absf = global float 0.0, align 4
+
 define void @main() #1 {
 entry:
-  %i = alloca i32, align 4
-  %absi = alloca i32, align 4
-  %f = alloca float, align 4
-  %absf = alloca float, align 4
-  %0 = load i32, ptr %i, align 4
+  %0 = load i32, ptr @i, align 4
 
 ; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] SAbs %[[#]]
   %elt.abs = call i32 @llvm.abs.i32(i32 %0, i1 false)
 
-  store i32 %elt.abs, ptr %absi, align 4
-  %1 = load float, ptr %f, align 4
+  store i32 %elt.abs, ptr @absi, align 4
+  %1 = load float, ptr @f, align 4
 
 ; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] FAbs %[[#]]
   %elt.abs1 = call float @llvm.fabs.f32(float %1)
 
-  store float %elt.abs1, ptr %absf, align 4
+  store float %elt.abs1, ptr @absf, align 4
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll
index 7583066c01cf8..dceaa8c209957 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll
@@ -7,21 +7,23 @@
 ; CHECK: %[[#v4float:]] = OpTypeVector %[[#float]] 4
 ; CHECK: %[[#float_0_30103001:]] = OpConstant %[[#float]] 0.30103000998497009
 
+@logf = global float 0.0, align 4
+@logf4 = global <4 x float> zeroinitializer, align 16
+
 define void @main(float %f, <4 x float> %f4) {
 entry:
 ; CHECK-DAG: %[[#f:]] = OpFunctionParameter %[[#float]]
 ; CHECK-DAG: %[[#f4:]] = OpFunctionParameter %[[#v4float]]
-  %logf = alloca float, align 4
-  %logf4 = alloca <4 x float>, align 16
-
 
 ; CHECK: %[[#log2:]] = OpExtInst %[[#float]] %[[#extinst]] Log2 %[[#f]]
 ; CHECK: %[[#res:]] = OpFMul %[[#float]] %[[#log2]] %[[#float_0_30103001]]
   %elt.log10 = call float @llvm.log10.f32(float %f)
+  store float %elt.log10, ptr @logf, align 4
 
 ; CHECK: %[[#log2:]] = OpExtInst %[[#v4float]] %[[#extinst]] Log2 %[[#f4]]
 ; CHECK: %[[#res:]] = OpVectorTimesScalar %[[#v4float]] %[[#log2]] %[[#float_0_30103001]]
   %elt.log101 = call <4 x float> @llvm.log10.v4f32(<4 x float> %f4)
+  store <4 x float> %elt.log101, ptr @logf4, align 16
 
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll b/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll
index b788f34bf7238..02825e3cbb599 100644
--- a/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll
+++ b/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll
@@ -4,25 +4,40 @@
 ; CHECK-LABEL: Begin function original_testcase
 define fastcc void @original_testcase() {
 top:
+  %0 = alloca [1 x ptr], align 4
   ; CHECK: OpCompositeInsert
-  %0 = insertvalue [1 x ptr] zeroinitializer, ptr poison, 0
+  %1 = insertvalue [1 x ptr] zeroinitializer, ptr poison, 0
+  store [1 x ptr] %1, ptr %0
   ret void
 }
 
 ; CHECK-LABEL: Begin function additional_testcases
 define fastcc void @additional_testcases() {
 top:
+  %0 = alloca [2 x ptr], align 4
+
+
   ; Test with different pointer types
   ; CHECK: OpCompositeInsert
   %1 = insertvalue [1 x ptr] zeroinitializer, ptr undef, 0
+  ; CHECK: OpStore
+  store [1 x ptr] %1, ptr %0
+
   ; CHECK-NEXT: OpCompositeInsert
   %2 = insertvalue {ptr, i32} zeroinitializer, ptr poison, 0
+  ; CHECK: OpStore
+  store {ptr, i32} %2, ptr %0
+
   ; CHECK-NEXT: OpCompositeInsert
   %3 = insertvalue {ptr, ptr} undef, ptr null, 0
+  ; CHECK: OpStore
+  store {ptr, ptr} %3, ptr %0
 
   ; Test with undef aggregate
   ; CHECK-NEXT: OpCompositeInsert
   %4 = insertvalue [1 x ptr] undef, ptr undef, 0
+  ; CHECK: OpStore
+  store [1 x ptr] %4, ptr %0
 
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll b/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll
index 6e6cd2f68a971..510c7954c78f8 100644
--- a/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll
+++ b/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll
@@ -13,13 +13,18 @@
 
 %struct = type { [3 x float] }
 
+@G = global float 0.0
+
 define spir_kernel void @bar(i1 %sw) {
 entry:
   %var1 = alloca %struct
+  store %struct zeroinitializer, ptr %var1
   %var2 = alloca %struct
+  store %struct zeroinitializer, ptr %var2
   %elem1 = getelementptr inbounds [3 x float], ptr %var1, i64 0, i64 0
   %elem2 = getelementptr inbounds [3 x float], ptr %var2, i64 0, i64 1
   %elem = select i1 %sw, ptr %elem1, ptr %elem2
   %res = load float, ptr %elem
+  store float %res, ptr @G
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll b/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll
deleted file mode 100644
index efde6a2c082fc..0000000000000
--- a/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll
+++ /dev/null
@@ -1,23 +0,0 @@
-; This test case ensures that cleaning of temporary constants doesn't purge tracked ones.
-
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-
-; CHECK-SPIRV-DAG: %[[#Int:]] = OpTypeInt 8 0
-; CHECK-SPIRV-DAG: %[[#C0:]] = OpConstantNull %[[#Int]]
-; CHECK-SPIRV-DAG: %[[#C1:]] = OpConstant %[[#Int]] 1{{$}}
-
-define spir_kernel void @foo() {
-entry:
-  %addr = alloca i32
-  %r1 = call i8 @_Z20__spirv_SpecConstantia(i32 0, i8 1)
-  ; The name '%conv17.i' is important for the test case,
-  ; because it includes i32 0 when encoded for SPIR-V usage.
-  %conv17.i = sext i8 %r1 to i64
-  %tobool = trunc i8 %r1 to i1
-  %r2 = zext i1 %tobool to i32
-  store i32 %r2, ptr %addr
-  ret void
-}
-
-declare i8 @_Z20__spirv_SpecConstantia(i32, i8)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
index 3d2080e0050b7..691325251f11d 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
@@ -8,14 +8,15 @@
 
 %class.anon = type { i8 }
 
-define spir_func void @_Z3fooi(i32 %x) {
+define spir_func i32 @_Z3fooi(i32 %x) {
 entry:
   %x.addr = alloca i32, align 4
   store i32 %x, i32* %x.addr, align 4
-  %0 = load i32, i32* %x.addr, align 4
+  %0 = load i32, ptr %x.addr, align 4
   %cmp = icmp ne i32 %0, 0
   call void @llvm.assume(i1 %cmp)
-  ret void
+  %retval = select i1 %cmp, i32 100, i32 10
+  ret i32 %retval
 }
 
 declare void @llvm.assume(i1)
@@ -45,9 +46,9 @@ entry:
   call void @llvm.lifetime.start.p0i8(i64 4, i8* %0)
   store i32 1, i32* %a, align 4
   %1 = load i32, i32* %a, align 4
-  call spir_func void @_Z3fooi(i32 %1)
-  %2 = bitcast i32* %a to i8*
-  call void @llvm.lifetime.end.p0i8(i64 4, i8* %2)
+  %2 = call spir_func i32 @_Z3fooi(i32 %1)
+  %3 = bitcast i32* %a to i8*
+  call void @llvm.lifetime.end.p0i8(i64 4, i8* %3)
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
index 438fff6e94f89..18856147896bb 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
@@ -7,20 +7,20 @@
 ; CHECK: OpCapability ArbitraryPrecisionIntegersINTEL
 ; CHECK: OpExtension "SPV_INTEL_arbitrary_precision_integers"
 
-; CHECK: %[[#I4:]] = OpTypeInt 4 0
-; CHECK: %[[#I2:]] = OpTypeInt 2 0
-; CHECK: %[[#Z4:]] = OpConstantNull %[[#I4]]
-; CHECK: %[[#Z2:]] = OpConstantNull %[[#I2]]
-; CHECK: %[[#V2I2:]] = OpTypeVector %[[#I2]] 2
-; CHECK: %[[#V2I4:]] = OpTypeVector %[[#I4]] 2
-; CHECK: %[[#V3I2:]] = OpTypeVector %[[#I2]] 3
-; CHECK: %[[#V3I4:]] = OpTypeVector %[[#I4]] 3
-; CHECK: %[[#V4I2:]] = OpTypeVector %[[#I2]] 4
-; CHECK: %[[#V4I4:]] = OpTypeVector %[[#I4]] 4
-; CHECK: %[[#V8I2:]] = OpTypeVector %[[#I2]] 8
-; CHECK: %[[#V8I4:]] = OpTypeVector %[[#I4]] 8
-; CHECK: %[[#V16I2:]] = OpTypeVector %[[#I2]] 16
-; CHECK: %[[#V16I4:]] = OpTypeVector %[[#I4]] 16
+; CHECK-DAG: %[[#I4:]] = OpTypeInt 4 0
+; CHECK-DAG: %[[#I2:]] = OpTypeInt 2 0
+; CHECK-DAG: %[[#Z4:]] = OpConstantNull %[[#I4]]
+; CHECK-DAG: %[[#Z2:]] = OpConstantNull %[[#I2]]
+; CHECK-DAG: %[[#V2I2:]] = OpTypeVector %[[#I2]] 2
+; CHECK-DAG: %[[#V2I4:]] = OpTypeVector %[[#I4]] 2
+; CHECK-DAG: %[[#V3I2:]] = OpTypeVector %[[#I2]] 3
+; CHECK-DAG: %[[#V3I4:]] = OpTypeVector %[[#I4]] 3
+; CHECK-DAG: %[[#V4I2:]] = OpTypeVector %[[#I2]] 4
+; CHECK-DAG: %[[#V4I4:]] = OpTypeVector %[[#I4]] 4
+; CHECK-DAG: %[[#V8I2:]] = OpTypeVector %[[#I2]] 8
+; CHECK-DAG: %[[#V8I4:]] = OpTypeVector %[[#I4]] 8
+; CHECK-DAG: %[[#V16I2:]] = OpTypeVector %[[#I2]] 16
+; CHECK-DAG: %[[#V16I4:]] = OpTypeVector %[[#I4]] 16
 
 
 ; CHECK: %[[#]] = OpBitReverse %[[#I2]] %[[#Z2]]
@@ -36,45 +36,70 @@
 ; CHECK: %[[#]] = OpBitReverse %[[#V16I2]] %[[#]]
 ; CHECK: %[[#]] = OpBitReverse %[[#V16I4]] %[[#]]
 
+@G_i2_res = global i2 0
+@G_i4_res = global i4 0
+@G_v2i2_res = global <2 x i2> zeroinitializer
+@G_v2i4_res = global <2 x i4> zeroinitializer
+@G_v3i2_res = global <3 x i2> zeroinitializer
+@G_v3i4_res = global <3 x i4> zeroinitializer
+@G_v4i2_res = global <4 x i2> zeroinitializer
+@G_v4i4_res = global <4 x i4> zeroinitializer
+@G_v8i2_res = global <8 x i2> zeroinitializer
+@G_v8i4_res = global <8 x i4> zeroinitializer
+@G_v16i2_res = global <16 x i2> zeroinitializer
+@G_v16i4_res = global <16 x i4> zeroinitializer
+
 define spir_kernel void @testBitRev() {
 entry:
   %call2 = call i2 @llvm.bitreverse.i2(i2 0)
+  store i2 %call2, i2* @G_i2_res
   %call4 = call i4 @llvm.bitreverse.i4(i4 0)
+  store i4 %call4, i4* @G_i4_res
   ret void
 }
 
 define spir_kernel void @testBitRevV2(<2 x i2> %a, <2 x i4> %b) {
 entry:
   %call2 = call <2 x i2> @llvm.bitreverse.v2i2(<2 x i2> %a)
+  store <2 x i2> %call2, <2 x i2>* @G_v2i2_res
   %call4 = call <2 x i4> @llvm.bitreverse.v2i4(<2 x i4> %b)
+  store <2 x i4> %call4, <2 x i4>* @G_v2i4_res
   ret void
 }
 
 define spir_kernel void @testBitRevV3(<3 x i2> %a, <3 x i4> %b) {
 entry:
   %call2 = call <3 x i2> @llvm.bitreverse.v3i2(<3 x i2> %a)
+  store <3 x i2> %call2, <3 x i2>* @G_v3i2_res
   %call4 = call <3 x i4> @llvm.bitreverse.v3i4(<3 x i4> %b)
+  store <3 x i4> %call4, <3 x i4>* @G_v3i4_res
   ret void
 }
 
 define spir_kernel void @testBitRevV4(<4 x i2> %a, <4 x i4> %b) {
 entry:
   %call2 = call <4 x i2> @llvm.bitreverse.v4i2(<4 x i2> %a)
+  store <4 x i2> %call2, <4 x i2>* @G_v4i2_res
   %call4 = call <4 x i4> @llvm.bitreverse.v4i4(<4 x i4> %b)
+  store <4 x i4> %call4, <4 x i4>* @G_v4i4_res
   ret void
 }
 
 define spir_kernel void @testBitRevV8(<8 x i2> %a, <8 x i4> %b) {
 entry:
   %call2 = call <8 x i2> @llvm.bitreverse.v8i2(<8 x i2> %a)
+  store <8 x i2> %call2, <8 x i2>* @G_v8i2_res
   %call4 = call <8 x i4> @llvm.bitreverse.v8i4(<8 x i4> %b)
+  store <8 x i4> %call4, <8 x i4>* @G_v8i4_res
   ret void
 }
 
 define spir_kernel void @testBitRevV16(<16 x i2> %a, <16 x i4> %b) {
 entry:
   %call2 = call <16 x i2> @llvm.bitreverse.v16i2(<16 x i2> %a)
+  store <16 x i2> %call2, <16 x i2>* @G_v16i2_res
   %call4 = call <16 x i4> @llvm.bitreverse.v16i4(<16 x i4> %b)
+  store <16 x i4> %call4, <16 x i4>* @G_v16i4_res
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll
index 11bedfa605f9b..8e8e4df8fabc6 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll
@@ -23,15 +23,28 @@
 ; CHECK: OpExtInst %[[#]] %[[#]] fma %[[#]] %[[#]] %[[#]]
 ; CHECK: OpFRem
 
+@G_r1 = global float 0.0
+@G_r2 = global float 0.0
+@G_r3 = global float 0.0
+@G_r4 = global float 0.0
+@G_r5 = global float 0.0
+@G_r6 = global float 0.0
+
 ; Function Attrs: norecurse nounwind strictfp
 define dso_local spir_kernel void @test(float %a, i32 %in, i32 %ui) {
 entry:
   %r1 = tail call float @llvm.experimental.constrained.fadd.f32(float %a, float %a, metadata !"round.tonearest", metadata !"fpexcept.strict")
+  store float %r1, ptr @G_r1
   %r2 = tail call float @llvm.experimental.constrained.fdiv.f32(float %a, float %a, metadata !"round.towardzero", metadata !"fpexcept.strict")
+  store float %r2, ptr @G_r2
   %r3 = tail call float @llvm.experimental.constrained.fsub.f32(float %a, float %a, metadata !"round.upward", metadata !"fpexcept.strict")
+  store float %r3, ptr @G_r3
   %r4 = tail call float @llvm.experimental.constrained.fmul.f32(float %a, float %a, metadata !"round.downward", metadata !"fpexcept.strict")
+  store float %r4, ptr @G_r4
   %r5 = tail call float @llvm.experimental.constrained.fma.f32(float %a, float %a, float %a, metadata !"round.dynamic", metadata !"fpexcept.strict")
+  store float %r5, ptr @G_r5
   %r6 = tail call float @llvm.experimental.constrained.frem.f32(float %a, float %a, metadata !"round.dynamic", metadata !"fpexcept.strict")
+  store float %r6, ptr @G_r6
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
index f83cd8ad1969c..375da5b32e232 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
@@ -18,19 +18,20 @@
 ; CL:      %[[#FooVar:]] = OpVariable
 ; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
 ; CL-NEXT: OpLifetimeStart %[[#Casted1]] 16
-; CL-NEXT: OpBitcast
-; CL-NEXT: OpInBoundsPtrAccessChain
-; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
+; CL: OpInBoundsPtrAccessChain
+; CL: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
 ; CL-NEXT: OpLifetimeStop %[[#Casted2]] 16
 
 ; VK:      OpFunction
 ; VK:      %[[#FooVar:]] = OpVariable
 ; VK-NEXT: OpInBoundsAccessChain
+; VK-NEXT: OpStore
 ; VK-NEXT: OpReturn
 define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
   %RoundedRangeKernel = alloca %tprange, align 8
   call void @llvm.lifetime.start.p0(ptr nonnull %RoundedRangeKernel)
   %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
+  store i64 zeroinitializer, ptr %KernelFunc, align 8
   call void @llvm.lifetime.end.p0(ptr nonnull %RoundedRangeKernel)
   ret void
 }
@@ -39,37 +40,40 @@ define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange)
 ; CL: %[[#BarVar:]] = OpVariable
 ; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]]
 ; CL-NEXT: OpLifetimeStart %[[#Casted1]] 16
-; CL-NEXT: OpBitcast
-; CL-NEXT: OpInBoundsPtrAccessChain
-; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]]
+; CL: OpInBoundsPtrAccessChain
+; CL: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]]
 ; CL-NEXT: OpLifetimeStop %[[#Casted2]] 16
 
 ; VK:      OpFunction
 ; VK:      %[[#BarVar:]] = OpVariable
 ; VK-NEXT: OpInBoundsAccessChain
+; VK-NEXT: OpStore
 ; VK-NEXT: OpReturn
 define spir_func void @bar(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
   %RoundedRangeKernel = alloca %tprange, align 8
   call void @llvm.lifetime.start.p0(ptr nonnull %RoundedRangeKernel)
   %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
+  store i64 zeroinitializer, ptr %KernelFunc, align 8
   call void @llvm.lifetime.end.p0(ptr nonnull %RoundedRangeKernel)
   ret void
 }
 
 ; CL: OpFunction
 ; CL: %[[#TestVar:]] = OpVariable
-; CL-NEXT: OpLifetimeStart %[[#TestVar]] 1
-; CL-NEXT: OpInBoundsPtrAccessChain
-; CL-NEXT: OpLifetimeStop %[[#TestVar]] 1
+; CL: OpLifetimeStart %[[#TestVar]] 1
+; CL: OpInBoundsPtrAccessChain
+; CL: OpLifetimeStop %[[#TestVar]] 1
 
 ; VK:      OpFunction
 ; VK:      %[[#Test:]] = OpVariable
 ; VK-NEXT: OpInBoundsAccessChain
+; VK-NEXT: OpStore
 ; VK-NEXT: OpReturn
 define spir_func void @test(ptr noundef align 8 %_arg) {
   %var = alloca i8, align 8
   call void @llvm.lifetime.start.p0(ptr nonnull %var)
   %KernelFunc = getelementptr inbounds i8, ptr %var, i64 1
+  store i8 0, ptr %KernelFunc, align 8
   call void @llvm.lifetime.end.p0(ptr nonnull %var)
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll
index 08f15c077fed9..db930d1b28ec3 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll
@@ -9,29 +9,55 @@
 ; CHECK-DAG: OpName %[[#Bar:]] "bar"
 ; CHECK: %[[#Foo]] = OpFunction
 ; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
 ; CHECK: %[[#Bar]] = OpFunction
 ; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
+
+@G_r1_foo = global i16 0
+@G_r2_foo = global i16 0
+@G_r3_foo = global i16 0
+@G_r4_foo = global i16 0
+@G_r1_bar = global <4 x i32> zeroinitializer
+@G_r2_bar = global <4 x i32> zeroinitializer
+@G_r3_bar = global <4 x i32> zeroinitializer
+@G_r4_bar = global <4 x i32> zeroinitializer
 
 define spir_func void @foo(i16 %x, i16 %y) {
 entry:
   %r1 = tail call i16 @llvm.uadd.sat.i16(i16 %x, i16 %y)
+  store i16 %r1, ptr @G_r1_foo
   %r2 = tail call i16 @llvm.usub.sat.i16(i16 %x, i16 %y)
+  store i16 %r2, ptr @G_r2_foo
   %r3 = tail call i16 @llvm.sadd.sat.i16(i16 %x, i16 %y)
+  store i16 %r3, ptr @G_r3_foo
   %r4 = tail call i16 @llvm.ssub.sat.i16(i16 %x, i16 %y)
+  store i16 %r4, ptr @G_r4_foo
   ret void
 }
 
 define spir_func void @bar(<4 x i32> %x, <4 x i32> %y) {
 entry:
   %r1 = tail call <4 x i32> @llvm.uadd.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+  store <4 x i32> %r1, ptr @G_r1_bar
   %r2 = tail call <4 x i32> @llvm.usub.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+  store <4 x i32> %r2, ptr @G_r2_bar
   %r3 = tail call <4 x i32> @llvm.sadd.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+  store <4 x i32> %r3, ptr @G_r3_bar
   %r4 = tail call <4 x i32> @llvm.ssub.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+  store <4 x i32> %r4, ptr @G_r4_bar
   ret void
 }
+
+declare i16 @llvm.uadd.sat.i16(i16, i16)
+declare i16 @llvm.usub.sat.i16(i16, i16)
+declare i16 @llvm.sadd.sat.i16(i16, i16)
+declare i16 @llvm.ssub.sat.i16(i16, i16)
+declare <4 x i32> @llvm.uadd.sat.v4i32(<4 x i32>, <4 x i32>)
+declare <4 x i32> @llvm.usub.sat.v4i32(<4 x i32>, <4 x i32>)
+declare <4 x i32> @llvm.sadd.sat.v4i32(<4 x i32>, <4 x i32>)
+declare <4 x i32> @llvm.ssub.sat.v4i32(<4 x i32>, <4 x i32>)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll
index 08e429f36827c..54cb096da8d89 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll
@@ -90,12 +90,13 @@ define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, ptr %p)
 ; CHECK: OpIAddCarry %[[StructLong]]
 ; CHECK: OpIAddCarry %[[StructLong]]
 ; CHECK: OpReturn
-define void @foo(i64 %a, i64 %b) {
+define i64 @foo(i64 %a, i64 %b) {
   %r1 = call { i64, i1 } @llvm.uadd.with.overflow.i64(i64 %a, i64 %b)
   %r2 = call { i64, i1 } @llvm.uadd.with.overflow.i64(i64 %a, i64 %b)
   %d1 = extractvalue  { i64, i1 } %r1, 0
   %d2 = extractvalue  { i64, i1 } %r2, 0
-  ret void
+  %sum = add i64 %d1, %d2
+  ret i64 %sum
 }
 
 declare {i8, i1} @llvm.uadd.with.overflow.i8(i8, i8)
diff --git a/llvm/test/CodeGen/SPIRV/logical-access-chain.ll b/llvm/test/CodeGen/SPIRV/logical-access-chain.ll
index d56678ecfc2c9..e96ebf777c28f 100644
--- a/llvm/test/CodeGen/SPIRV/logical-access-chain.ll
+++ b/llvm/test/CodeGen/SPIRV/logical-access-chain.ll
@@ -2,6 +2,7 @@
 
 ; CHECK-DAG:      [[uint:%[0-9]+]] = OpTypeInt 32 0
 ; CHECK-DAG:     [[uint2:%[0-9]+]] = OpTypeVector [[uint]] 2
+; CHECK-DAG:    [[uint_0:%[0-9]+]] = OpConstant [[uint]] 0
 ; CHECK-DAG:    [[uint_1:%[0-9]+]] = OpConstant [[uint]] 1
 ; CHECK-DAG:  [[ptr_uint:%[0-9]+]] = OpTypePointer Function [[uint]]
 ; CHECK-DAG: [[ptr_uint2:%[0-9]+]] = OpTypePointer Function [[uint2]]
@@ -12,7 +13,9 @@ entry:
 ; CHECK: [[var:%[0-9]+]] = OpVariable [[ptr_uint2]] Function
 
   %1 = getelementptr <2 x i32>, ptr %0, i32 0, i32 1
-; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[var]] [[uint_1]]
+; CHECK: [[gep:%[0-9]+]] = OpAccessChain [[ptr_uint]] [[var]] [[uint_1]]
+  store i32 0, ptr %1
+; CHECK: OpStore [[gep]] [[uint_0]]
 
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/logical-struct-access.ll b/llvm/test/CodeGen/SPIRV/logical-struct-access.ll
index 66337b1ba2b37..518e011bf0be2 100644
--- a/llvm/test/CodeGen/SPIRV/logical-struct-access.ll
+++ b/llvm/test/CodeGen/SPIRV/logical-struct-access.ll
@@ -1,5 +1,4 @@
-; RUN: llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - -print-after-all | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - | FileCheck %s
 
 ; CHECK-DAG: [[uint:%[0-9]+]] = OpTypeInt 32 0
 
@@ -24,35 +23,85 @@
 ; CHECK-DAG:    [[ptr_A:%[0-9]+]] = OpTypePointer Function [[A]]
 ; CHECK-DAG:    [[ptr_B:%[0-9]+]] = OpTypePointer Function [[B]]
 
-define void @main() #1 {
-entry:
-  %0 = alloca %B, align 4
-; CHECK: [[tmp:%[0-9]+]] = OpVariable [[ptr_B]] Function
-
-  %1 = getelementptr %B, ptr %0, i32 0, i32 0
+define internal ptr @gep_B_0(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_A]] [[tmp]] [[uint_0]]
-  %2 = getelementptr inbounds %B, ptr %0, i32 0, i32 0
+  %res = getelementptr %B, ptr %base, i32 0, i32 0
+  ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_0(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_A]] [[tmp]] [[uint_0]]
+  %res = getelementptr inbounds %B, ptr %base, i32 0, i32 0
+  ret ptr %res
+}
 
-  %3 = getelementptr %B, ptr %0, i32 0, i32 1
+define internal ptr @gep_B_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[tmp]] [[uint_1]]
-  %4 = getelementptr inbounds %B, ptr %0, i32 0, i32 1
+  %res = getelementptr %B, ptr %base, i32 0, i32 1
+  ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_uint]] [[tmp]] [[uint_1]]
+  %res = getelementptr inbounds %B, ptr %base, i32 0, i32 1
+  ret ptr %res
+}
 
-  %5 = getelementptr %B, ptr %0, i32 0, i32 2
+define internal ptr @gep_B_2(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_A]] [[tmp]] [[uint_2]]
-  %6 = getelementptr inbounds %B, ptr %0, i32 0, i32 2
+  %res = getelementptr %B, ptr %base, i32 0, i32 2
+  ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_2(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_A]] [[tmp]] [[uint_2]]
+  %res = getelementptr inbounds %B, ptr %base, i32 0, i32 2
+  ret ptr %res
+}
 
-  %7 = getelementptr %B, ptr %0, i32 0, i32 2, i32 1
+define internal ptr @gep_B_2_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[tmp]] [[uint_2]] [[uint_1]]
-  %8 = getelementptr inbounds %B, ptr %0, i32 0, i32 2, i32 1
+  %res = getelementptr %B, ptr %base, i32 0, i32 2, i32 1
+  ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_2_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_uint]] [[tmp]] [[uint_2]] [[uint_1]]
+  %res = getelementptr inbounds %B, ptr %base, i32 0, i32 2, i32 1
+  ret ptr %res
+}
 
-  %9 = getelementptr %B, ptr %0, i32 0, i32 2
-  %10 = getelementptr %A, ptr %9, i32 0, i32 1
+define internal ptr @gep_B_2_A_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
 ; CHECK: [[x:%[0-9]+]] = OpAccessChain [[ptr_A]] [[tmp]] [[uint_2]]
 ; CHECK:   {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[x]] [[uint_1]]
+  %x = getelementptr %B, ptr %base, i32 0, i32 2
+  %res = getelementptr %A, ptr %x, i32 0, i32 1
+  ret ptr %res
+}
+
+define void @main() #1 {
+entry:
+  %0 = alloca %B, align 4
+; CHECK: [[tmp:%[0-9]+]] = OpVariable [[ptr_B]] Function
+
+  %1 = call ptr @gep_B_0(ptr %0)
+  %2 = call ptr @gep_inbounds_B_0(ptr %0)
+  %3 = call ptr @gep_B_1(ptr %0)
+  %4 = call ptr @gep_inbounds_B_1(ptr %0)
+  %5 = call ptr @gep_B_2(ptr %0)
+  %6 = call ptr @gep_inbounds_B_2(ptr %0)
+  %7 = call ptr @gep_B_2_1(ptr %0)
+  %8 = call ptr @gep_inbounds_B_2_1(ptr %0)
+  %10 = call ptr @gep_B_2_A_1(ptr %0)
 
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/phi-insert-point.ll b/llvm/test/CodeGen/SPIRV/phi-insert-point.ll
index 70d121cdf4b3a..a34186d491257 100644
--- a/llvm/test/CodeGen/SPIRV/phi-insert-point.ll
+++ b/llvm/test/CodeGen/SPIRV/phi-insert-point.ll
@@ -36,9 +36,18 @@ ok:
   br label %exit
 
 exit:
+  store i64 %r1, ptr @g1
+  store i64 %r2, ptr @g2
+  store ptr addrspace(4) %r3, ptr @g3
+  store ptr addrspace(4) %r4, ptr @g4
   ret void
 }
 
+@g1 = internal global i64 0
+@g2 = internal global i64 0
+@g3 = internal global ptr addrspace(4) null
+@g4 = internal global ptr addrspace(4) null
+
 define spir_kernel void @bar(i64 %arg_val, i64 %arg_val_def, ptr addrspace(4) byval(%struct) %arg_ptr, ptr addrspace(4) %arg_ptr_def) {
 entry:
   %fl = icmp eq i64 %arg_val, 0
@@ -55,5 +64,9 @@ ok:
   br label %exit
 
 exit:
+  store i64 %r1, ptr @g1
+  store i64 %r2, ptr @g2
+  store ptr addrspace(4) %r3, ptr @g3
+  store ptr addrspace(4) %r4, ptr @g4
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll b/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll
index bc090ce55fbec..c250ebae12746 100644
--- a/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll
+++ b/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll
@@ -20,11 +20,14 @@
 ; CHECK: %[[#Case1]] = OpFunction
 define spir_func void @case1(i1 %b1, i1 %b2, i1 %b3) {
 entry:
+  %tmp.1 = alloca i8, align 1
 ; CHECK: OpBranchConditional %[[#]] %[[#l1:]] %[[#l2:]]
   br i1 %b1, label %l1, label %l2
 
 l1:
   %str = phi ptr addrspace(1) [ @.str.1, %entry ], [ @.str.2, %l2 ], [ @.str.2, %l3 ]
+  %v1 = load i8, ptr addrspace(1) %str, align 1
+  store i8 %v1, ptr %tmp.1, align 1
   br label %exit
 
 ; CHECK: %[[#l2]] = OpLabel
@@ -51,11 +54,14 @@ exit:
 ; CHECK: %[[#Case2]] = OpFunction
 define spir_func void @case2(i1 %b1, i1 %b2, i1 %b3, ptr addrspace(1) byval(%struct1) %str1, ptr addrspace(1) byval(%struct2) %str2) {
 entry:
+  %tmp.2 = alloca i8, align 1
 ; CHECK: OpBranchConditional %[[#]] %[[#l1:]] %[[#l2:]]
   br i1 %b1, label %l1, label %l2
 
 l1:
   %str = phi ptr addrspace(1) [ %str1, %entry ], [ %str2, %l2 ], [ %str2, %l3 ]
+  %v2 = load i8, ptr addrspace(1) %str, align 1
+  store i8 %v2, ptr %tmp.2, align 1
   br label %exit
 
 ; CHECK: %[[#l2]] = OpLabel
@@ -83,10 +89,13 @@ define spir_func void @case3(i1 %b1, i1 %b2, i1 %b3, ptr addrspace(1) byval(%str
 
 ; CHECK: OpBranchConditional %[[#]] %[[#l1:]] %[[#l2:]]
 entry:
+  %tmp.3 = alloca i8, align 1
   br i1 %b1, label %l1, label %l2
 
 l1:
   %str = phi ptr addrspace(1) [ %_arg_str1, %entry ], [ %str2, %l2 ], [ %str3, %l3 ]
+  %v3 = load i8, ptr addrspace(1) %str, align 1
+  store i8 %v3, ptr %tmp.3, align 1
   br label %exit
 
 ; CHECK: %[[#l2]] = OpLabel
diff --git a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll
index 7db1eed84bf7d..3382987bbd581 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll
@@ -26,9 +26,13 @@
 %struct.S = type { i32 }
 %struct.__wrapper_class = type { [7 x %struct.S] }
 
+@G_elem = global ptr null
+@G_data = global i64 0
+
 define spir_kernel void @foo1(ptr noundef byval(%struct.__wrapper_class) align 4 %_arg_Arr) {
 entry:
   %elem = getelementptr inbounds i8, ptr %_arg_Arr, i64 0
+  store ptr %elem, ptr @G_elem
   ret void
 }
 
@@ -36,5 +40,6 @@ define spir_kernel void @foo2(ptr noundef byval(%struct.__wrapper_class) align 4
 entry:
   %elem = getelementptr inbounds %struct.__wrapper_class, ptr %_arg_Arr, i64 0
   %data = load i64, ptr %elem
+  store i64 %data, ptr @G_data
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll
index d6a0071167cef..ed5652a750582 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll
@@ -14,8 +14,11 @@
 %struct.S = type { i32 }
 %struct.__wrapper_class = type { [7 x %struct.S] }
 
+@G = global i32 0
+
 define spir_kernel void @foo(ptr noundef byval(%struct.__wrapper_class) align 4 %_arg_Arr) {
 entry:
   %val = load i32, ptr %_arg_Arr
+  store i32 %val, ptr @G
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll b/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll
index 0e2730e18bf38..e47aa61a8acd7 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll
@@ -30,6 +30,8 @@
 %"class.std::complex" = type { { double, double } }
 %class.anon = type { i32, ptr addrspace(4), [2 x [2 x %"class.std::complex"]] }
 
+@G = global ptr addrspace(4) null
+
 define weak_odr dso_local spir_kernel void @foo(i32 noundef %_arg_N, ptr addrspace(1) noundef align 8 %_arg_p) {
 entry:
   %Kernel = alloca %class.anon, align 8
@@ -38,5 +40,6 @@ entry:
   %r0 = addrspacecast ptr addrspace(1) %_arg_p to ptr addrspace(4)
   store ptr addrspace(4) %r0, ptr %p, align 8
   %r3 = load ptr addrspace(4), ptr %p, align 8
+  store ptr addrspace(4) %r3, ptr @G
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll
index 7a09ac973b590..0e397ec51caaa 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll
@@ -7,9 +7,14 @@
 ; CHECK:  %[[#]] = OpInBoundsPtrAccessChain %[[#PTR1]] %[[#]] %[[#]]
 ; CHECK:  %[[#]] = OpInBoundsPtrAccessChain %[[#PTR2]] %[[#]] %[[#]]
 
+@G_c = global ptr addrspace(1) null
+@G_d = global ptr addrspace(2) null
+
 define spir_kernel void @foo(ptr addrspace(1) %a, ptr addrspace(2) %b) {
 entry:
   %c = getelementptr inbounds i8, ptr addrspace(1) %a, i32 1
+  store ptr addrspace(1) %c, ptr @G_c
   %d = getelementptr inbounds i8, ptr addrspace(2) %b, i32 2
+  store ptr addrspace(2) %d, ptr @G_d
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll
index c822dbc5d6c0e..e12a809125248 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll
@@ -7,9 +7,12 @@
 ; CHECK:  %[[#GEP:]] = OpInBoundsPtrAccessChain %[[#PTR]] %[[#ARG]] %[[#]]
 ; CHECK:  %[[#]] = OpLoad %[[#FLOAT32]] %[[#GEP]] Aligned 4
 
+@G = global float 0.0
+
 define spir_kernel void @test1(ptr addrspace(1) %arg1) !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_type_qual !4 {
   %a = getelementptr inbounds float, ptr addrspace(1) %arg1, i64 1
   %b = load float, ptr addrspace(1) %a, align 4
+  store float %b, ptr @G
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll
index 1d846a35a65aa..859253e5b18d9 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll
@@ -7,6 +7,9 @@
 ; CHECK-DAG: %[[#PTR_VEC3:]] = OpTypePointer CrossWorkgroup %[[#VEC3]]
 ; CHECK-DAG: %[[#PTR_VEC4:]] = OpTypePointer CrossWorkgroup %[[#VEC4]]
 
+@G_loadv1 = global <4 x i8> zeroinitializer
+@G_loadv2 = global <4 x i8> zeroinitializer
+
 ; CHECK: %[[#AC1:]] = OpInBoundsPtrAccessChain %[[#PTR_VEC3]] %[[#]] %[[#]]
 ; CHECK: %[[#BC1:]] = OpBitcast %[[#PTR_VEC4]] %[[#AC1]]
 ; CHECK: %[[#LD1:]] = OpLoad %[[#VEC4]] %[[#BC1]] Aligned 4
@@ -15,6 +18,7 @@
 define spir_kernel void @foo(ptr addrspace(1) %a, i64 %b) {
   %index = getelementptr inbounds <3 x i8>, ptr addrspace(1) %a, i64 %b
   %loadv = load <4 x i8>, ptr addrspace(1) %index, align 4
+  store <4 x i8> %loadv, ptr @G_loadv1
   ret void
 }
 
@@ -29,5 +33,6 @@ define spir_kernel void @bar(ptr addrspace(1) %a, i64 %b) {
 ; from older LLVM IR with typed pointers.
   %cast = bitcast ptr addrspace(1) %index to ptr addrspace(1)
   %loadv = load <4 x i8>, ptr addrspace(1) %cast, align 4
+  store <4 x i8> %loadv, ptr @G_loadv2
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll
index a5e891dae6f11..3ae03edf5200f 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll
@@ -7,11 +7,15 @@
 ; CHECK-DAG: %[[#PTRINT8:]] = OpTypePointer Workgroup %[[#INT8]]
 ; CHECK-DAG: %[[#CONST:]] = OpConstant %[[#INT64]] 1
 
+@G_gep1 = global ptr addrspace(3) null
+@G_gep2 = global ptr addrspace(3) null
+
 ; CHECK: %[[#PARAM1:]] = OpFunctionParameter %[[#PTRINT8]]
 define spir_kernel void @test1(ptr addrspace(3) %address) {
 ; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#PTRINT8]] %[[#PARAM1]] %[[#CONST]]
   %cast = bitcast ptr addrspace(3) %address to ptr addrspace(3)
   %gep = getelementptr inbounds i8, ptr addrspace(3) %cast, i64 1
+  store ptr addrspace(3) %gep, ptr @G_gep1
   ret void
 }
 
@@ -19,5 +23,6 @@ define spir_kernel void @test1(ptr addrspace(3) %address) {
 define spir_kernel void @test2(ptr addrspace(3) %address) {
 ; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#PTRINT8]] %[[#PARAM2]] %[[#CONST]]
   %gep = getelementptr inbounds i8, ptr addrspace(3) %address, i64 1
+  store ptr addrspace(3) %gep, ptr @G_gep2
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
index 19451d23c6830..39563aecafec4 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
@@ -7,13 +7,16 @@
 ; CHECK-DAG: %[[#value:]] = OpConstant %[[#type]] 456
 ; CHECK-DAG:   %[[#var:]] = OpVariable %[[#ptrty]] Private %[[#value]]
 
+@G = internal global i32 0
+
 define hidden spir_func void @Foo() {
   %p = addrspacecast ptr addrspace(10) @PrivInternal to ptr
   %v = load i32, ptr %p, align 4
+  store i32 %v, ptr @G
   ret void
 ; CHECK:      OpLabel
-; CHECK-NEXT: OpLoad %[[#type]] %[[#var]] Aligned 4
-; CHECK-Next: OpReturn
+; CHECK: OpLoad %[[#type]] %[[#var]] Aligned 4
+; CHECK: OpReturn
 }
 
 define void @main() #1 {
diff --git a/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll b/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll
index b3c68d22f9bdd..681fb70ad706d 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll
@@ -9,9 +9,14 @@
 ; CHECK:  %[[#]] = OpLoad %[[#INT8]] %[[#FNP1]] Aligned 1
 ; CHECK:  %[[#]] = OpLoad %[[#INT8]] %[[#FNP2]] Aligned 1
 
+@G_c = global i8 0
+@G_d = global i8 0
+
 define spir_kernel void @foo(ptr addrspace(1) %a, ptr addrspace(2) %b) {
 entry:
   %c = load i8, ptr addrspace(1) %a
+  store i8 %c, ptr @G_c
   %d = load i8, ptr addrspace(2) %b
+  store i8 %d, ptr @G_d
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll b/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll
index a9e79df259c4f..44134f83cfec3 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll
@@ -51,6 +51,7 @@ l1:
 l2:
   %val2 = phi ptr addrspace(4) [ %p, %l1 ], [ %val3, %l3 ]
   %val1 = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @G1 to ptr addrspace(4)), %l1 ], [ %val2, %l3 ]
+  store i16 0, ptr addrspace(4) %val1, align 2
   br i1 %f2, label %l3, label %exit
 
 l3:
@@ -75,6 +76,7 @@ l1:
 l2:
   %val1 = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @G1 to ptr addrspace(4)), %l1 ], [ %val2, %l3 ]
   %val2 = phi ptr addrspace(4) [ %p, %l1 ], [ %val3, %l3 ]
+  store i16 0, ptr addrspace(4) %val1, align 2
   br i1 %f2, label %l3, label %exit
 
 exit:
diff --git a/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll
index 4d5549dfab8d9..123daa411810b 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll
@@ -10,6 +10,7 @@
 ; CHECK-DAG:  OpName %[[#func_chain:]] "chain"
 
 @global = internal addrspace(10) global i32 zeroinitializer
+@G = global i32 0
 
 define void @simple() {
 ; CHECK: %[[#func_simple]] = OpFunction
@@ -17,6 +18,7 @@ entry:
   %ptr = getelementptr i32, ptr addrspace(10) @global, i32 0
   %casted = addrspacecast ptr addrspace(10) %ptr to ptr
   %val = load i32, ptr %casted
+  store i32 %val, ptr @G
 ; CHECK: %{{.*}} = OpLoad %[[#uint]] %[[#var]] Aligned 4
   ret void
 }
@@ -31,6 +33,7 @@ entry:
   %e = addrspacecast ptr addrspace(10) %d to ptr
 
   %val = load i32, ptr %e
+  store i32 %val, ptr @G
 ; CHECK: %{{.*}} = OpLoad %[[#uint]] %[[#var]] Aligned 4
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll b/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll
index 876cd3c20cf35..80ee36cfe15d2 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll
@@ -15,6 +15,9 @@
 ; CHECK: OpGenericCastToPtr
 ; CHECK: OpPtrEqual
 
+@G_b1 = global i1 0
+@G_b2 = global i1 0
+
 define spir_kernel void @foo(ptr addrspace(3) align 4 %_arg_local, ptr addrspace(1) align 4 %_arg_global) {
 entry:
   %p1 = getelementptr inbounds i32, ptr addrspace(1) %_arg_global, i64 0
@@ -24,9 +27,12 @@ entry:
   %p4 = addrspacecast ptr addrspace(1) %p3 to ptr addrspace(4)
   %p5 = tail call spir_func ptr addrspace(3) @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPvi(ptr addrspace(4) %p4, i32 4)
   %b1 = icmp eq ptr addrspace(3) %p5, null
+  store i1 %b1, ptr @G_b1
   %p6 = getelementptr inbounds i32, ptr addrspace(3) %p5, i64 0
   %p7 = tail call spir_func ptr addrspace(3) @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPvi(ptr addrspace(4) %p4, i32 4)
   %b2 = icmp eq ptr addrspace(3) %p7, null
+  store i1 %b2, ptr @G_b2
+  store ptr addrspace(3) %p6, ptr addrspace(3) %p2
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll b/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll
index 7548f4757dbe6..6fc03a386d14d 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll
@@ -4,18 +4,23 @@
 
 @.str = private unnamed_addr constant [7 x i8] c"buffer\00", align 1
 
+; The i64 values in the extracts will be turned
+; into immidiate values. There should be no 64-bit
+; integers in the module.
+; CHECK-NOT:  OpTypeInt 64 0
+
 define void @main() "hlsl.shader"="pixel"  {
-; CHECK:         %24 = OpFunction %2 None %3 ; -- Begin function main
-; CHECK-NEXT:     %1 = OpLabel
-; CHECK-NEXT:    %25 = OpVariable %13 Function %22
-; CHECK-NEXT:    %26 = OpLoad %7 %23
-; CHECK-NEXT:    %27 = OpImageRead %5 %26 %15
-; CHECK-NEXT:    %28 = OpCompositeExtract %4 %27 0
-; CHECK-NEXT:    %29 = OpCompositeExtract %4 %27 1
-; CHECK-NEXT:    %30 = OpFAdd %4 %29 %28
-; CHECK-NEXT:    %31 = OpCompositeInsert %5 %30 %27 0
-; CHECK-NEXT:    %32 = OpLoad %7 %23
-; CHECK-NEXT:          OpImageWrite %32 %15 %31
+; CHECK:         %[[FUNC:[0-9]+]] = OpFunction %[[VOID:[0-9]+]] None %[[FNTYPE:[0-9]+]] ; -- Begin function main
+; CHECK-NEXT:     %[[LABEL:[0-9]+]] = OpLabel
+; CHECK-NEXT:    %[[VAR:[0-9]+]] = OpVariable %[[PTR_FN:[a-zA-Z0-9_]+]] Function %[[INIT:[a-zA-Z0-9_]+]]
+; CHECK-NEXT:    %[[LOAD1:[0-9]+]] = OpLoad %[[IMG_TYPE:[a-zA-Z0-9_]+]] %[[IMG_VAR:[a-zA-Z0-9_]+]]
+; CHECK-NEXT:    %[[READ:[0-9]+]] = OpImageRead %[[VEC4:[a-zA-Z0-9_]+]] %[[LOAD1]] %[[COORD:[a-zA-Z0-9_]+]]
+; CHECK-NEXT:    %[[EXTRACT1:[0-9]+]] = OpCompositeExtract %[[FLOAT:[a-zA-Z0-9_]+]] %[[READ]] 0
+; CHECK-NEXT:    %[[EXTRACT2:[0-9]+]] = OpCompositeExtract %[[FLOAT]] %[[READ]] 1
+; CHECK-NEXT:    %[[ADD:[0-9]+]] = OpFAdd %[[FLOAT]] %[[EXTRACT2]] %[[EXTRACT1]]
+; CHECK-NEXT:    %[[INSERT:[0-9]+]] = OpCompositeInsert %[[VEC4]] %[[ADD]] %[[READ]] 0
+; CHECK-NEXT:    %[[LOAD2:[0-9]+]] = OpLoad %[[IMG_TYPE]] %[[IMG_VAR]]
+; CHECK-NEXT:          OpImageWrite %[[LOAD2]] %[[COORD]] %[[INSERT]]
 ; CHECK-NEXT:          OpReturn
 ; CHECK-NEXT:          OpFunctionEnd
 entry:
diff --git a/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll b/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll
index 101116f437811..7409b3db51948 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll
@@ -34,6 +34,8 @@
 
 %class.CustomType = type { i64 }
 
+@G = global ptr addrspace(4) null
+
 define linkonce_odr dso_local spir_func void @bar(ptr addrspace(4) noundef %first) {
 entry:
   %first.addr = alloca ptr addrspace(4)
@@ -44,6 +46,7 @@ entry:
   call spir_func void @foo(i64 noundef 100, ptr addrspace(4) noundef dereferenceable(8) %first.addr.ascast, ptr addrspace(4) noundef dereferenceable(8) %temp.ascast)
   call spir_func void @foo(i64 noundef 100, ptr addrspace(4) noundef dereferenceable(8) %temp.ascast, ptr addrspace(4) noundef dereferenceable(8) %first.addr.ascast)
   %var = alloca ptr addrspace(4), align 8
+  store ptr addrspace(4) null, ptr %var
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll b/llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll
new file mode 100644
index 0000000000000..6bd640f813142
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll
@@ -0,0 +1,31 @@
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - -filetype=obj | spirv-val %}
+
+%A = type {
+  i32,
+  i32
+}
+
+%B = type {
+  %A,
+  i32,
+  %A
+}
+
+; Make sure all struct types are removed.
+; CHECK-NOT: OpTypeStruct
+
+; Make sure the GEPs and the function scope variable are removed.
+; CHECK: OpFunction
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: OpReturn
+; CHECK-NEXT: OpFunctionEnd
+define void @main() #1 {
+entry:
+  %0 = alloca %B, align 4
+  %1 = getelementptr %B, ptr %0, i32 0, i32 2
+  %2 = getelementptr %A, ptr %1, i32 0, i32 1
+  ret void
+}
+
+attributes #1 = { "hlsl.numthreads"="4,8,16" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll
index 481bad9a26b7b..280f586891717 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll
@@ -19,10 +19,15 @@
 ; TODO: Add a check to ensure that there's no behavior change of bitreverse operation
 ;       between the LLVM-IR and SPIR-V for i2 and i4
 
+@G_res2 = global i2 0
+@G_res4 = global i4 0
+
 define spir_func void @foo(i2 %a, i4 %b) {
 entry:
   %res2 = tail call i2 @llvm.bitreverse.i2(i2 %a)
+  store i2 %res2, ptr @G_res2
   %res4 = tail call i4 @llvm.bitreverse.i4(i4 %b)
+  store i4 %res4, ptr @G_res4
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll
index 119dbe14446c1..68f33510b6a8d 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll
@@ -45,6 +45,12 @@ entry:
   %GE = call spir_func ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) %var1, i32 5)
   %LE = call spir_func ptr addrspace(3) @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPvi(ptr addrspace(4) %var2, i32 4)
   %PE = call spir_func ptr @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePvi(ptr addrspace(4) %var3, i32 7)
+  store i32 0, ptr addrspace(1) %G, align 4
+  store i8 0, ptr addrspace(3) %L, align 1
+  store i32 0, ptr %P, align 4
+  store i32 0, ptr addrspace(1) %GE, align 4
+  store i8 0, ptr addrspace(3) %LE, align 1
+  store i32 0, ptr %PE, align 4
   ret void
 }
 
@@ -70,6 +76,9 @@ entry:
   %G = call spir_func ptr addrspace(1) @_Z9to_globalPv(ptr addrspace(4) %var1)
   %L = call spir_func ptr addrspace(3) @_Z8to_localPv(ptr addrspace(4) %var2)
   %P = call spir_func ptr @_Z10to_privatePv(ptr addrspace(4) %var3)
+  store i32 0, ptr addrspace(1) %G, align 4
+  store i8 0, ptr addrspace(3) %L, align 1
+  store i32 0, ptr %P, align 4
   ret void
 }
 
@@ -114,6 +123,12 @@ entry:
   %GE = call spir_func ptr addrspace(1) @__spirv_GenericCastToPtrExplicit_ToGlobal(ptr addrspace(4) %var1, i32 5)
   %LE = call spir_func ptr addrspace(3) @__spirv_GenericCastToPtrExplicit_ToLocal(ptr addrspace(4) %var2, i32 4)
   %PE = call spir_func ptr @__spirv_GenericCastToPtrExplicit_ToPrivate(ptr addrspace(4) %var3, i32 7)
+  store i32 0, ptr addrspace(1) %G, align 4
+  store i8 0, ptr addrspace(3) %L, align 1
+  store i32 0, ptr %P, align 4
+  store i32 0, ptr addrspace(1) %GE, align 4
+  store i8 0, ptr addrspace(3) %LE, align 1
+  store i32 0, ptr %PE, align 4
   ret void
 }
 
@@ -139,6 +154,9 @@ entry:
   %G = call spir_func ptr addrspace(1) @to_global(ptr addrspace(4) %var1)
   %L = call spir_func ptr addrspace(3) @to_local(ptr addrspace(4) %var2)
   %P = call spir_func ptr @to_private(ptr addrspace(4) %var3)
+  store i32 0, ptr addrspace(1) %G, align 4
+  store i8 0, ptr addrspace(3) %L, align 1
+  store i32 0, ptr %P, align 4
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll
index 818243ab19e41..9f08a65c16866 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll
@@ -16,9 +16,13 @@
 ; CHECK-SPIRV: OpGenericCastToPtr %[[#LocalCharPtr]] %[[#Ptr2]]
 ; CHECK-SPIRV: OpFunctionEnd
 
+@G_p = global ptr addrspace(3) null
+@G_p2 = global ptr addrspace(3) null
+
 define spir_kernel void @foo(ptr addrspace(1) %arg) {
 entry:
   %p = addrspacecast ptr addrspace(1) %arg to ptr addrspace(3)
+  store ptr addrspace(3) %p, ptr @G_p
   ret void
 }
 
@@ -26,5 +30,6 @@ define spir_kernel void @bar(ptr addrspace(1) %arg) {
 entry:
   %p1 = addrspacecast ptr addrspace(1) %arg to ptr addrspace(4)
   %p2 = addrspacecast ptr addrspace(4) %p1 to ptr addrspace(3)
+  store ptr addrspace(3) %p2, ptr @G_p2
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll b/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll
index 46eaba9d5ceb1..c752e278927a9 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll
@@ -184,6 +184,8 @@
 ; CHECK-SPIRV: %[[#r89]] = OpUnordered %[[#bool]]
 ; CHECK-SPIRV: %[[#r90]] = OpUnordered %[[#bool]]
 
+@G = global [90 x i1] zeroinitializer
+
 define spir_kernel void @testFCmp(float %a, float %b) local_unnamed_addr {
 entry:
   %r1 = fcmp oeq float %a, %b
@@ -276,5 +278,185 @@ entry:
   %r88 = fcmp uno float %a, %b
   %r89 = fcmp ninf uno float %a, %b
   %r90 = fcmp nsz uno float %a, %b
+  %p1 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 0
+  store i1 %r1, ptr %p1
+  %p2 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 1
+  store i1 %r2, ptr %p2
+  %p3 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 2
+  store i1 %r3, ptr %p3
+  %p4 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 3
+  store i1 %r4, ptr %p4
+  %p5 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 4
+  store i1 %r5, ptr %p5
+  %p6 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 5
+  store i1 %r6, ptr %p6
+  %p7 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 6
+  store i1 %r7, ptr %p7
+  %p8 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 7
+  store i1 %r8, ptr %p8
+  %p9 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 8
+  store i1 %r9, ptr %p9
+  %p10 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 9
+  store i1 %r10, ptr %p10
+  %p11 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 10
+  store i1 %r11, ptr %p11
+  %p12 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 11
+  store i1 %r12, ptr %p12
+  %p13 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 12
+  store i1 %r13, ptr %p13
+  %p14 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 13
+  store i1 %r14, ptr %p14
+  %p15 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 14
+  store i1 %r15, ptr %p15
+  %p16 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 15
+  store i1 %r16, ptr %p16
+  %p17 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 16
+  store i1 %r17, ptr %p17
+  %p18 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 17
+  store i1 %r18, ptr %p18
+  %p19 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 18
+  store i1 %r19, ptr %p19
+  %p20 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 19
+  store i1 %r20, ptr %p20
+  %p21 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 20
+  store i1 %r21, ptr %p21
+  %p22 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 21
+  store i1 %r22, ptr %p22
+  %p23 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 22
+  store i1 %r23, ptr %p23
+  %p24 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 23
+  store i1 %r24, ptr %p24
+  %p25 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 24
+  store i1 %r25, ptr %p25
+  %p26 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 25
+  store i1 %r26, ptr %p26
+  %p27 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 26
+  store i1 %r27, ptr %p27
+  %p28 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 27
+  store i1 %r28, ptr %p28
+  %p29 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 28
+  store i1 %r29, ptr %p29
+  %p30 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 29
+  store i1 %r30, ptr %p30
+  %p31 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 30
+  store i1 %r31, ptr %p31
+  %p32 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 31
+  store i1 %r32, ptr %p32
+  %p33 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 32
+  store i1 %r33, ptr %p33
+  %p34 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 33
+  store i1 %r34, ptr %p34
+  %p35 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 34
+  store i1 %r35, ptr %p35
+  %p36 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 35
+  store i1 %r36, ptr %p36
+  %p37 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 36
+  store i1 %r37, ptr %p37
+  %p38 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 37
+  store i1 %r38, ptr %p38
+  %p39 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 38
+  store i1 %r39, ptr %p39
+  %p40 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 39
+  store i1 %r40, ptr %p40
+  %p41 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 40
+  store i1 %r41, ptr %p41
+  %p42 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 41
+  store i1 %r42, ptr %p42
+  %p43 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 42
+  store i1 %r43, ptr %p43
+  %p44 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 43
+  store i1 %r44, ptr %p44
+  %p45 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 44
+  store i1 %r45, ptr %p45
+  %p46 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 45
+  store i1 %r46, ptr %p46
+  %p47 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 46
+  store i1 %r47, ptr %p47
+  %p48 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 47
+  store i1 %r48, ptr %p48
+  %p49 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 48
+  store i1 %r49, ptr %p49
+  %p50 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 49
+  store i1 %r50, ptr %p50
+  %p51 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 50
+  store i1 %r51, ptr %p51
+  %p52 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 51
+  store i1 %r52, ptr %p52
+  %p53 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 52
+  store i1 %r53, ptr %p53
+  %p54 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 53
+  store i1 %r54, ptr %p54
+  %p55 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 54
+  store i1 %r55, ptr %p55
+  %p56 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 55
+  store i1 %r56, ptr %p56
+  %p57 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 56
+  store i1 %r57, ptr %p57
+  %p58 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 57
+  store i1 %r58, ptr %p58
+  %p59 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 58
+  store i1 %r59, ptr %p59
+  %p60 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 59
+  store i1 %r60, ptr %p60
+  %p61 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 60
+  store i1 %r61, ptr %p61
+  %p62 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 61
+  store i1 %r62, ptr %p62
+  %p63 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 62
+  store i1 %r63, ptr %p63
+  %p64 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 63
+  store i1 %r64, ptr %p64
+  %p65 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 64
+  store i1 %r65, ptr %p65
+  %p66 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 65
+  store i1 %r66, ptr %p66
+  %p67 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 66
+  store i1 %r67, ptr %p67
+  %p68 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 67
+  store i1 %r68, ptr %p68
+  %p69 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 68
+  store i1 %r69, ptr %p69
+  %p70 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 69
+  store i1 %r70, ptr %p70
+  %p71 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 70
+  store i1 %r71, ptr %p71
+  %p72 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 71
+  store i1 %r72, ptr %p72
+  %p73 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 72
+  store i1 %r73, ptr %p73
+  %p74 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 73
+  store i1 %r74, ptr %p74
+  %p75 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 74
+  store i1 %r75, ptr %p75
+  %p76 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 75
+  store i1 %r76, ptr %p76
+  %p77 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 76
+  store i1 %r77, ptr %p77
+  %p78 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 77
+  store i1 %r78, ptr %p78
+  %p79 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 78
+  store i1 %r79, ptr %p79
+  %p80 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 79
+  store i1 %r80, ptr %p80
+  %p81 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 80
+  store i1 %r81, ptr %p81
+  %p82 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 81
+  store i1 %r82, ptr %p82
+  %p83 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 82
+  store i1 %r83, ptr %p83
+  %p84 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 83
+  store i1 %r84, ptr %p84
+  %p85 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 84
+  store i1 %r85, ptr %p85
+  %p86 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 85
+  store i1 %r86, ptr %p86
+  %p87 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 86
+  store i1 %r87, ptr %p87
+  %p88 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 87
+  store i1 %r88, ptr %p88
+  %p89 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 88
+  store i1 %r89, ptr %p89
+  %p90 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 89
+  store i1 %r90, ptr %p90
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll b/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll
index c8691c32710ad..7658362773218 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll
@@ -31,9 +31,12 @@
 
 %StructEvent = type { target("spirv.Event") }
 
+@G_r = global target("spirv.Event") poison
+
 define spir_kernel void @test_half(ptr addrspace(3) %_arg1, ptr addrspace(1) %_arg2) {
 entry:
   %r = tail call spir_func target("spirv.Event") @_Z22__spirv_GroupAsyncCopyjPU3AS3Dv2_DF16_PU3AS1KS_mm9ocl_event(i32 2, ptr addrspace(3) %_arg1, ptr addrspace(1) %_arg2, i64 16, i64 10, target("spirv.Event") zeroinitializer)
+  store target("spirv.Event") %r, ptr @G_r
   ret void
 }
 
@@ -42,7 +45,6 @@ declare dso_local spir_func target("spirv.Event") @_Z22__spirv_GroupAsyncCopyjPU
 ; CHECK: OpFunction
 ; CHECK: OpFunctionParameter
 ; CHECK: %[[#Src:]] = OpFunctionParameter
-; CHECK: OpVariable %[[#TyStructPtr]] Function
 ; CHECK: %[[#EventVar:]] = OpVariable %[[#TyEventPtr]] Function
 ; CHECK: %[[#Dest:]] = OpInBoundsPtrAccessChain
 ; CHECK: %[[#CopyRes:]] = OpGroupAsyncCopy %[[#TyEvent]] %[[#]] %[[#Dest]] %[[#Src]] %[[#]] %[[#]] %[[#ConstEvent]]
diff --git a/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll b/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
index 46668645f418b..9c8b4070d834d 100644
--- a/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
+++ b/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
@@ -68,6 +68,27 @@
 ; SPV-DAG: %[[#ones_64:]] = OpConstantComposite %[[#vec_64]] %[[#one_64]] %[[#one_64]]
 ; SPV-DAG: %[[#pointer:]] = OpTypePointer CrossWorkgroup %[[#float]]
 
+@G_s1 = global i8 0
+@G_s2 = global i16 0
+@G_s3 = global i32 0
+@G_s4 = global i64 0
+@G_s5 = global <2 x i8> zeroinitializer
+@G_s6 = global <2 x i16> zeroinitializer
+@G_s7 = global <2 x i32> zeroinitializer
+@G_s8 = global <2 x i64> zeroinitializer
+@G_z1 = global i8 0
+@G_z2 = global i16 0
+@G_z3 = global i32 0
+@G_z4 = global i64 0
+@G_z5 = global <2 x i8> zeroinitializer
+@G_z6 = global <2 x i16> zeroinitializer
+@G_z7 = global <2 x i32> zeroinitializer
+@G_z8 = global <2 x i64> zeroinitializer
+@G_ufp1 = global float 0.0
+@G_ufp2 = global <2 x float> zeroinitializer
+@G_sfp1 = global float 0.0
+@G_sfp2 = global <2 x float> zeroinitializer
+
 ; SPV-DAG: OpFunction
 ; SPV-DAG: %[[#A:]] = OpFunctionParameter %[[#pointer]]
 ; SPV-DAG: %[[#B:]] = OpFunctionParameter %[[#]]
@@ -87,47 +108,67 @@ entry:
 
 ; SPV-DAG: %[[#s1]] = OpSelect %[[#int_8]] %[[#i1s]] %[[#mone_8]] %[[#zero_8]]
   %s1 = sext i1 %i1s to i8
+  store i8 %s1, ptr @G_s1
 ; SPV-DAG: %[[#s2]] = OpSelect %[[#int_16]] %[[#i1s]] %[[#mone_16]] %[[#zero_16]]
   %s2 = sext i1 %i1s to i16
+  store i16 %s2, ptr @G_s2
 ; SPV-DAG: %[[#s3]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#mone_32]] %[[#zero_32]]
   %s3 = sext i1 %i1s to i32
+  store i32 %s3, ptr @G_s3
 ; SPV-DAG: %[[#s4]] = OpSelect %[[#int_64]] %[[#i1s]] %[[#mone_64]] %[[#zero_64]]
   %s4 = sext i1 %i1s to i64
+  store i64 %s4, ptr @G_s4
 ; SPV-DAG: %[[#s5]] = OpSelect %[[#vec_8]] %[[#i1v]] %[[#mones_8]] %[[#zeros_8]]
   %s5 = sext <2 x i1> %i1v to <2 x i8>
+  store <2 x i8> %s5, ptr @G_s5
 ; SPV-DAG: %[[#s6]] = OpSelect %[[#vec_16]] %[[#i1v]] %[[#mones_16]] %[[#zeros_16]]
   %s6 = sext <2 x i1> %i1v to <2 x i16>
+  store <2 x i16> %s6, ptr @G_s6
 ; SPV-DAG: %[[#s7]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#mones_32]] %[[#zeros_32]]
   %s7 = sext <2 x i1> %i1v to <2 x i32>
+  store <2 x i32> %s7, ptr @G_s7
 ; SPV-DAG: %[[#s8]] = OpSelect %[[#vec_64]] %[[#i1v]] %[[#mones_64]] %[[#zeros_64]]
   %s8 = sext <2 x i1> %i1v to <2 x i64>
+  store <2 x i64> %s8, ptr @G_s8
 ; SPV-DAG: %[[#z1]] = OpSelect %[[#int_8]] %[[#i1s]] %[[#one_8]] %[[#zero_8]]
   %z1 = zext i1 %i1s to i8
+  store i8 %z1, ptr @G_z1
 ; SPV-DAG: %[[#z2]] = OpSelect %[[#int_16]] %[[#i1s]] %[[#one_16]] %[[#zero_16]]
   %z2 = zext i1 %i1s to i16
+  store i16 %z2, ptr @G_z2
 ; SPV-DAG: %[[#z3]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
   %z3 = zext i1 %i1s to i32
+  store i32 %z3, ptr @G_z3
 ; SPV-DAG: %[[#z4]] = OpSelect %[[#int_64]] %[[#i1s]] %[[#one_64]] %[[#zero_64]]
   %z4 = zext i1 %i1s to i64
+  store i64 %z4, ptr @G_z4
 ; SPV-DAG: %[[#z5]] = OpSelect %[[#vec_8]] %[[#i1v]] %[[#ones_8]] %[[#zeros_8]]
   %z5 = zext <2 x i1> %i1v to <2 x i8>
+  store <2 x i8> %z5, ptr @G_z5
 ; SPV-DAG: %[[#z6]] = OpSelect %[[#vec_16]] %[[#i1v]] %[[#ones_16]] %[[#zeros_16]]
   %z6 = zext <2 x i1> %i1v to <2 x i16>
+  store <2 x i16> %z6, ptr @G_z6
 ; SPV-DAG: %[[#z7]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
   %z7 = zext <2 x i1> %i1v to <2 x i32>
+  store <2 x i32> %z7, ptr @G_z7
 ; SPV-DAG: %[[#z8]] = OpSelect %[[#vec_64]] %[[#i1v]] %[[#ones_64]] %[[#zeros_64]]
   %z8 = zext <2 x i1> %i1v to <2 x i64>
+  store <2 x i64> %z8, ptr @G_z8
 ; SPV-DAG: %[[#ufp1_res:]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
 ; SPV-DAG: %[[#ufp1]] = OpConvertUToF %[[#float]] %[[#ufp1_res]]
   %ufp1 = uitofp i1 %i1s to float
+  store float %ufp1, ptr @G_ufp1
 ; SPV-DAG: %[[#ufp2_res:]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
 ; SPV-DAG: %[[#ufp2]] = OpConvertUToF %[[#vec_float]] %[[#ufp2_res]]
   %ufp2 = uitofp <2 x i1> %i1v to <2 x float>
+  store <2 x float> %ufp2, ptr @G_ufp2
 ; SPV-DAG: %[[#sfp1_res:]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
 ; SPV-DAG: %[[#sfp1]] = OpConvertSToF %[[#float]] %[[#sfp1_res]]
   %sfp1 = sitofp i1 %i1s to float
+  store float %sfp1, ptr @G_sfp1
 ; SPV-DAG: %[[#sfp2_res:]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
 ; SPV-DAG: %[[#sfp2]] = OpConvertSToF %[[#vec_float]] %[[#sfp2_res]]
   %sfp2 = sitofp <2 x i1> %i1v to <2 x float>
+  store <2 x float> %sfp2, ptr @G_sfp2
   ret void
 }