diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 9f55330cb7ac6..2c3e0876b757d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -21,7 +21,9 @@
 #include "llvm/IR/InstIterator.h"
 #include "llvm/IR/InstVisitor.h"
 #include "llvm/IR/IntrinsicsSPIRV.h"
+#include "llvm/IR/PatternMatch.h"
 #include "llvm/IR/TypedPointerType.h"
+#include "llvm/Transforms/Utils/Local.h"
 
 #include <queue>
 #include <unordered_set>
@@ -187,6 +189,8 @@ class SPIRVEmitIntrinsics
 
   void applyDemangledPtrArgTypes(IRBuilder<> &B);
 
+  GetElementPtrInst *simplifyZeroLengthArrayGepInst(GetElementPtrInst *GEP);
+
   bool runOnFunction(Function &F);
   bool postprocessTypes(Module &M);
   bool processFunctionPointers(Module &M);
@@ -2561,6 +2565,30 @@ void SPIRVEmitIntrinsics::applyDemangledPtrArgTypes(IRBuilder<> &B) {
   }
 }
 
+GetElementPtrInst *
+SPIRVEmitIntrinsics::simplifyZeroLengthArrayGepInst(GetElementPtrInst *GEP) {
+  // getelementptr [0 x T], P, 0 (zero), I -> getelementptr T, P, I.
+  // If type is 0-length array and first index is 0 (zero), drop both the
+  // 0-length array type and the first index. This is a common pattern in the
+  // IR, e.g. when using a zero-length array as a placeholder for a flexible
+  // array such as unbound arrays.
+  assert(GEP && "GEP is null");
+  Type *SrcTy = GEP->getSourceElementType();
+  SmallVector<Value *, 8> Indices(GEP->indices());
+  ArrayType *ArrTy = dyn_cast<ArrayType>(SrcTy);
+  if (ArrTy && ArrTy->getNumElements() == 0 &&
+      PatternMatch::match(Indices[0], PatternMatch::m_Zero())) {
+    IRBuilder<> Builder(GEP);
+    Indices.erase(Indices.begin());
+    SrcTy = ArrTy->getElementType();
+    Value *NewGEP = Builder.CreateGEP(SrcTy, GEP->getPointerOperand(), Indices,
+                                      "", GEP->getNoWrapFlags());
+    assert(llvm::isa<GetElementPtrInst>(NewGEP) && "NewGEP should be a GEP");
+    return cast<GetElementPtrInst>(NewGEP);
+  }
+  return nullptr;
+}
+
 bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
   if (Func.isDeclaration())
     return false;
@@ -2578,14 +2606,30 @@ bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
   AggrConstTypes.clear();
   AggrStores.clear();
 
-  // fix GEP result types ahead of inference
+  // Fix GEP result types ahead of inference, and simplify if possible.
+  // Data structure for dead instructions that were simplified and replaced.
+  SmallPtrSet<Instruction *, 4> DeadInsts;
   for (auto &I : instructions(Func)) {
     auto *Ref = dyn_cast<GetElementPtrInst>(&I);
     if (!Ref || GR->findDeducedElementType(Ref))
       continue;
+
+    GetElementPtrInst *NewGEP = simplifyZeroLengthArrayGepInst(Ref);
+    if (NewGEP) {
+      Ref->replaceAllUsesWith(NewGEP);
+      if (isInstructionTriviallyDead(Ref))
+        DeadInsts.insert(Ref);
+
+      Ref = NewGEP;
+    }
     if (Type *GepTy = getGEPType(Ref))
       GR->addDeducedElementType(Ref, normalizeType(GepTy));
   }
+  // Remove dead instructions that were simplified and replaced.
+  for (auto *I : DeadInsts) {
+    assert(I->use_empty() && "Dead instruction should not have any uses left");
+    I->eraseFromParent();
+  }
 
   processParamTypesByFunHeader(CurrF, B);
 
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 83fccdc2bdba3..960eb2ef93a9e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -828,6 +828,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
          "Invalid array element type");
   SPIRVType *SpvTypeInt32 = getOrCreateSPIRVIntegerType(32, MIRBuilder);
   SPIRVType *ArrayType = nullptr;
+  const SPIRVSubtarget &ST =
+      cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
   if (NumElems != 0) {
     Register NumElementsVReg =
         buildConstantInt(NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
@@ -838,6 +840,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
           .addUse(NumElementsVReg);
     });
   } else {
+    if (!ST.isShader())
+      return nullptr;
     ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
       return MIRBuilder.buildInstr(SPIRV::OpTypeRuntimeArray)
           .addDef(createTypeVReg(MIRBuilder))
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index 0cd9d7882a52a..000e2d7105c8d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -744,8 +744,14 @@ void SPIRV::RequirementHandler::checkSatisfiable(
     IsSatisfiable = false;
   }
 
+  AvoidCapabilitiesSet AvoidCaps;
+  if (!ST.isShader())
+    AvoidCaps.S.insert(SPIRV::Capability::Shader);
+  else
+    AvoidCaps.S.insert(SPIRV::Capability::Kernel);
+
   for (auto Cap : MinimalCaps) {
-    if (AvailableCaps.contains(Cap))
+    if (AvailableCaps.contains(Cap) && !AvoidCaps.S.contains(Cap))
       continue;
     LLVM_DEBUG(dbgs() << "Capability not supported: "
                       << getSymbolicOperandMnemonic(
@@ -1865,6 +1871,10 @@ void addInstrRequirements(const MachineInstr &MI,
     Reqs.addCapability(SPIRV::Capability::TernaryBitwiseFunctionINTEL);
     break;
   }
+  case SPIRV::OpCopyMemorySized: {
+    Reqs.addCapability(SPIRV::Capability::Addresses);
+    // TODO: Add UntypedPointersKHR when implemented.
+  }
 
   default:
     break;
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
new file mode 100644
index 0000000000000..251b48f8bf629
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -0,0 +1,78 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
+
+; CHECK: OpCapability Kernel
+; CHECK-NOT: OpCapability Shader
+; CHECK-DAG: %[[#float16:]] = OpTypeFloat 16
+; CHECK-DAG: %[[#SyclHalfTy:]] = OpTypeStruct %[[#float16]]
+; CHECK-DAG: %[[#i16:]] = OpTypeInt 16
+; CHECK-DAG: %[[#i32:]] = OpTypeInt 32
+; CHECK-DAG: %[[#i64:]] = OpTypeInt 64
+; CHECK-DAG: %[[#ConstNull:]] = OpConstantNull %[[#i64]]
+; CHECK-DAG: %[[#ConstOne:]] = OpConstant %[[#i64]] 1
+; CHECK-DAG: %[[#ConstFive:]] = OpConstant %[[#i16]] 5
+; CHECK-DAG: %[[#SyclHalfTyPtr:]] = OpTypePointer Function %[[#SyclHalfTy]]
+; CHECK-DAG: %[[#i32Ptr:]] = OpTypePointer Function %[[#i32]]
+; CHECK-DAG: %[[#StorePtrTy:]] = OpTypePointer Function %[[#i16]]
+
+%"class.sycl::_V1::detail::half_impl::half" = type { half }
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+  %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 0, i64 0
+  store i16 5, ptr %0
+  ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo2(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+  %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 1, i64 1
+  store i16 5, ptr %0
+  ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo3(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+  %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 0, i64 0
+  store i16 5, ptr %0
+  ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo4(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+  %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 1, i64 1
+  store i16 5, ptr %0
+  ret void
+}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
index 40e2aff0d755a..7adb039464c4f 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
@@ -1,18 +1,11 @@
 ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
 
 ; CHECK-EXTENSION:      OpCapability BitInstructions
 ; CHECK-EXTENSION-NEXT: OpExtension "SPV_KHR_bit_instructions"
 ; CHECK-EXTENSION-NOT:  OpCabilitity Shader
-; CHECK-NO-EXTENSION:     OpCapability Shader
-; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions
-; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions"
-
-
 ; CHECK-EXTENSION: %[[#int:]] = OpTypeInt 32
 ; CHECK-EXTENSION: OpBitReverse %[[#int]]
-; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32
-; CHECK-NO-EXTENSION: OpBitReverse %[[#int]]
 
 define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr {
 entry:
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
index 65cccc83a3e02..3bd1bd633c258 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only.ll
@@ -1,12 +1,8 @@
 ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val %} 
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - -filetype=obj | spirv-val --target-env opencl2.2 %} 
 ;
 ; CHECK-EXTENSION: Capability BitInstructions
 ; CHECK-EXTENSION: Extension "SPV_KHR_bit_instructions"
-; CHECK-NO-EXTENSION-NOT: Capability BitInstructions 
-; CHECK-NO-EXTENSION-NOT: Extension "SPV_KHR_bit_instructions"
-; CHECK-NO-EXTENSION: Capability Shader 
 ;
 ; CHECK-EXTENSION: %[[#]] = OpFunction %[[#]] None %[[#]]
 ; CHECK-EXTENSION: %[[#reversebase:]] = OpFunctionParameter %[[#]]
@@ -15,24 +11,11 @@
 ; kernel void testBitReverse_SPIRVFriendly(long4 b, global long4 *res) {
 ;   *res = bit_reverse(b);
 ; }
-define spir_kernel void @testBitReverse_SPIRVFriendly(<4 x i64> %b, ptr addrspace(1) nocapture align 32 %res) #3 {
+define spir_kernel void @testBitReverse_SPIRVFriendly(<4 x i64> %b, ptr addrspace(1) %res) {
 entry:
   %call = call <4 x i64> @llvm.bitreverse.v4i64(<4 x i64> %b)
-  store <4 x i64> %call, ptr addrspace(1) %res, align 32
+  store <4 x i64> %call, ptr addrspace(1) %res
   ret void
 }
 
-declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>) #4
-
-
-attributes #3 = { nounwind }
-attributes #4 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-
-!llvm.module.flags = !{!0}
-!opencl.ocl.version = !{!1}
-!opencl.spir.version = !{!1}
-!llvm.ident = !{!2}
-
-!0 = !{i32 1, !"wchar_size", i32 4}
-!1 = !{i32 2, i32 0}
-!2 = !{!"clang version 20.0.0git (https://github.com/llvm/llvm-project.git cc61409d353a40f62d3a137f3c7436aa00df779d)"}
+declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>)
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll
new file mode 100644
index 0000000000000..61ef273a25734
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions/cl_khr_extended_bit_ops_spv-friendly_only_no_extension.ll
@@ -0,0 +1,16 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} 
+;
+; CHECK-NO-EXTENSION-NOT: Capability BitInstructions 
+; CHECK-NO-EXTENSION-NOT: Extension "SPV_KHR_bit_instructions"
+; CHECK-NO-EXTENSION: Capability Shader 
+
+define internal spir_func void @testBitReverse_SPIRVFriendly() #3 {
+entry:
+  %call = call <4 x i64> @llvm.bitreverse.v4i64(<4 x i64> <i64 1, i64 2, i64 3, i64 4>)
+  ret void
+}
+
+declare <4 x i64> @llvm.bitreverse.v4i64(<4 x i64>)
+
+attributes #3 = { nounwind "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
new file mode 100644
index 0000000000000..452df0a64c063
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
@@ -0,0 +1,23 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
+
+
+; CHECK-NO-EXTENSION:     OpCapability Shader
+; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions
+; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions"
+; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32
+; CHECK-NO-EXTENSION: OpBitReverse %[[#int]]
+
+define hidden spir_func void @testBitRev(i32 %a, i32 %b, i32 %c, ptr %res) local_unnamed_addr {
+entry:
+  %call = tail call i32 @llvm.bitreverse.i32(i32 %b)
+  store i32 %call, ptr %res, align 4
+  ret void
+}
+
+define void @main() #1 {
+  ret void
+}
+
+declare i32 @llvm.bitreverse.i32(i32)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
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 77b8c5118e59b..438fff6e94f89 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
@@ -1,8 +1,8 @@
 ;; Check that llvm.bitreverse.* intrinsics are lowered for
 ;; 2/4-bit scalar and vector types.
 
-; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers %s -o - | FileCheck %s
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers,+SPV_KHR_bit_instructions %s -o - | FileCheck %s
+; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_INTEL_arbitrary_precision_integers,+SPV_KHR_bit_instructions %s -o - -filetype=obj | spirv-val %}
 
 ; CHECK: OpCapability ArbitraryPrecisionIntegersINTEL
 ; CHECK: OpExtension "SPV_INTEL_arbitrary_precision_integers"
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
index 9d07b63b49a52..483d7077a66f9 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
@@ -18,7 +18,6 @@
 ; CL:      %[[#FooVar:]] = OpVariable
 ; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
 ; CL-NEXT: OpLifetimeStart %[[#Casted1]], 72
-; CL-NEXT: OpCopyMemorySized
 ; CL-NEXT: OpBitcast
 ; CL-NEXT: OpInBoundsPtrAccessChain
 ; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
@@ -26,13 +25,11 @@
 
 ; VK:      OpFunction
 ; VK:      %[[#FooVar:]] = OpVariable
-; VK-NEXT: OpCopyMemorySized
 ; VK-NEXT: OpInBoundsAccessChain
 ; 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(i64 72, ptr nonnull %RoundedRangeKernel)
-  call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
   %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
   call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %RoundedRangeKernel)
   ret void
@@ -41,20 +38,17 @@ define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange)
 ; CL: OpFunction
 ; CL: %[[#BarVar:]] = OpVariable
 ; CL-NEXT: OpLifetimeStart %[[#BarVar]], 0
-; CL-NEXT: OpCopyMemorySized
 ; CL-NEXT: OpBitcast
 ; CL-NEXT: OpInBoundsPtrAccessChain
 ; CL-NEXT: OpLifetimeStop %[[#BarVar]], 0
 
 ; VK:      OpFunction
 ; VK:      %[[#BarVar:]] = OpVariable
-; VK-NEXT: OpCopyMemorySized
 ; VK-NEXT: OpInBoundsAccessChain
 ; 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(i64 -1, ptr nonnull %RoundedRangeKernel)
-  call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
   %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
   call void @llvm.lifetime.end.p0(i64 -1, ptr nonnull %RoundedRangeKernel)
   ret void
@@ -63,20 +57,17 @@ define spir_func void @bar(ptr noundef byval(%tprange) align 8 %_arg_UserRange)
 ; CL: OpFunction
 ; CL: %[[#TestVar:]] = OpVariable
 ; CL-NEXT: OpLifetimeStart %[[#TestVar]], 1
-; CL-NEXT: OpCopyMemorySized
 ; CL-NEXT: OpInBoundsPtrAccessChain
 ; CL-NEXT: OpLifetimeStop %[[#TestVar]], 1
 
 ; VK:      OpFunction
 ; VK:      %[[#Test:]] = OpVariable
-; VK-NEXT: OpCopyMemorySized
 ; VK-NEXT: OpInBoundsAccessChain
 ; VK-NEXT: OpReturn
 define spir_func void @test(ptr noundef align 8 %_arg) {
   %var = alloca i8, align 8
   call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %var)
-  call void @llvm.memcpy.p0.p0.i64(ptr align 8 %var, ptr align 8 %_arg, i64 1, i1 false)
-  %KernelFunc = getelementptr inbounds i8, ptr %var, i64 0
+  %KernelFunc = getelementptr inbounds i8, ptr %var, i64 1
   call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %var)
   ret void
 }
diff --git a/llvm/test/CodeGen/SPIRV/memory_model_md.ll b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
index e52343cbbb7e4..684a163397ca8 100644
--- a/llvm/test/CodeGen/SPIRV/memory_model_md.ll
+++ b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
@@ -1,6 +1,7 @@
 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
 
-; SPV: OpMemoryModel Physical32 Simple
+; SPV: OpMemoryModel Physical32 OpenCL
 define dso_local dllexport void @k_no_fc(i32 %ibuf, i32 %obuf) local_unnamed_addr {
 entry:
   ret void
@@ -8,4 +9,4 @@ entry:
 
 !spirv.MemoryModel = !{!0}
 
-!0 = !{i32 1, i32 0}
+!0 = !{i32 1, i32 2}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
index 544c657da8488..19451d23c6830 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
@@ -1,5 +1,5 @@
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
 
 @PrivInternal = internal addrspace(10) global i32 456
 ; CHECK-DAG:  %[[#type:]] = OpTypeInt 32 0
@@ -7,7 +7,7 @@
 ; CHECK-DAG: %[[#value:]] = OpConstant %[[#type]] 456
 ; CHECK-DAG:   %[[#var:]] = OpVariable %[[#ptrty]] Private %[[#value]]
 
-define spir_kernel void @Foo() {
+define hidden spir_func void @Foo() {
   %p = addrspacecast ptr addrspace(10) @PrivInternal to ptr
   %v = load i32, ptr %p, align 4
   ret void
@@ -15,3 +15,9 @@ define spir_kernel void @Foo() {
 ; CHECK-NEXT: OpLoad %[[#type]] %[[#var]] Aligned 4
 ; CHECK-Next: OpReturn
 }
+
+define void @main() #1 {
+  ret void
+}
+
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll
new file mode 100644
index 0000000000000..51db12046e9de
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class-private.ll
@@ -0,0 +1,21 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3%}
+
+; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0
+
+; CHECK-DAG: %[[#VAL:]] = OpConstant %[[#U32]] 456
+; CHECK-DAG: %[[#VTYPE:]] = OpTypePointer Private %[[#U32]]
+; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]]
+@PrivInternal = internal addrspace(10) global i32 456
+
+define hidden spir_func void @Foo() {
+  %tmp = load i32, ptr addrspace(10) @PrivInternal
+  ret void
+}
+
+define void @main() #1 {
+  ret void
+}
+
+declare void @llvm.memcpy.p1.p2.i64(ptr addrspace(1) noalias nocapture writeonly, ptr addrspace(2) noalias nocapture readonly, i64, i1 immarg)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
index a1ded0569d67e..6914f4faebdcd 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
@@ -1,5 +1,5 @@
 ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
 
 ; CHECK-DAG: %[[#U8:]] = OpTypeInt 8 0
 ; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0
@@ -15,12 +15,7 @@
 ; CHECK-DAG: %[[#INIT:]] = OpVariable %[[#VTYPE]] UniformConstant %[[#VAL]]
 @Init = private addrspace(2) constant i32 123
 
-; CHECK-DAG: %[[#VAL:]] = OpConstant %[[#U32]] 456
-; CHECK-DAG: %[[#VTYPE:]] = OpTypePointer Private %[[#U32]]
-; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]]
-@PrivInternal = internal addrspace(10) global i32 456
-
-define spir_kernel void @Foo() {
+define internal spir_func void @Foo() {
   ; CHECK: %[[#]] = OpLoad %[[#]] %[[#PTR]] Aligned 8
   %l = load ptr addrspace(1), ptr addrspace(1) @Ptr, align 8
   ; CHECK: OpCopyMemorySized %[[#]] %[[#INIT]] %[[#]] Aligned 4
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
index f396b5a01ae91..838c55172d3fc 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
@@ -1,13 +1,19 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
 
 ; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32
 ; CHECK-SPIRV: OpBitReverse %[[#int]]
 
-define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr {
+define hidden spir_func void @testBitRev(i32 %a, i32 %b, i32 %c, ptr %res) {
 entry:
   %call = tail call i32 @llvm.bitreverse.i32(i32 %b)
-  store i32 %call, i32 addrspace(1)* %res, align 4
+  store i32 %call, ptr %res, align 4
+  ret void
+}
+
+define void @main() #1 {
   ret void
 }
 
 declare i32 @llvm.bitreverse.i32(i32)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
deleted file mode 100644
index 8f04929fdd587..0000000000000
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
+++ /dev/null
@@ -1,14 +0,0 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-
-; CHECK-SPIRV: %[[#short:]] = OpTypeInt 16
-; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2
-; CHECK-SPIRV: OpBitReverse %[[#short2]]
-
-define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> addrspace(1)* nocapture %res) local_unnamed_addr {
-entry:
-  %call = tail call <2 x i16> @llvm.bitreverse.v2i16(<2 x i16> %b)
-  store <2 x i16> %call, <2 x i16> addrspace(1)* %res, align 4
-  ret void
-}
-
-declare <2 x i16> @llvm.bitreverse.v2i16(<2 x i16>)
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll
new file mode 100644
index 0000000000000..3e2ed8b373e4d
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i32.ll
@@ -0,0 +1,20 @@
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val --target-env vulkan1.3 %}
+
+; CHECK-SPIRV: %[[#short:]] = OpTypeInt 32
+; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2
+; CHECK-SPIRV: OpBitReverse %[[#short2]]
+
+define hidden spir_func void @testBitRev(<2 x i32> %a, <2 x i32> %b, <2 x i32> %c, ptr %res) {
+entry:
+  %call = tail call <2 x i32> @llvm.bitreverse.v2i32(<2 x i32> %b)
+  store <2 x i32> %call, ptr %res, align 4
+  ret void
+}
+
+define void @main() #1 {
+  ret void
+}
+
+declare <2 x i32> @llvm.bitreverse.v2i32(<2 x i32>)
+attributes #1 = { "hlsl.numthreads"="8,1,1" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/zero-length-array.ll b/llvm/test/CodeGen/SPIRV/zero-length-array.ll
index 668bf2018dec7..666176c87adb6 100644
--- a/llvm/test/CodeGen/SPIRV/zero-length-array.ll
+++ b/llvm/test/CodeGen/SPIRV/zero-length-array.ll
@@ -1,9 +1,8 @@
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan-compute %s -o - -filetype=obj | spirv-val %}
 
 ; CHECK: %[[#type:]] = OpTypeInt 32 0
-; CHECK: %[[#ext:]] = OpTypeRuntimeArray %[[#type]]
-; CHECK: %[[#]] = OpTypePointer Function %[[#ext]]
+; CHECK: %[[#ext:]] = OpConstant %[[#type]] 0
 
 define spir_func void @_Z3foov() {
 entry: