[SPIRV] Do not use OpTypeRuntimeArray in Kernel env.

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);
 

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))

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;

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
+}

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:

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>)

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" }

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" }

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"

llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll

@@ -18,21 +18,18 @@
 ; 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]]
 ; CL-NEXT: OpLifetimeStop %[[#Casted2]], 72
 
 ; 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
 }

llvm/test/CodeGen/SPIRV/memory_model_md.ll

@@ -1,11 +1,12 @@
 ; 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
 }
 
 !spirv.MemoryModel = !{!0}
 
-!0 = !{i32 1, i32 0}
+!0 = !{i32 1, i32 2}