diff --git a/lib/SPIRV/SPIRVInternal.h b/lib/SPIRV/SPIRVInternal.h index 062609ff57..b4439649fe 100644 --- a/lib/SPIRV/SPIRVInternal.h +++ b/lib/SPIRV/SPIRVInternal.h @@ -239,7 +239,6 @@ inline void SPIRVMap::init() { add(Attribute::ByVal, FunctionParameterAttributeByVal); add(Attribute::StructRet, FunctionParameterAttributeSret); add(Attribute::NoAlias, FunctionParameterAttributeNoAlias); - add(Attribute::NoCapture, FunctionParameterAttributeNoCapture); add(Attribute::ReadOnly, FunctionParameterAttributeNoWrite); add(Attribute::ReadNone, FunctionParameterAttributeNoReadWrite); } diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp index d5cc452d63..e87fa296c4 100644 --- a/lib/SPIRV/SPIRVReader.cpp +++ b/lib/SPIRV/SPIRVReader.cpp @@ -3223,6 +3223,11 @@ void SPIRVToLLVM::transFunctionAttrs(SPIRVFunction *BF, Function *F) { // OpenCL metadata if (Kind == FunctionParameterAttributeRuntimeAlignedINTEL) return; + if (Kind == FunctionParameterAttributeNoCapture) { + I->addAttr(Attribute::getWithCaptureInfo(F->getContext(), + CaptureInfo::none())); + return; + } Attribute::AttrKind LLVMKind = SPIRSPIRVFuncParamAttrMap::rmap(Kind); if (IllegalAttrs.contains(LLVMKind)) return; diff --git a/lib/SPIRV/SPIRVRegularizeLLVM.cpp b/lib/SPIRV/SPIRVRegularizeLLVM.cpp index b81ca006e6..38bb34075b 100644 --- a/lib/SPIRV/SPIRVRegularizeLLVM.cpp +++ b/lib/SPIRV/SPIRVRegularizeLLVM.cpp @@ -101,7 +101,7 @@ void SPIRVRegularizeLLVMBase::lowerIntrinsicToFunction( Intrinsic->setCalledFunction(F); return; } - // TODO copy arguments attributes: nocapture writeonly. + // TODO copy arguments attributes: captures(none) writeonly. FunctionCallee FC = M->getOrInsertFunction(FuncName, Intrinsic->getFunctionType()); auto IntrinsicID = Intrinsic->getIntrinsicID(); diff --git a/test/DebugInfo/Generic/2010-10-01-crash.ll b/test/DebugInfo/Generic/2010-10-01-crash.ll index 7b29d5238e..d7250f7551 100644 --- a/test/DebugInfo/Generic/2010-10-01-crash.ll +++ b/test/DebugInfo/Generic/2010-10-01-crash.ll @@ -15,7 +15,7 @@ entry: declare void @llvm.dbg.declare(metadata, metadata, metadata) nounwind readnone -declare void @llvm.memcpy.p0.p0.i32(ptr nocapture, ptr nocapture, i32, i1) nounwind +declare void @llvm.memcpy.p0.p0.i32(ptr captures(none), ptr captures(none), i32, i1) nounwind !llvm.dbg.cu = !{!2} diff --git a/test/DebugInfo/Generic/c-and-cpp-mixed.ll b/test/DebugInfo/Generic/c-and-cpp-mixed.ll index c90751f30a..078873c39a 100644 --- a/test/DebugInfo/Generic/c-and-cpp-mixed.ll +++ b/test/DebugInfo/Generic/c-and-cpp-mixed.ll @@ -18,16 +18,16 @@ ; CHECK-LLVM: define spir_func void @foo() #0 !dbg ![[#Func1:]] { ; CHECK-LLVM: entry: -; CHECK-LLVM: %puts = call spir_func i32 @puts(ptr addrspace(1) nocapture @str) #0, !dbg ![[#Puts1Loc:]] +; CHECK-LLVM: %puts = call spir_func i32 @puts(ptr addrspace(1) captures(none) @str) #0, !dbg ![[#Puts1Loc:]] ; CHECK-LLVM: ret void, !dbg ![[#Ret1:]] ; CHECK-LLVM: } -; CHECK-LLVM: define spir_func i32 @main(i32 %argc, ptr nocapture %argv) #0 !dbg ![[#Func2:]] { +; CHECK-LLVM: define spir_func i32 @main(i32 %argc, ptr captures(none) %argv) #0 !dbg ![[#Func2:]] { ; CHECK-LLVM: entry: ; CHECK-LLVM: #dbg_value(i32 %argc, ![[#Fun2Param1:]], !DIExpression(), ![[#Fun2Param1Loc:]]) ; CHECK-LLVM: #dbg_value(ptr %argv, ![[#Fun2Param2:]], !DIExpression(DW_OP_deref, DW_OP_deref), ![[#Fun2Param2Loc:]]) ; CHECK-LLVM: %0 = bitcast ptr addrspace(1) @str1 to ptr addrspace(1), !dbg ![[#Puts2Loc:]] -; CHECK-LLVM: %puts = call spir_func i32 @puts(ptr addrspace(1) nocapture %0) #0, !dbg ![[#Puts2Loc]] +; CHECK-LLVM: %puts = call spir_func i32 @puts(ptr addrspace(1) captures(none) %0) #0, !dbg ![[#Puts2Loc]] ; CHECK-LLVM: call spir_func void @foo() #0, !dbg ![[#CallFoo:]] ; CHECK-LLVM: ret i32 0, !dbg ![[#Ret2:]] ; CHECK-LLVM: } @@ -74,9 +74,9 @@ entry: ret void, !dbg !25 } -declare i32 @puts(ptr addrspace(1) nocapture) nounwind +declare i32 @puts(ptr addrspace(1) captures(none)) nounwind -define i32 @main(i32 %argc, ptr nocapture %argv) nounwind !dbg !12 { +define i32 @main(i32 %argc, ptr captures(none) %argv) nounwind !dbg !12 { entry: tail call void @llvm.dbg.value(metadata i32 %argc, metadata !21, metadata !DIExpression()), !dbg !26 ; Avoid talking about the pointer size in debug info because that's target dependent diff --git a/test/DebugInfo/Generic/debug-info-eis-option.ll b/test/DebugInfo/Generic/debug-info-eis-option.ll index 20c25d679b..2338792c89 100644 --- a/test/DebugInfo/Generic/debug-info-eis-option.ll +++ b/test/DebugInfo/Generic/debug-info-eis-option.ll @@ -15,7 +15,7 @@ target triple = "spir64-unknown-unknown" source_filename = "linear-dbg-value.ll" ; Function Attrs: nounwind readonly uwtable -define i32 @foo(ptr nocapture readonly %a, i32 %N) local_unnamed_addr #0 !dbg !6 { +define i32 @foo(ptr captures(none) readonly %a, i32 %N) local_unnamed_addr #0 !dbg !6 { entry: %cmp6 = icmp sgt i32 %N, 0, !dbg !11 br i1 %cmp6, label %for.body.preheader, label %for.cond.cleanup, !dbg !15 diff --git a/test/DebugInfo/Generic/linear-dbg-value.ll b/test/DebugInfo/Generic/linear-dbg-value.ll index b4bd79b374..9b3285455a 100644 --- a/test/DebugInfo/Generic/linear-dbg-value.ll +++ b/test/DebugInfo/Generic/linear-dbg-value.ll @@ -15,7 +15,7 @@ target triple = "spir64-unknown-unknown" source_filename = "linear-dbg-value.ll" ; Function Attrs: nounwind readonly uwtable -define i32 @foo(ptr nocapture readonly %a, i32 %N) local_unnamed_addr #0 !dbg !6 { +define i32 @foo(ptr captures(none) readonly %a, i32 %N) local_unnamed_addr #0 !dbg !6 { entry: %cmp6 = icmp sgt i32 %N, 0, !dbg !11 br i1 %cmp6, label %for.body.preheader, label %for.cond.cleanup, !dbg !15 diff --git a/test/DebugInfo/Generic/two-cus-from-same-file.ll b/test/DebugInfo/Generic/two-cus-from-same-file.ll index 9ddd9f8133..514b71a9bb 100644 --- a/test/DebugInfo/Generic/two-cus-from-same-file.ll +++ b/test/DebugInfo/Generic/two-cus-from-same-file.ll @@ -26,9 +26,9 @@ entry: ret void, !dbg !25 } -declare i32 @puts(ptr addrspace(1) nocapture) nounwind +declare i32 @puts(ptr addrspace(1) captures(none)) nounwind -define i32 @main(i32 %argc, ptr nocapture %argv) nounwind !dbg !12 { +define i32 @main(i32 %argc, ptr captures(none) %argv) nounwind !dbg !12 { entry: tail call void @llvm.dbg.value(metadata i32 %argc, metadata !21, metadata !DIExpression()), !dbg !26 ; Avoid talking about the pointer size in debug info because that's target dependent diff --git a/test/DebugInfo/NonSemantic/Shader200/FortranComplex.ll b/test/DebugInfo/NonSemantic/Shader200/FortranComplex.ll index cb4a661d3f..4513c1978c 100644 --- a/test/DebugInfo/NonSemantic/Shader200/FortranComplex.ll +++ b/test/DebugInfo/NonSemantic/Shader200/FortranComplex.ll @@ -73,10 +73,10 @@ alloca_0: ret void, !dbg !16 } -declare !llfort.intrin_id !17 !llfort.type_idx !18 i32 @for_set_fpe_(ptr addrspace(1) nocapture readonly) local_unnamed_addr +declare !llfort.intrin_id !17 !llfort.type_idx !18 i32 @for_set_fpe_(ptr addrspace(1) captures(none) readonly) local_unnamed_addr ; Function Attrs: nofree -declare !llfort.intrin_id !19 !llfort.type_idx !20 i32 @for_set_reentrancy(ptr addrspace(1) nocapture readonly) local_unnamed_addr +declare !llfort.intrin_id !19 !llfort.type_idx !20 i32 @for_set_reentrancy(ptr addrspace(1) captures(none) readonly) local_unnamed_addr ; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none) declare void @llvm.dbg.value(metadata, metadata, metadata) diff --git a/test/DebugInfo/RecursiveDebugInfo.ll b/test/DebugInfo/RecursiveDebugInfo.ll index fab380b628..82fa84074b 100644 --- a/test/DebugInfo/RecursiveDebugInfo.ll +++ b/test/DebugInfo/RecursiveDebugInfo.ll @@ -74,7 +74,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nounwind readnone speculatable willreturn declare void @llvm.dbg.declare(metadata, metadata, metadata) #2 @@ -99,7 +99,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nounwind define linkonce_odr dso_local spir_func void @_ZN8iteratorI6vectorEC2Ev(ptr addrspace(4) %this) unnamed_addr #4 comdat align 2 !dbg !62 { diff --git a/test/DebugInfo/X86/dbg-value-location.ll b/test/DebugInfo/X86/dbg-value-location.ll index 44dc48c6f3..bab3ddcef9 100644 --- a/test/DebugInfo/X86/dbg-value-location.ll +++ b/test/DebugInfo/X86/dbg-value-location.ll @@ -56,7 +56,7 @@ if.else: ; preds = %entry ret i32 0 } -declare hidden fastcc i32 @bar(i32, ptr nocapture) nounwind optsize ssp +declare hidden fastcc i32 @bar(i32, ptr captures(none)) nounwind optsize ssp declare hidden fastcc i32 @bar2(i32) nounwind optsize ssp declare hidden fastcc i32 @bar3(i32) nounwind optsize ssp declare void @llvm.dbg.value(metadata, metadata, metadata) nounwind readnone diff --git a/test/DebugInfo/X86/dbg-value-range.ll b/test/DebugInfo/X86/dbg-value-range.ll index a7ac9bbaee..218ecf3e5a 100644 --- a/test/DebugInfo/X86/dbg-value-range.ll +++ b/test/DebugInfo/X86/dbg-value-range.ll @@ -16,7 +16,7 @@ target triple = "spir64-unknown-unknown" %struct.a = type { i32 } -define i32 @bar(ptr nocapture %b) nounwind ssp !dbg !0 { +define i32 @bar(ptr captures(none) %b) nounwind ssp !dbg !0 { entry: tail call void @llvm.dbg.value(metadata ptr %b, metadata !6, metadata !DIExpression()), !dbg !13 %tmp2 = load i32, ptr %b, align 4, !dbg !14 diff --git a/test/DebugInfo/bitfields_packed.ll b/test/DebugInfo/bitfields_packed.ll index 60730341cc..0decbdc2e6 100644 --- a/test/DebugInfo/bitfields_packed.ll +++ b/test/DebugInfo/bitfields_packed.ll @@ -79,7 +79,7 @@ define dso_local spir_func noundef i32 @_Z3fooii(i32 noundef %0, i32 noundef %1) declare void @llvm.dbg.declare(metadata, metadata, metadata) #1 ; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite) -declare void @llvm.memcpy.p4.p1.i64(ptr addrspace(4) noalias nocapture writeonly, ptr addrspace(1) noalias nocapture readonly, i64, i1 immarg) #2 +declare void @llvm.memcpy.p4.p1.i64(ptr addrspace(4) noalias captures(none) writeonly, ptr addrspace(1) noalias captures(none) readonly, i64, i1 immarg) #2 attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "approx-func-fp-math"="true" "frame-pointer"="all" "min-legal-vector-width"="0" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="bitfields-packed.cpp" "unsafe-fp-math"="true" } attributes #1 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) } diff --git a/test/SpecConstants/long-spec-const-composite.ll b/test/SpecConstants/long-spec-const-composite.ll index 2847e517ef..ab224d7c7f 100644 --- a/test/SpecConstants/long-spec-const-composite.ll +++ b/test/SpecConstants/long-spec-const-composite.ll @@ -65600,10 +65600,10 @@ entry: } ; Function Attrs: argmemonly nofree nosync nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nofree nosync nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 declare float @_Z20__spirv_SpecConstantif(i32, float) diff --git a/test/capbility-kernel.ll b/test/capbility-kernel.ll index d7010f2aef..266d7d8680 100644 --- a/test/capbility-kernel.ll +++ b/test/capbility-kernel.ll @@ -12,7 +12,7 @@ target triple = "spir64-unknown-unknown" ; CHECK-DAG: {{[0-9]*}} Capability Linkage ; Function Attrs: nounwind -define spir_func void @func_export(ptr addrspace(1) nocapture %a) #0 { +define spir_func void @func_export(ptr addrspace(1) captures(none) %a) #0 { entry: ; CHECK-DAG: {{[0-9]*}} Capability Int64 %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #3 diff --git a/test/constexpr_phi.ll b/test/constexpr_phi.ll index f20a64aa06..c4d0d498d8 100644 --- a/test/constexpr_phi.ll +++ b/test/constexpr_phi.ll @@ -41,7 +41,7 @@ define dso_local i32 @_Z2f2i(i32 %0) { ret i32 %2 } -define dso_local i64 @_Z3runiiPi(i32 %0, i32 %1, ptr nocapture %2) local_unnamed_addr { +define dso_local i64 @_Z3runiiPi(i32 %0, i32 %1, ptr captures(none) %2) local_unnamed_addr { %4 = icmp slt i32 %0, 10 br i1 %4, label %5, label %7 diff --git a/test/custom_class.ll b/test/custom_class.ll index 9265da6272..ae3d267999 100644 --- a/test/custom_class.ll +++ b/test/custom_class.ll @@ -49,7 +49,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlvE_clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -63,7 +63,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind define dso_local spir_func void @_ZN11CustomClassC2Ei(ptr addrspace(4) %this, i32 %value) unnamed_addr #3 align 2 { diff --git a/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/non-constant-printf.ll b/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/non-constant-printf.ll index 38b4f6d10a..38bfa22123 100644 --- a/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/non-constant-printf.ll +++ b/test/extensions/EXT/SPV_EXT_relaxed_printf_string_address_space/non-constant-printf.ll @@ -82,7 +82,7 @@ declare spir_func i32 @_Z18__spirv_ocl_printfPU3AS3c(ptr addrspace(3)) #0 declare spir_func i32 @_Z18__spirv_ocl_printfPU3AS4c(ptr addrspace(4)) #0 ; Function Attrs: nounwind -declare void @llvm.memcpy.p0.p2.i64(ptr nocapture, ptr addrspace(2) nocapture readonly, i64, i1) #0 +declare void @llvm.memcpy.p0.p2.i64(ptr captures(none), ptr addrspace(2) captures(none) readonly, i64, i1) #0 attributes #0 = { nounwind } diff --git a/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll b/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll index 041874085b..3825ec2328 100644 --- a/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll +++ b/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_fixed_point/capability-arbitrary-precision-fixed-point-numbers.ll @@ -240,7 +240,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -264,7 +264,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind define linkonce_odr dso_local spir_func void @_Z4sqrtILi13ELi5ELb0ELi2ELi2EEvv() #3 comdat { diff --git a/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll b/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll index 4391da3ec9..344aa26c57 100644 --- a/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll +++ b/test/extensions/INTEL/SPV_INTEL_arbitrary_precision_floating_point/capability-arbitrary-precision-floating-point.ll @@ -493,7 +493,7 @@ define dso_local spir_kernel void @_ZTSZ4mainE15kernel_function() #0 !kernel_arg } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %0) #2 align 2 { @@ -548,7 +548,7 @@ define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %0) #2 al } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind define linkonce_odr dso_local spir_func void @_Z13ap_float_castILi11ELi28ELi9ELi30EEvv() #3 { diff --git a/test/extensions/INTEL/SPV_INTEL_complex_float_mul_div/complex-operations.ll b/test/extensions/INTEL/SPV_INTEL_complex_float_mul_div/complex-operations.ll index b8490f2861..5daef2cc48 100644 --- a/test/extensions/INTEL/SPV_INTEL_complex_float_mul_div/complex-operations.ll +++ b/test/extensions/INTEL/SPV_INTEL_complex_float_mul_div/complex-operations.ll @@ -30,7 +30,7 @@ target triple = "spir-unknown-unknown" %structtype = type { float, float } ; Function Attrs: nounwind -define spir_func void @_Z19cmul_kernel_complexPSt7complexIfES1_S1_(ptr noalias nocapture readonly %a, ptr noalias nocapture readonly %b, ptr noalias nocapture %c) #0 { +define spir_func void @_Z19cmul_kernel_complexPSt7complexIfES1_S1_(ptr noalias captures(none) readonly %a, ptr noalias captures(none) readonly %b, ptr noalias captures(none) %c) #0 { entry: %0 = load <2 x float>, ptr %a, align 4 %1 = load <2 x float>, ptr %b, align 4 diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_buffer_location/FPGABufferLocation.ll b/test/extensions/INTEL/SPV_INTEL_fpga_buffer_location/FPGABufferLocation.ll index 3429cd0bbd..78ce87dc8b 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_buffer_location/FPGABufferLocation.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_buffer_location/FPGABufferLocation.ll @@ -119,7 +119,7 @@ entry: declare ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4), ptr addrspace(1), ptr addrspace(1), i32, ptr addrspace(1)) #1 ; Function Attrs: argmemonly nofree nounwind willreturn -declare void @llvm.memcpy.p4.p0(ptr addrspace(4) noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #2 +declare void @llvm.memcpy.p4.p0(ptr addrspace(4) noalias captures(none) writeonly, ptr noalias captures(none) readonly, i64, i1 immarg) #2 !opencl.enable.FP_CONTRACT = !{} !opencl.ocl.version = !{!0} diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll index f6f53a5566..0f822fb41c 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttr.ll @@ -113,7 +113,7 @@ define dso_local spir_kernel void @_ZTSZ4mainE15kernel_function() #0 !kernel_arg } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* captures(none)) #1 ; Function Attrs: inlinehint nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class._ZTSZ4mainE3$_0.anon" addrspace(4)* %0) #2 align 2 { @@ -126,7 +126,7 @@ define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class._ZTSZ4mainE3$_0.an } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* captures(none)) #1 ; CHECK-SPIRV: Function ; Function Attrs: nounwind diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll index 71322fb469..9e8883e84f 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll @@ -136,7 +136,7 @@ arrayinit.end8: ; preds = %arrayinit.body3 } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* captures(none)) #1 ; CHECK-SPIRV: Function {{.*}} [[TYPE_EMB_FUNC]] ; CHECK-LLVM: define internal spir_func void {{.*}}(ptr addrspace(4) %this) @@ -266,7 +266,7 @@ for.end20: ; preds = %for.cond.cleanup } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* captures(none)) #1 ; Function Attrs: convergent norecurse uwtable define dso_local spir_kernel void @_ZTSZ4mainE18VaryingSafelenOnPointersTest(i32 addrspace(1)* %_arg_, i32 addrspace(1)* %_arg_1) #0 !kernel_arg_buffer_location !4 { diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopAttr.ll b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopAttr.ll index 4721ddd575..38d4dec795 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopAttr.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopAttr.ll @@ -248,9 +248,9 @@ for.end11: ; preds = %for.cond.cleanup4 ret void } -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopMergeInst.ll b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopMergeInst.ll index 60935b8ef5..bf0f896f47 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopMergeInst.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGALoopMergeInst.ll @@ -137,7 +137,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -150,7 +150,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nounwind define dso_local spir_func void @_Z3foov() #3 { diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/intel_multiple_fpga_loop_attrs.ll b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/intel_multiple_fpga_loop_attrs.ll index 4e86c6d499..d9322834a6 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/intel_multiple_fpga_loop_attrs.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/intel_multiple_fpga_loop_attrs.ll @@ -76,10 +76,10 @@ for.end: ; preds = %for.cond.cleanup } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind define i32 @main() #2 { diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll b/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll index 876b7e9c11..7405618de6 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll @@ -96,7 +96,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1 +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* captures(none)) #1 ; Function Attrs: inlinehint norecurse nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class._ZTSZ4mainE3$_0.anon" addrspace(4)* %this) #2 align 2 { @@ -126,7 +126,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1 +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* captures(none)) #1 ; CHECK-LLVM: define spir_func void @{{.*}}foo ; Function Attrs: norecurse nounwind diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll b/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll index 6cfd5c5be7..b6d67a7e03 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStaticVar.ll @@ -87,7 +87,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr %this) #2 align 2 { @@ -103,7 +103,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; CHECK-LLVM: void @_Z13numbanks_stati(i32 %a) ; Function Attrs: norecurse nounwind diff --git a/test/extensions/INTEL/SPV_INTEL_fpga_reg/IntelFPGAReg.ll b/test/extensions/INTEL/SPV_INTEL_fpga_reg/IntelFPGAReg.ll index 367608b876..2d623b5d99 100644 --- a/test/extensions/INTEL/SPV_INTEL_fpga_reg/IntelFPGAReg.ll +++ b/test/extensions/INTEL/SPV_INTEL_fpga_reg/IntelFPGAReg.ll @@ -116,7 +116,7 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -129,7 +129,7 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nounwind define spir_func void @_Z3foov() #3 { @@ -303,10 +303,10 @@ entry: declare i32 @llvm.annotation.i32.p1(i32, ptr addrspace(1), ptr addrspace(1), i32) #4 ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p0.p0.i64(ptr nocapture writeonly, ptr nocapture readonly, i64, i1 immarg) #1 +declare void @llvm.memcpy.p0.p0.i64(ptr captures(none) writeonly, ptr captures(none) readonly, i64, i1 immarg) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p0.p1.i64(ptr nocapture writeonly, ptr addrspace(1) nocapture readonly, i64, i1 immarg) #1 +declare void @llvm.memcpy.p0.p1.i64(ptr captures(none) writeonly, ptr addrspace(1) captures(none) readonly, i64, i1 immarg) #1 ; Function Attrs: nounwind declare ptr @llvm.ptr.annotation.p0(ptr, ptr, ptr, i32, ptr) #4 diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll index 9e07baf650..9c7e2d38a8 100644 --- a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/decor-func-ptr-arg-attr.ll @@ -16,7 +16,7 @@ ; CHECK-SPIRV: FunctionPointerCallINTEL ; CHECK-SPIRV-SAME: [[#TargetId]] -; CHECK-LLVM: call spir_func addrspace(9) void %cond.i.i(ptr noalias nocapture byval(%multi_ptr) %agg.tmp.i.i) +; CHECK-LLVM: call spir_func addrspace(9) void %cond.i.i(ptr noalias captures(none) byval(%multi_ptr) %agg.tmp.i.i) ; ModuleID = 'sycl_test.cpp' target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" @@ -31,7 +31,7 @@ target triple = "spir64-unknown-unknown" $RoundedRangeKernel = comdat any ; Function Attrs: nounwind -define spir_func void @inc_function(ptr byval(%"multi_ptr") noalias nocapture %ptr) #0 { +define spir_func void @inc_function(ptr byval(%"multi_ptr") noalias captures(none) %ptr) #0 { entry: ret void } @@ -42,7 +42,7 @@ define weak_odr dso_local spir_kernel void @RoundedRangeKernel(ptr byval(%"range entry: %agg.tmp.i.i = alloca %"multi_ptr", align 8 %cond.i.i = select i1 %_arg_, ptr @inc_function, ptr null - call spir_func void %cond.i.i(ptr nonnull byval(%"multi_ptr") align 8 noalias nocapture %agg.tmp.i.i) #1, !callees !7 + call spir_func void %cond.i.i(ptr nonnull byval(%"multi_ptr") align 8 noalias captures(none) %agg.tmp.i.i) #1, !callees !7 ret void } diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll index 67eebd988e..cd5f3a7886 100644 --- a/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/CodeSectionINTEL/select.ll @@ -88,10 +88,10 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind readnone define dso_local spir_func i32 @_Z3barii(i32 %a, i32 %b) local_unnamed_addr #2 { diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/decor-func-ptr-arg-attr.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/decor-func-ptr-arg-attr.ll index df955ff093..29b71f65c7 100644 --- a/test/extensions/INTEL/SPV_INTEL_function_pointers/decor-func-ptr-arg-attr.ll +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/decor-func-ptr-arg-attr.ll @@ -16,7 +16,7 @@ ; CHECK-SPIRV: FunctionPointerCallINTEL ; CHECK-SPIRV-SAME: [[#TargetId]] -; CHECK-LLVM: call spir_func void %cond.i.i(ptr noalias nocapture byval(%multi_ptr) %agg.tmp.i.i) +; CHECK-LLVM: call spir_func void %cond.i.i(ptr noalias captures(none) byval(%multi_ptr) %agg.tmp.i.i) ; ModuleID = 'sycl_test.cpp' target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" @@ -31,7 +31,7 @@ target triple = "spir64-unknown-unknown" $RoundedRangeKernel = comdat any ; Function Attrs: nounwind -define spir_func void @inc_function(ptr byval(%"multi_ptr") noalias nocapture %ptr) #0 { +define spir_func void @inc_function(ptr byval(%"multi_ptr") noalias captures(none) %ptr) #0 { entry: ret void } @@ -42,7 +42,7 @@ define weak_odr dso_local spir_kernel void @RoundedRangeKernel(ptr byval(%"range entry: %agg.tmp.i.i = alloca %"multi_ptr", align 8 %cond.i.i = select i1 %_arg_, ptr @inc_function, ptr null - call spir_func void %cond.i.i(ptr nonnull byval(%"multi_ptr") align 8 noalias nocapture %agg.tmp.i.i) #1, !callees !7 + call spir_func void %cond.i.i(ptr nonnull byval(%"multi_ptr") align 8 noalias captures(none) %agg.tmp.i.i) #1, !callees !7 ret void } diff --git a/test/extensions/INTEL/SPV_INTEL_function_pointers/select.ll b/test/extensions/INTEL/SPV_INTEL_function_pointers/select.ll index d9a2861807..392d467500 100644 --- a/test/extensions/INTEL/SPV_INTEL_function_pointers/select.ll +++ b/test/extensions/INTEL/SPV_INTEL_function_pointers/select.ll @@ -88,10 +88,10 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind readnone define dso_local spir_func i32 @_Z3barii(i32 %a, i32 %b) local_unnamed_addr #2 { diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/array_of_matrices.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/array_of_matrices.ll index cb6a0e586c..9bfed66fcc 100644 --- a/test/extensions/INTEL/SPV_INTEL_joint_matrix/array_of_matrices.ll +++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/array_of_matrices.ll @@ -299,10 +299,10 @@ _ZZZ12joint_matmulILj256ELj256ELj256ELj256ELj2EN4sycl3_V13ext6oneapi8bfloat16EfL } ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) declare void @llvm.assume(i1 noundef) #2 diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_apply.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_apply.ll index 59451d0c72..5298bba95b 100644 --- a/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_apply.ll +++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_apply.ll @@ -127,10 +127,10 @@ entry: } ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) ; Function Attrs: convergent nounwind declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) @_Z26__spirv_CompositeConstruct(ptr noundef byval(%"class.sycl::_V1::ext::oneapi::bfloat16") align 2) local_unnamed_addr diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_checked.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_checked.ll index 1cebf53c10..4188d75318 100644 --- a/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_checked.ll +++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_checked.ll @@ -157,10 +157,10 @@ declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3, declare dso_local spir_func void @_Z42__spirv_CooperativeMatrixStoreCheckedINTEL(ptr addrspace(4) noundef, i32 noundef, i32 noundef, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef, i32 noundef, i32 noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3 attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" } attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) } diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_prefetch.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_prefetch.ll index 480832d666..39468f13f7 100644 --- a/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_prefetch.ll +++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_prefetch.ll @@ -165,10 +165,10 @@ declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3, declare dso_local spir_func void @_Z33__spirv_CooperativeMatrixStoreKHR(ptr addrspace(4) noundef, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3 attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" } attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) } diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix.ll index 066ea5aeb9..81d21ab9b0 100644 --- a/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix.ll +++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix.ll @@ -148,10 +148,10 @@ declare dso_local spir_func noundef target("spirv.JointMatrixINTEL", i32, 12, 12 declare dso_local spir_func void @_Z29__spirv_JointMatrixStoreINTELIiLm12ELm12ELN5__spv9MatrixUseE2ELNS0_12MatrixLayoutE3ELNS0_5Scope4FlagE3EEvPT_PNS0_24__spirv_JointMatrixINTELIS5_XT0_EXT1_EXT3_EXT4_EXT2_EEEmS2_S4_i(ptr addrspace(4) noundef, target("spirv.JointMatrixINTEL", i32, 12, 12, 3, 3, 2) noundef, i64 noundef, i32 noundef, i32 noundef, i32 noundef) local_unnamed_addr #2 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3 attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" } attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) } diff --git a/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_checked.ll b/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_checked.ll index 509b6815f5..c8186062cf 100644 --- a/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_checked.ll +++ b/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_checked.ll @@ -153,10 +153,10 @@ declare dso_local spir_func noundef target("spirv.JointMatrixINTEL", i32, 12, 12 declare dso_local spir_func void @_Z42__spirv_CooperativeMatrixStoreCheckedINTEL(ptr addrspace(4) noundef, i32 noundef, i32 noundef, target("spirv.JointMatrixINTEL", i32, 12, 12, 3, 3, 2) noundef, i32 noundef, i32 noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3 attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" } attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) } diff --git a/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll b/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll index 67fa12212b..b8e56a78ed 100644 --- a/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll +++ b/test/extensions/INTEL/SPV_INTEL_kernel_attributes/intel_fpga_function_attributes.ll @@ -102,7 +102,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nounwind define linkonce_odr spir_func void @_ZN3FooclEv(ptr addrspace(4) %this) #2 comdat align 2 { @@ -114,7 +114,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nounwind define spir_kernel void @_ZTSZ3barvE12kernel_name2() #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !4 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !4 { entry: diff --git a/test/extensions/INTEL/SPV_INTEL_memory_access_aliasing/intel-alias-lifetime.ll b/test/extensions/INTEL/SPV_INTEL_memory_access_aliasing/intel-alias-lifetime.ll index 6d42f83d7f..3dffbc303d 100644 --- a/test/extensions/INTEL/SPV_INTEL_memory_access_aliasing/intel-alias-lifetime.ll +++ b/test/extensions/INTEL/SPV_INTEL_memory_access_aliasing/intel-alias-lifetime.ll @@ -20,7 +20,7 @@ define spir_kernel void @lifetime_simple() } ; Function Attrs: nounwind -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #0 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #0 attributes #0 = { nounwind } diff --git a/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_12.ll b/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_12.ll index dcb61da3d4..9591808a86 100644 --- a/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_12.ll +++ b/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_12.ll @@ -58,7 +58,7 @@ target triple = "spir64" ; CHECK-LLVM-LABEL: define spir_kernel void @test ; Function Attrs: convergent norecurse nounwind -define dso_local spir_kernel void @test(ptr addrspace(1) nocapture noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @test(ptr addrspace(1) captures(none) noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { tail call spir_func void @_Z31intel_work_group_barrier_arrivej(i32 noundef 1) #2 ; CHECK-LLVM: call spir_func void @_Z31intel_work_group_barrier_arrivej(i32 1) tail call spir_func void @_Z29intel_work_group_barrier_waitj(i32 noundef 1) #2 diff --git a/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll b/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll index 36a196ecfc..eaa0fa27b8 100644 --- a/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll +++ b/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll @@ -100,7 +100,7 @@ target triple = "spir64" ; CHECK-LLVM-LABEL: define spir_kernel void @test ; Function Attrs: convergent norecurse nounwind -define dso_local spir_kernel void @test(ptr addrspace(1) nocapture noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @test(ptr addrspace(1) captures(none) noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { tail call spir_func void @_Z31intel_work_group_barrier_arrivej(i32 noundef 1) #2 ; CHECK-LLVM: call spir_func void @_Z31intel_work_group_barrier_arrivej12memory_scope(i32 1, i32 1) tail call spir_func void @_Z29intel_work_group_barrier_waitj(i32 noundef 1) #2 diff --git a/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_spirv.ll b/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_spirv.ll index 0d0fa1835c..f64d9010d6 100644 --- a/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_spirv.ll +++ b/test/extensions/INTEL/SPV_INTEL_split_barrier/split_work_group_barrier_spirv.ll @@ -101,7 +101,7 @@ target triple = "spir64" ; CHECK-LLVM-LABEL: define spir_kernel void @test ; Function Attrs: convergent norecurse nounwind -define dso_local spir_kernel void @test(ptr addrspace(1) nocapture noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @test(ptr addrspace(1) captures(none) noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { tail call spir_func void @_Z33__spirv_ControlBarrierArriveINTELiii(i32 noundef 2, i32 noundef 2, i32 noundef 260) #2 ; CHECK-LLVM: call spir_func void @_Z33__spirv_ControlBarrierArriveINTELiii(i32 2, i32 2, i32 260) #1 tail call spir_func void @_Z31__spirv_ControlBarrierWaitINTELiii(i32 noundef 2, i32 noundef 2, i32 noundef 258) #2 diff --git a/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/InfiniteLoopMetadataPlacement.ll b/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/InfiniteLoopMetadataPlacement.ll index d3ba17e04b..0edc362faf 100644 --- a/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/InfiniteLoopMetadataPlacement.ll +++ b/test/extensions/INTEL/SPV_INTEL_unstructured_loop_controls/InfiniteLoopMetadataPlacement.ll @@ -78,10 +78,10 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 attributes #0 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { argmemonly nounwind } diff --git a/test/extensions/INTEL/SPV_INTEL_usm_storage_classes/intel_usm_addrspaces.ll b/test/extensions/INTEL/SPV_INTEL_usm_storage_classes/intel_usm_addrspaces.ll index 8043f4a7ad..437b04eaea 100644 --- a/test/extensions/INTEL/SPV_INTEL_usm_storage_classes/intel_usm_addrspaces.ll +++ b/test/extensions/INTEL/SPV_INTEL_usm_storage_classes/intel_usm_addrspaces.ll @@ -48,7 +48,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -60,7 +60,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind define spir_func void @_Z6usagesv() #3 { diff --git a/test/extensions/INTEL/SPV_INTEL_variable_length_array/basic.ll b/test/extensions/INTEL/SPV_INTEL_variable_length_array/basic.ll index a41de9d8ba..e89b9ff958 100644 --- a/test/extensions/INTEL/SPV_INTEL_variable_length_array/basic.ll +++ b/test/extensions/INTEL/SPV_INTEL_variable_length_array/basic.ll @@ -76,13 +76,13 @@ entry: ret i32 %add5 } -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 declare ptr @llvm.stacksave.p0() #2 declare void @llvm.stackrestore.p0(ptr) #2 -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 attributes #0 = { nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { argmemonly nounwind willreturn } diff --git a/test/extensions/INTEL/SPV_INTEL_variable_length_array/vla_spec_const.ll b/test/extensions/INTEL/SPV_INTEL_variable_length_array/vla_spec_const.ll index a024cbfdfa..ed02851892 100644 --- a/test/extensions/INTEL/SPV_INTEL_variable_length_array/vla_spec_const.ll +++ b/test/extensions/INTEL/SPV_INTEL_variable_length_array/vla_spec_const.ll @@ -78,7 +78,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlvE_clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -104,7 +104,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse define linkonce_odr dso_local spir_func i64 @_ZNK2cl4sycl12experimental13spec_constantIm13MyUInt64ConstE3getEv(ptr addrspace(4) %this) #3 comdat align 2 { diff --git a/test/extensions/KHR/SPV_KHR_cooperative_matrix/array_of_matrices.ll b/test/extensions/KHR/SPV_KHR_cooperative_matrix/array_of_matrices.ll index 7df3d61f6e..a5b4df2bed 100644 --- a/test/extensions/KHR/SPV_KHR_cooperative_matrix/array_of_matrices.ll +++ b/test/extensions/KHR/SPV_KHR_cooperative_matrix/array_of_matrices.ll @@ -300,10 +300,10 @@ _ZZZ12joint_matmulILj256ELj256ELj256ELj256ELj2EN4sycl3_V13ext6oneapi8bfloat16EfL } ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write) declare void @llvm.assume(i1 noundef) #2 diff --git a/test/extensions/KHR/SPV_KHR_cooperative_matrix/cooperative_matrix.ll b/test/extensions/KHR/SPV_KHR_cooperative_matrix/cooperative_matrix.ll index 4e99ab4ccb..9e055c7810 100644 --- a/test/extensions/KHR/SPV_KHR_cooperative_matrix/cooperative_matrix.ll +++ b/test/extensions/KHR/SPV_KHR_cooperative_matrix/cooperative_matrix.ll @@ -165,10 +165,10 @@ declare dso_local spir_func void @_Z33__spirv_CooperativeMatrixStoreKHR(ptr addr declare dso_local spir_func void @_Z33__spirv_CooperativeMatrixStoreKHR_2(ptr addrspace(1) noundef, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3 ; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #3 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3 attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" } attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) } diff --git a/test/extensions/KHR/SPV_KHR_expect_assume/assume.ll b/test/extensions/KHR/SPV_KHR_expect_assume/assume.ll index 4c28501195..d7bd9c121d 100644 --- a/test/extensions/KHR/SPV_KHR_expect_assume/assume.ll +++ b/test/extensions/KHR/SPV_KHR_expect_assume/assume.ll @@ -84,10 +84,10 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #4 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #4 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #4 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #4 attributes #0 = { nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { nounwind willreturn } diff --git a/test/extensions/KHR/SPV_KHR_expect_assume/expect.ll b/test/extensions/KHR/SPV_KHR_expect_assume/expect.ll index b3110b2f01..9e7c23077a 100644 --- a/test/extensions/KHR/SPV_KHR_expect_assume/expect.ll +++ b/test/extensions/KHR/SPV_KHR_expect_assume/expect.ll @@ -89,7 +89,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr addrspace(4) %this) #2 align 2 { @@ -111,7 +111,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: norecurse nounwind define spir_func i32 @_Z12expect_consti(i32 %x) #3 { diff --git a/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-all-function-attributes-crash.ll b/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-all-function-attributes-crash.ll index 8cbc6b5851..a18b844755 100644 --- a/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-all-function-attributes-crash.ll +++ b/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-all-function-attributes-crash.ll @@ -15,7 +15,7 @@ define spir_kernel void @test_array(ptr addrspace(1) %in, ptr addrspace(1) %out) } ; Function Attrs: nounwind -declare void @llvm.memmove.p1.p1.i32(ptr addrspace(1) nocapture, ptr addrspace(1) nocapture readonly, i32, i1) #0 +declare void @llvm.memmove.p1.p1.i32(ptr addrspace(1) captures(none), ptr addrspace(1) captures(none) readonly, i32, i1) #0 ; CHECK-SPIRV: Name [[#ID:]] "llvm.memmove.p1.p1.i32" ; CHECK-LLVM-NOT: llvm.memmove diff --git a/test/group-decorate.spt b/test/group-decorate.spt index ae1d2a6a97..c327c20050 100644 --- a/test/group-decorate.spt +++ b/test/group-decorate.spt @@ -35,4 +35,4 @@ ; RUN: llvm-spirv -r %t.spv -o %t.bc ; RUN: llvm-dis < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM -; CHECK-LLVM: define spir_kernel void @test(ptr addrspace(1) nocapture %src1, ptr addrspace(1) nocapture %src2, ptr addrspace(1) nocapture %dst) {{.*}} +; CHECK-LLVM: define spir_kernel void @test(ptr addrspace(1) captures(none) %src1, ptr addrspace(1) captures(none) %src2, ptr addrspace(1) captures(none) %dst) {{.*}} diff --git a/test/link-attribute.ll b/test/link-attribute.ll index 5d258096dc..36b91c4d60 100644 --- a/test/link-attribute.ll +++ b/test/link-attribute.ll @@ -19,7 +19,7 @@ target triple = "spir64-unknown-unknown" @imageSampler = addrspace(2) constant i32 36, align 4 ; Function Attrs: nounwind -define spir_kernel void @sample_kernel(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, ptr addrspace(1) nocapture %xOffsets, ptr addrspace(1) nocapture %yOffsets, ptr addrspace(1) nocapture %results) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @sample_kernel(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, ptr addrspace(1) captures(none) %xOffsets, ptr addrspace(1) captures(none) %yOffsets, ptr addrspace(1) captures(none) %results) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { %1 = tail call spir_func i64 @_Z13get_global_idj(i32 0) #1 %2 = trunc i64 %1 to i32 %3 = tail call spir_func i64 @_Z13get_global_idj(i32 1) #1 diff --git a/test/llvm-intrinsics/bitreverse.ll b/test/llvm-intrinsics/bitreverse.ll index 3537d7e7d4..1b112d2ba6 100644 --- a/test/llvm-intrinsics/bitreverse.ll +++ b/test/llvm-intrinsics/bitreverse.ll @@ -111,7 +111,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir-unknown-unknown" ; Function Attrs: convergent nounwind writeonly -define spir_kernel void @testBitRev(i8 %a, i16 %b, i32 %c, i64 %d, ptr addrspace(1) nocapture %res) { +define spir_kernel void @testBitRev(i8 %a, i16 %b, i32 %c, i64 %d, ptr addrspace(1) captures(none) %res) { entry: %call8 = call i8 @llvm.bitreverse.i8(i8 %a) store i8 %call8, ptr addrspace(1) %res, align 2, !tbaa !7 @@ -128,7 +128,7 @@ entry: ret void } -define spir_kernel void @testBitRevV2(<2 x i8> %a, <2 x i16> %b, <2 x i32> %c, <2 x i64> %d, ptr addrspace(1) nocapture %res) { +define spir_kernel void @testBitRevV2(<2 x i8> %a, <2 x i16> %b, <2 x i32> %c, <2 x i64> %d, ptr addrspace(1) captures(none) %res) { entry: %call8 = call <2 x i8> @llvm.bitreverse.v2i8(<2 x i8> %a) store <2 x i8> %call8, ptr addrspace(1) %res, align 2, !tbaa !7 @@ -145,7 +145,7 @@ entry: ret void } -define spir_kernel void @testBitRevV3(<3 x i8> %a, <3 x i16> %b, <3 x i32> %c, <3 x i64> %d, ptr addrspace(1) nocapture %res) { +define spir_kernel void @testBitRevV3(<3 x i8> %a, <3 x i16> %b, <3 x i32> %c, <3 x i64> %d, ptr addrspace(1) captures(none) %res) { entry: %call8 = call <3 x i8> @llvm.bitreverse.v3i8(<3 x i8> %a) store <3 x i8> %call8, ptr addrspace(1) %res, align 2, !tbaa !7 @@ -162,7 +162,7 @@ entry: ret void } -define spir_kernel void @testBitRevV4(<4 x i8> %a, <4 x i16> %b, <4 x i32> %c, <4 x i64> %d, ptr addrspace(1) nocapture %res) { +define spir_kernel void @testBitRevV4(<4 x i8> %a, <4 x i16> %b, <4 x i32> %c, <4 x i64> %d, ptr addrspace(1) captures(none) %res) { entry: %call8 = call <4 x i8> @llvm.bitreverse.v4i8(<4 x i8> %a) store <4 x i8> %call8, ptr addrspace(1) %res, align 2, !tbaa !7 @@ -179,7 +179,7 @@ entry: ret void } -define spir_kernel void @testBitRevV8(<8 x i8> %a, <8 x i16> %b, <8 x i32> %c, <8 x i64> %d, ptr addrspace(1) nocapture %res) { +define spir_kernel void @testBitRevV8(<8 x i8> %a, <8 x i16> %b, <8 x i32> %c, <8 x i64> %d, ptr addrspace(1) captures(none) %res) { entry: %call8 = call <8 x i8> @llvm.bitreverse.v8i8(<8 x i8> %a) store <8 x i8> %call8, ptr addrspace(1) %res, align 2, !tbaa !7 @@ -196,7 +196,7 @@ entry: ret void } -define spir_kernel void @testBitRevV16(<16 x i8> %a, <16 x i16> %b, <16 x i32> %c, <16 x i64> %d, ptr addrspace(1) nocapture %res) { +define spir_kernel void @testBitRevV16(<16 x i8> %a, <16 x i16> %b, <16 x i32> %c, <16 x i64> %d, ptr addrspace(1) captures(none) %res) { entry: %call8 = call <16 x i8> @llvm.bitreverse.v16i8(<16 x i8> %a) store <16 x i8> %call8, ptr addrspace(1) %res, align 2, !tbaa !7 diff --git a/test/llvm-intrinsics/dynamic-memmove.ll b/test/llvm-intrinsics/dynamic-memmove.ll index 04bd2333c1..335aef815d 100644 --- a/test/llvm-intrinsics/dynamic-memmove.ll +++ b/test/llvm-intrinsics/dynamic-memmove.ll @@ -11,7 +11,7 @@ target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" target triple = "spir64-unknown-unknown" -declare void @llvm.memmove.p1.p1.i64(ptr addrspace(1) nocapture, ptr addrspace(1) nocapture readonly, i64, i1) +declare void @llvm.memmove.p1.p1.i64(ptr addrspace(1) captures(none), ptr addrspace(1) captures(none) readonly, i64, i1) define spir_func void @memmove_caller(ptr addrspace(1) %dst, ptr addrspace(1) %src, i64 %n) { entry: diff --git a/test/llvm-intrinsics/instrprof.ll b/test/llvm-intrinsics/instrprof.ll index 910eb646f5..ccf4e93b1a 100644 --- a/test/llvm-intrinsics/instrprof.ll +++ b/test/llvm-intrinsics/instrprof.ll @@ -46,7 +46,7 @@ declare void @llvm.instrprof.increment.step(ptr, i64, i32, i32, i64) #1 declare void @llvm.instrprof.value.profile(ptr, i64, i64, i32, i32) #1 ; Function Attrs: nounwind -declare void @llvm.memcpy.p0.p1.i64(ptr nocapture, ptr addrspace(1) nocapture readonly, i64, i1) #1 +declare void @llvm.memcpy.p0.p1.i64(ptr captures(none), ptr addrspace(1) captures(none) readonly, i64, i1) #1 attributes #0 = { convergent mustprogress norecurse } attributes #1 = { nounwind } diff --git a/test/llvm-intrinsics/invariant.ll b/test/llvm-intrinsics/invariant.ll index 8c31938cbb..34efc154f7 100644 --- a/test/llvm-intrinsics/invariant.ll +++ b/test/llvm-intrinsics/invariant.ll @@ -14,10 +14,10 @@ target triple = "spir64-unknown-linux" @WGSharedVar = internal addrspace(3) constant i64 0, align 8 ; Function Attrs: argmemonly nounwind -declare ptr @llvm.invariant.start.p3(i64 immarg, ptr addrspace(3) nocapture) #0 +declare ptr @llvm.invariant.start.p3(i64 immarg, ptr addrspace(3) captures(none)) #0 ; Function Attrs: argmemonly nounwind -declare void @llvm.invariant.end.p3(ptr, i64 immarg, ptr addrspace(3) nocapture) #0 +declare void @llvm.invariant.end.p3(ptr, i64 immarg, ptr addrspace(3) captures(none)) #0 define linkonce_odr dso_local spir_func void @func() { store i64 2, ptr addrspace(3) @WGSharedVar diff --git a/test/llvm-intrinsics/lifetime.ll b/test/llvm-intrinsics/lifetime.ll index c1fd8b939a..e1f1fce175 100644 --- a/test/llvm-intrinsics/lifetime.ll +++ b/test/llvm-intrinsics/lifetime.ll @@ -61,7 +61,7 @@ target triple = "spir64-unknown-unknown" %class.anon = type { i8 } ; Function Attrs: nounwind -define spir_kernel void @lifetime_simple(i32 addrspace(1)* nocapture %res, i32 addrspace(1)* nocapture %lhs, i32 addrspace(1)* nocapture %rhs) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @lifetime_simple(i32 addrspace(1)* captures(none) %res, i32 addrspace(1)* captures(none) %lhs, i32 addrspace(1)* captures(none) %rhs) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { %1 = alloca i32 %2 = call spir_func i64 @_Z13get_global_idj(i32 0) #1 %3 = shl i64 %2, 32 @@ -95,10 +95,10 @@ entry: declare spir_func void @foo(%class.anon* %this) #0 ; Function Attrs: nounwind -declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #0 +declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* captures(none)) #0 ; Function Attrs: nounwind -declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #0 +declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* captures(none)) #0 ; Function Attrs: nounwind readnone declare spir_func i64 @_Z13get_global_idj(i32) #1 @@ -117,10 +117,10 @@ entry: declare spir_func void @boo(%class.anon addrspace(4)* %this) #0 ; Function Attrs: nounwind -declare void @llvm.lifetime.start.p4i8(i64 immarg, i8 addrspace(4)* nocapture) #0 +declare void @llvm.lifetime.start.p4i8(i64 immarg, i8 addrspace(4)* captures(none)) #0 ; Function Attrs: nounwind -declare void @llvm.lifetime.end.p4i8(i64 immarg, i8 addrspace(4)* nocapture) #0 +declare void @llvm.lifetime.end.p4i8(i64 immarg, i8 addrspace(4)* captures(none)) #0 attributes #0 = { nounwind } diff --git a/test/llvm-intrinsics/memcpy.align.ll b/test/llvm-intrinsics/memcpy.align.ll index c886c3326e..bca3d485b2 100644 --- a/test/llvm-intrinsics/memcpy.align.ll +++ b/test/llvm-intrinsics/memcpy.align.ll @@ -73,16 +73,16 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0i8(i64, i8* nocapture) #1 +declare void @llvm.lifetime.start.p0i8(i64, i8* captures(none)) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p0i8.p2i8.i32(i8* nocapture writeonly, i8 addrspace(2)* nocapture readonly, i32, i1) #1 +declare void @llvm.memcpy.p0i8.p2i8.i32(i8* captures(none) writeonly, i8 addrspace(2)* captures(none) readonly, i32, i1) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p0i8.p0i8.i32(i8* nocapture writeonly, i8* nocapture readonly, i32, i1) #1 +declare void @llvm.memcpy.p0i8.p0i8.i32(i8* captures(none) writeonly, i8* captures(none) readonly, i32, i1) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0i8(i64, i8* nocapture) #1 +declare void @llvm.lifetime.end.p0i8(i64, i8* captures(none)) #1 ; Function Attrs: convergent nounwind define spir_func void @bar(%struct.B* noalias sret(%struct.B) %agg.result) #0 { diff --git a/test/llvm-intrinsics/memmove.ll b/test/llvm-intrinsics/memmove.ll index 21472a2e4e..33a18a8daa 100644 --- a/test/llvm-intrinsics/memmove.ll +++ b/test/llvm-intrinsics/memmove.ll @@ -143,14 +143,14 @@ target triple = "spir-unknown-unknown" @"func_object1" = internal addrspace(3) global %class.kfunc zeroinitializer, align 8 ; Function Attrs: nounwind -define spir_kernel void @test_full_move(%struct.SomeStruct addrspace(1)* nocapture readonly %in, %struct.SomeStruct addrspace(1)* nocapture %out) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @test_full_move(%struct.SomeStruct addrspace(1)* captures(none) readonly %in, %struct.SomeStruct addrspace(1)* captures(none) %out) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { %1 = bitcast %struct.SomeStruct addrspace(1)* %in to i8 addrspace(1)* %2 = bitcast %struct.SomeStruct addrspace(1)* %out to i8 addrspace(1)* call void @llvm.memmove.p1i8.p1i8.i32(i8 addrspace(1)* align 64 %2, i8 addrspace(1)* align 64 %1, i32 128, i1 false) ret void } -define spir_kernel void @test_partial_move(%struct.SomeStruct addrspace(1)* nocapture readonly %in, %struct.SomeStruct addrspace(4)* nocapture %out) { +define spir_kernel void @test_partial_move(%struct.SomeStruct addrspace(1)* captures(none) readonly %in, %struct.SomeStruct addrspace(4)* captures(none) %out) { %1 = bitcast %struct.SomeStruct addrspace(1)* %in to i8 addrspace(1)* %2 = bitcast %struct.SomeStruct addrspace(4)* %out to i8 addrspace(4)* %3 = addrspacecast i8 addrspace(4)* %2 to i8 addrspace(1)* @@ -189,10 +189,10 @@ merge: ; preds = %entry.merge_crit_ed } ; Function Attrs: argmemonly nofree nosync nounwind willreturn -declare void @llvm.memmove.p4i8.p4i8.i64(i8 addrspace(4)* nocapture writeonly, i8 addrspace(4)* nocapture readonly, i64, i1 immarg) +declare void @llvm.memmove.p4i8.p4i8.i64(i8 addrspace(4)* captures(none) writeonly, i8 addrspace(4)* captures(none) readonly, i64, i1 immarg) ; Function Attrs: nounwind -declare void @llvm.memmove.p1i8.p1i8.i32(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i32, i1) #1 +declare void @llvm.memmove.p1i8.p1i8.i32(i8 addrspace(1)* captures(none), i8 addrspace(1)* captures(none) readonly, i32, i1) #1 attributes #0 = { nounwind "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no_infs-fp-math"="false" "no-nans-fp-math"="false" "no-realign-stack" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { nounwind } diff --git a/test/llvm-intrinsics/memset-align.ll b/test/llvm-intrinsics/memset-align.ll index 37458e2c18..afafd8ff1b 100644 --- a/test/llvm-intrinsics/memset-align.ll +++ b/test/llvm-intrinsics/memset-align.ll @@ -33,6 +33,6 @@ entry: } ; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write) -declare void @llvm.memset.p4.i64(ptr addrspace(4) nocapture writeonly, i8, i64, i1 immarg) #0 +declare void @llvm.memset.p4.i64(ptr addrspace(4) captures(none) writeonly, i8, i64, i1 immarg) #0 attributes #0 = { nocallback nofree nounwind willreturn memory(argmem: write) } diff --git a/test/llvm-intrinsics/memset-opaque.ll b/test/llvm-intrinsics/memset-opaque.ll index 6bf32da53e..2bd235e0cb 100644 --- a/test/llvm-intrinsics/memset-opaque.ll +++ b/test/llvm-intrinsics/memset-opaque.ll @@ -72,7 +72,7 @@ target triple = "spir" ; CHECK-LLVM-OPAQUE: internal unnamed_addr addrspace(2) constant [4 x i8] c"\15\15\15\15" ; Function Attrs: nounwind -define spir_func void @_Z5foo11v(ptr addrspace(4) noalias nocapture sret(%struct.S1) %agg.result, i32 %s1, i64 %s2, i8 %v) #0 { +define spir_func void @_Z5foo11v(ptr addrspace(4) noalias captures(none) sret(%struct.S1) %agg.result, i32 %s1, i64 %s2, i8 %v) #0 { %x = alloca [4 x i8] tail call void @llvm.memset.p4.i32(ptr addrspace(4) align 4 %agg.result, i8 0, i32 12, i1 false) ; CHECK-LLVM-OPAQUE: call void @llvm.memcpy.p4.p2.i32(ptr addrspace(4) align 4 %1, ptr addrspace(2) align 4 %2, i32 12, i1 false) @@ -102,10 +102,10 @@ define spir_func void @_Z5foo11v(ptr addrspace(4) noalias nocapture sret(%struct } ; Function Attrs: nounwind -declare void @llvm.memset.p4.i32(ptr addrspace(4) nocapture, i8, i32, i1) #1 +declare void @llvm.memset.p4.i32(ptr addrspace(4) captures(none), i8, i32, i1) #1 ; Function Attrs: nounwind -declare void @llvm.memset.p0.i32(ptr nocapture, i8, i32, i1) #1 +declare void @llvm.memset.p0.i32(ptr captures(none), i8, i32, i1) #1 ; Function Attrs: nounwind declare void @llvm.memset.p3.i32(ptr addrspace(3), i8, i32, i1) #1 diff --git a/test/llvm-intrinsics/sadd.with.overflow.ll b/test/llvm-intrinsics/sadd.with.overflow.ll index 7897b30677..5192b9cb67 100644 --- a/test/llvm-intrinsics/sadd.with.overflow.ll +++ b/test/llvm-intrinsics/sadd.with.overflow.ll @@ -27,7 +27,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir" ; Function Attrs: nounwind -define dso_local spir_func void @foo32(i32 %a, i32 %b, ptr nocapture %c) local_unnamed_addr #0 { +define dso_local spir_func void @foo32(i32 %a, i32 %b, ptr captures(none) %c) local_unnamed_addr #0 { entry: %0 = tail call { i32, i1 } @llvm.sadd.with.overflow.i32(i32 %a, i32 %b), !nosanitize !2 %1 = extractvalue { i32, i1 } %0, 1, !nosanitize !2 @@ -50,7 +50,7 @@ declare { i32, i1 } @llvm.sadd.with.overflow.i32(i32, i32) #1 declare void @llvm.trap() #2 ; Function Attrs: nounwind -define dso_local spir_func void @foo64(i64 %a, i64 %b, ptr nocapture %c) local_unnamed_addr #0 { +define dso_local spir_func void @foo64(i64 %a, i64 %b, ptr captures(none) %c) local_unnamed_addr #0 { entry: %0 = tail call { i64, i1 } @llvm.sadd.with.overflow.i64(i64 %a, i64 %b), !nosanitize !2 %1 = extractvalue { i64, i1 } %0, 1, !nosanitize !2 diff --git a/test/llvm-intrinsics/umul.with.overflow.ll b/test/llvm-intrinsics/umul.with.overflow.ll index 6507b0a7d3..cddbd972ec 100644 --- a/test/llvm-intrinsics/umul.with.overflow.ll +++ b/test/llvm-intrinsics/umul.with.overflow.ll @@ -16,7 +16,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir" ; Function Attrs: nofree nounwind writeonly -define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, ptr nocapture %c) local_unnamed_addr #0 { +define dso_local spir_func void @_Z4foo8hhPh(i8 zeroext %a, i8 zeroext %b, ptr captures(none) %c) local_unnamed_addr #0 { entry: ; CHECK-LLVM: call { i8, i1 } @llvm.umul.with.overflow.i8 ; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_8]] @@ -39,7 +39,7 @@ entry: ; CHECK-SPIRV: ReturnValue [[INSERT_RES_1]] ; Function Attrs: nofree nounwind writeonly -define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, ptr nocapture %c) local_unnamed_addr #0 { +define dso_local spir_func void @_Z5foo32jjPj(i32 %a, i32 %b, ptr captures(none) %c) local_unnamed_addr #0 { entry: ; CHECK-LLVM: call { i32, i1 } @llvm.umul.with.overflow.i32 ; CHECK-SPIRV: FunctionCall [[#]] [[#]] [[NAME_UMUL_FUNC_32]] diff --git a/test/long-type-struct.ll b/test/long-type-struct.ll index 604cf290b3..7e426f2426 100644 --- a/test/long-type-struct.ll +++ b/test/long-type-struct.ll @@ -55,10 +55,10 @@ entry: } ; Function Attrs: argmemonly nofree nosync nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nofree nosync nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: convergent declare dso_local spir_func void @_Z3foo1A(ptr byval(%struct._ZTS1A.A) align 8) local_unnamed_addr #2 diff --git a/test/sitofp-with-bool.ll b/test/sitofp-with-bool.ll index cfdf8b1209..44270c0cd8 100644 --- a/test/sitofp-with-bool.ll +++ b/test/sitofp-with-bool.ll @@ -21,7 +21,7 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64" ; Function Attrs: nofree norecurse nounwind writeonly -define dso_local spir_kernel void @K(ptr addrspace(1) nocapture %A, i32 %B) local_unnamed_addr #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define dso_local spir_kernel void @K(ptr addrspace(1) captures(none) %A, i32 %B) local_unnamed_addr #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: %cmp = icmp sgt i32 %B, 0 %conv = sitofp i1 %cmp to float diff --git a/test/transcoding/BitReversePref.ll b/test/transcoding/BitReversePref.ll index d0e89bf980..418eeaf95a 100644 --- a/test/transcoding/BitReversePref.ll +++ b/test/transcoding/BitReversePref.ll @@ -20,10 +20,10 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: inlinehint norecurse nounwind define linkonce_odr dso_local spir_func i32 @_Z10BitReversei(i32 %value) #3 comdat { diff --git a/test/transcoding/BuildNDRange_2.ll b/test/transcoding/BuildNDRange_2.ll index 85f50255c2..c064d54c7a 100644 --- a/test/transcoding/BuildNDRange_2.ll +++ b/test/transcoding/BuildNDRange_2.ll @@ -92,7 +92,7 @@ entry: } ; Function Attrs: nounwind -declare void @llvm.memcpy.p1.p1.i64(ptr addrspace(1) nocapture, ptr addrspace(1) nocapture readonly, i64, i1) #2 +declare void @llvm.memcpy.p1.p1.i64(ptr addrspace(1) captures(none), ptr addrspace(1) captures(none) readonly, i64, i1) #2 declare spir_func void @_Z10ndrange_2DPKm(ptr sret(%struct.ndrange_t), ptr addrspace(1)) #1 diff --git a/test/transcoding/CreatePipeFromPipeStorage.ll b/test/transcoding/CreatePipeFromPipeStorage.ll index fbe0176cd2..21d9915e1c 100644 --- a/test/transcoding/CreatePipeFromPipeStorage.ll +++ b/test/transcoding/CreatePipeFromPipeStorage.ll @@ -80,7 +80,7 @@ entry: ; CHECK-SPV-IR: %[[PIPE_STORAGE_1_BC:[0-9]+]] = bitcast ptr addrspace(1) %[[PIPE_STORAGE_1]] to ptr addrspace(1) ; CHECK-SPV-IR: %[[WRITE_PIPE:[0-9]+]] = call spir_func target("spirv.Pipe", 1) @_Z39__spirv_CreatePipeFromPipeStorage_writePU3AS119__spirv_PipeStorage(ptr addrspace(1) %[[PIPE_STORAGE_1_BC]]) ; CHECK-SPV-IR: %[[WRITE_PIPE_WRAPPER:[0-9]+]] = addrspacecast ptr %mywpipe to ptr addrspace(4) - ; CHECK-SPV-IR: call spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) nocapture %[[WRITE_PIPE_WRAPPER]], target("spirv.Pipe", 1) %[[WRITE_PIPE]]) + ; CHECK-SPV-IR: call spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) captures(none) %[[WRITE_PIPE_WRAPPER]], target("spirv.Pipe", 1) %[[WRITE_PIPE]]) ; CHECK-SPIRV: Load {{[0-9]+}} [[PIPE_STORAGE_ID0:[0-9]+]] [[SPIRV1]] 2 4 ; CHECK-SPIRV: Bitcast {{[0-9]+}} [[PIPE_STORAGE_ID0_BC:[0-9]+]] [[PIPE_STORAGE_ID0]] @@ -98,7 +98,7 @@ entry: ; CHECK-SPV-IR: %[[PIPE_STORAGE_2_BC:[0-9]+]] = bitcast ptr addrspace(1) %[[PIPE_STORAGE_2]] to ptr addrspace(1) ; CHECK-SPV-IR: %[[READ_PIPE:[0-9]+]] = call spir_func target("spirv.Pipe", 0) @_Z38__spirv_CreatePipeFromPipeStorage_readPU3AS119__spirv_PipeStorage(ptr addrspace(1) %[[PIPE_STORAGE_2_BC]]) ; CHECK-SPV-IR: %[[READ_PIPE_WRAPPER:[0-9]+]] = addrspacecast ptr %myrpipe to ptr addrspace(4) - ; CHECK-SPV-IR: call spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) nocapture %[[READ_PIPE_WRAPPER]], target("spirv.Pipe", 0) %[[READ_PIPE]]) + ; CHECK-SPV-IR: call spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) captures(none) %[[READ_PIPE_WRAPPER]], target("spirv.Pipe", 0) %[[READ_PIPE]]) ; CHECK-SPIRV: Load {{[0-9]+}} [[PIPE_STORAGE_ID1:[0-9]+]] [[SPIRV1]] 2 4 ; CHECK-SPIRV: Bitcast {{[0-9]+}} [[PIPE_STORAGE_ID1_BC:[0-9]+]] [[PIPE_STORAGE_ID1]] @@ -116,14 +116,14 @@ entry: } ; Function Attrs: nounwind -define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) nocapture %this, target("spirv.Pipe", 0) %handle) unnamed_addr align 2 { +define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) captures(none) %this, target("spirv.Pipe", 0) %handle) unnamed_addr align 2 { entry: tail call spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC2EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) %this, target("spirv.Pipe", 0) %handle) ret void } ; Function Attrs: nounwind -define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC2EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) nocapture %this, target("spirv.Pipe", 0) %handle) unnamed_addr align 2 { +define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE0EEC2EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE0EEE(ptr addrspace(4) captures(none) %this, target("spirv.Pipe", 0) %handle) unnamed_addr align 2 { entry: %_handle = getelementptr inbounds %"class.cl::pipe", ptr addrspace(4) %this, i32 0, i32 0 store target("spirv.Pipe", 0) %handle, target("spirv.Pipe", 0) addrspace(4)* %_handle, align 4, !tbaa !11 @@ -134,14 +134,14 @@ entry: declare spir_func target("spirv.Pipe", 0) @_Z38__spirv_CreatePipeFromPipeStorage_readPU3AS1K19__spirv_PipeStorage(ptr addrspace(1)) ; Function Attrs: nounwind -define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) nocapture %this, target("spirv.Pipe", 1) %handle) unnamed_addr align 2 { +define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC1EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) captures(none) %this, target("spirv.Pipe", 1) %handle) unnamed_addr align 2 { entry: tail call spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC2EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) %this, target("spirv.Pipe", 1) %handle) ret void } ; Function Attrs: nounwind -define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC2EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) nocapture %this, target("spirv.Pipe", 1) %handle) unnamed_addr align 2 { +define linkonce_odr spir_func void @_ZNU3AS42cl4pipeIDv4_iLNS_11pipe_accessE1EEC2EPU3AS1NS_7__spirv10OpTypePipeILNS3_15AccessQualifierE1EEE(ptr addrspace(4) captures(none) %this, target("spirv.Pipe", 1) %handle) unnamed_addr align 2 { entry: %_handle = getelementptr inbounds %"class.cl::pipe", ptr addrspace(4) %this, i32 0, i32 0 store target("spirv.Pipe", 1) %handle, target("spirv.Pipe", 1) addrspace(4)* %_handle, align 4, !tbaa !13 diff --git a/test/transcoding/KernelArgTypeInOpString2.ll b/test/transcoding/KernelArgTypeInOpString2.ll index db5ad32679..6fc54036dd 100644 --- a/test/transcoding/KernelArgTypeInOpString2.ll +++ b/test/transcoding/KernelArgTypeInOpString2.ll @@ -53,7 +53,7 @@ target triple = "spir" ; Function Attrs: convergent noinline nounwind optnone -define dso_local spir_kernel void @foo(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @foo(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { store float 0.000000e+00, ptr addrspace(1) %0, align 4, !tbaa !8 ret void } diff --git a/test/transcoding/OpAllAny.ll b/test/transcoding/OpAllAny.ll index da15786586..109dbcd5e9 100644 --- a/test/transcoding/OpAllAny.ll +++ b/test/transcoding/OpAllAny.ll @@ -43,7 +43,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir" ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_vector(ptr addrspace(4) nocapture writeonly %out, <2 x i8> %c, <2 x i16> %s, <2 x i32> %i, <2 x i64> %l) local_unnamed_addr #0 { +define dso_local spir_func void @test_vector(ptr addrspace(4) captures(none) writeonly %out, <2 x i8> %c, <2 x i16> %s, <2 x i32> %i, <2 x i64> %l) local_unnamed_addr #0 { entry: %call = tail call spir_func i32 @_Z3anyDv2_c(<2 x i8> %c) #2 %call1 = tail call spir_func i32 @_Z3anyDv2_s(<2 x i16> %s) #2 diff --git a/test/transcoding/OpBitReverse_i32.ll b/test/transcoding/OpBitReverse_i32.ll index d1fe019fb8..becd3ae103 100644 --- a/test/transcoding/OpBitReverse_i32.ll +++ b/test/transcoding/OpBitReverse_i32.ll @@ -13,7 +13,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir-unknown-unknown" ; Function Attrs: convergent nounwind writeonly -define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, ptr addrspace(1) nocapture %res) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, ptr addrspace(1) captures(none) %res) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { entry: %call = tail call i32 @llvm.bitreverse.i32(i32 %b) store i32 %call, ptr addrspace(1) %res, align 4, !tbaa !7 diff --git a/test/transcoding/OpBitReverse_v2i16.ll b/test/transcoding/OpBitReverse_v2i16.ll index 3d945f8b51..4357f738bf 100644 --- a/test/transcoding/OpBitReverse_v2i16.ll +++ b/test/transcoding/OpBitReverse_v2i16.ll @@ -14,7 +14,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir-unknown-unknown" ; Function Attrs: convergent nounwind writeonly -define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, ptr addrspace(1) nocapture %res) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, ptr addrspace(1) captures(none) %res) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { entry: %call = tail call <2 x i16> @llvm.bitreverse.v2i16(<2 x i16> %b) store <2 x i16> %call, ptr addrspace(1) %res, align 4, !tbaa !7 diff --git a/test/transcoding/OpGenericPtrMemSemantics.ll b/test/transcoding/OpGenericPtrMemSemantics.ll index b630d570c4..f25d343ac6 100644 --- a/test/transcoding/OpGenericPtrMemSemantics.ll +++ b/test/transcoding/OpGenericPtrMemSemantics.ll @@ -48,7 +48,7 @@ entry: declare spir_func i32 @_Z9get_fencePU3AS4v(ptr addrspace(4)) #2 ; Function Attrs: nounwind -define spir_kernel void @testKernel(ptr addrspace(1) nocapture %results) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @testKernel(ptr addrspace(1) captures(none) %results) #1 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: %call = tail call spir_func i32 @_Z13get_global_idj(i32 0) #3 %0 = load i32, ptr addrspace(1) @gint, align 4 diff --git a/test/transcoding/OpGroupAllAny.ll b/test/transcoding/OpGroupAllAny.ll index 3d27d22154..18d723f155 100644 --- a/test/transcoding/OpGroupAllAny.ll +++ b/test/transcoding/OpGroupAllAny.ll @@ -19,7 +19,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir-unknown-unknown" ; Function Attrs: nounwind -define spir_kernel void @test(ptr addrspace(1) nocapture readnone %i) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @test(ptr addrspace(1) captures(none) readnone %i) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: %call = tail call spir_func i32 @_Z14work_group_alli(i32 5) #2 %call1 = tail call spir_func i32 @_Z14work_group_anyi(i32 5) #2 diff --git a/test/transcoding/OpImageQuerySize.ll b/test/transcoding/OpImageQuerySize.ll index 9bb27363ab..05924bf9f4 100644 --- a/test/transcoding/OpImageQuerySize.ll +++ b/test/transcoding/OpImageQuerySize.ll @@ -45,7 +45,7 @@ target triple = "spir64-unknown-unknown" ; CHECK-SPIRV-NOT: {{[0-9]*}} ExtInst {{[0-9]*}} {{[0-9]*}} {{[0-9]*}} get_image_array_size ; Function Attrs: nounwind -define spir_kernel void @test_image1d(ptr addrspace(1) nocapture %sizes, ptr addrspace(1) %img, ptr addrspace(1) %buffer, ptr addrspace(1) %array) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @test_image1d(ptr addrspace(1) captures(none) %sizes, ptr addrspace(1) %img, ptr addrspace(1) %buffer, ptr addrspace(1) %array) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image1d_ro(ptr addrspace(1) %img) #1 %2 = tail call spir_func i32 @_Z15get_image_width21ocl_image1d_buffer_ro(ptr addrspace(1) %buffer) #1 %3 = tail call spir_func i32 @_Z15get_image_width20ocl_image1d_array_ro(ptr addrspace(1) %array) #1 @@ -103,7 +103,7 @@ declare spir_func i64 @_Z20get_image_array_size20ocl_image1d_array_ro(ptr addrsp ; CHECK: shufflevector <3 x i32> {{.*}} <2 x i32> ; Function Attrs: nounwind -define spir_kernel void @test_image2d(ptr addrspace(1) nocapture %sizes, ptr addrspace(1) %img, ptr addrspace(1) nocapture %img_depth, ptr addrspace(1) %array, ptr addrspace(1) nocapture %array_depth) #0 !kernel_arg_addr_space !7 !kernel_arg_access_qual !8 !kernel_arg_type !9 !kernel_arg_base_type !11 !kernel_arg_type_qual !10 { +define spir_kernel void @test_image2d(ptr addrspace(1) captures(none) %sizes, ptr addrspace(1) %img, ptr addrspace(1) captures(none) %img_depth, ptr addrspace(1) %array, ptr addrspace(1) captures(none) %array_depth) #0 !kernel_arg_addr_space !7 !kernel_arg_access_qual !8 !kernel_arg_type !9 !kernel_arg_base_type !11 !kernel_arg_type_qual !10 { %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image2d_ro(ptr addrspace(1) %img) #1 %2 = tail call spir_func i32 @_Z16get_image_height14ocl_image2d_ro(ptr addrspace(1) %img) #1 %3 = tail call spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(ptr addrspace(1) %img) #1 @@ -168,7 +168,7 @@ declare spir_func <2 x i32> @_Z13get_image_dim20ocl_image2d_array_ro(ptr addrspa ; CHECK: shufflevector <3 x i32> {{.*}} <4 x i32> ; Function Attrs: nounwind -define spir_kernel void @test_image3d(ptr addrspace(1) nocapture %sizes, ptr addrspace(1) %img) #0 !kernel_arg_addr_space !13 !kernel_arg_access_qual !14 !kernel_arg_type !15 !kernel_arg_base_type !17 !kernel_arg_type_qual !16 { +define spir_kernel void @test_image3d(ptr addrspace(1) captures(none) %sizes, ptr addrspace(1) %img) #0 !kernel_arg_addr_space !13 !kernel_arg_access_qual !14 !kernel_arg_type !15 !kernel_arg_base_type !17 !kernel_arg_type_qual !16 { %1 = tail call spir_func i32 @_Z15get_image_width14ocl_image3d_ro(ptr addrspace(1) %img) #1 %2 = tail call spir_func i32 @_Z16get_image_height14ocl_image3d_ro(ptr addrspace(1) %img) #1 %3 = tail call spir_func i32 @_Z15get_image_depth14ocl_image3d_ro(ptr addrspace(1) %img) #1 @@ -216,7 +216,7 @@ declare spir_func <4 x i32> @_Z13get_image_dim14ocl_image3d_ro(ptr addrspace(1)) ; CHECK: extractelement <3 x i32> {{.*}} 1 ; Function Attrs: nounwind -define spir_kernel void @test_image2d_array_depth_t(ptr addrspace(1) nocapture %sizes, ptr addrspace(1) %array) #0 !kernel_arg_addr_space !27 !kernel_arg_access_qual !28 !kernel_arg_type !29 !kernel_arg_base_type !31 !kernel_arg_type_qual !30 { +define spir_kernel void @test_image2d_array_depth_t(ptr addrspace(1) captures(none) %sizes, ptr addrspace(1) %array) #0 !kernel_arg_addr_space !27 !kernel_arg_access_qual !28 !kernel_arg_type !29 !kernel_arg_base_type !31 !kernel_arg_type_qual !30 { %1 = tail call spir_func i32 @_Z15get_image_width26ocl_image2d_array_depth_ro(ptr addrspace(1) %array) #1 %2 = tail call spir_func i32 @_Z16get_image_height26ocl_image2d_array_depth_ro(ptr addrspace(1) %array) #1 %3 = tail call spir_func i64 @_Z20get_image_array_size26ocl_image2d_array_depth_ro(ptr addrspace(1) %array) #1 diff --git a/test/transcoding/OpImageReadMS.ll b/test/transcoding/OpImageReadMS.ll index 28eac4fb1f..6d78e4a32f 100644 --- a/test/transcoding/OpImageReadMS.ll +++ b/test/transcoding/OpImageReadMS.ll @@ -15,7 +15,7 @@ target triple = "spir-unknown-unknown" %opencl.image2d_msaa_ro_t = type opaque ; Function Attrs: nounwind -define spir_kernel void @sample_test(ptr addrspace(1) %source, i32 %sampler, ptr addrspace(1) nocapture %results) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @sample_test(ptr addrspace(1) %source, i32 %sampler, ptr addrspace(1) captures(none) %results) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: %call = tail call spir_func i32 @_Z13get_global_idj(i32 0) #2 %call1 = tail call spir_func i32 @_Z13get_global_idj(i32 1) #2 diff --git a/test/transcoding/OpVectorExtractDynamic.ll b/test/transcoding/OpVectorExtractDynamic.ll index 116135b878..ba6486997c 100644 --- a/test/transcoding/OpVectorExtractDynamic.ll +++ b/test/transcoding/OpVectorExtractDynamic.ll @@ -21,7 +21,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir-unknown-unknown" ; Function Attrs: nounwind -define spir_kernel void @test(ptr addrspace(1) nocapture %out, <2 x float> %vec, i32 %index) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @test(ptr addrspace(1) captures(none) %out, <2 x float> %vec, i32 %index) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: %res = extractelement <2 x float> %vec, i32 %index store float %res, ptr addrspace(1) %out, align 4 diff --git a/test/transcoding/OpenCL/atomic_syncscope_test.ll b/test/transcoding/OpenCL/atomic_syncscope_test.ll index 9ff8de90fd..46ebd96603 100644 --- a/test/transcoding/OpenCL/atomic_syncscope_test.ll +++ b/test/transcoding/OpenCL/atomic_syncscope_test.ll @@ -52,7 +52,7 @@ target triple = "spir64" ; CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 3) ; CHECK-LLVM: call spir_func i32 @_Z20atomic_load_explicitPU3AS4VU7_Atomici12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 4) -define dso_local void @fi1(ptr addrspace(4) nocapture noundef readonly %i) local_unnamed_addr #0 { +define dso_local void @fi1(ptr addrspace(4) captures(none) noundef readonly %i) local_unnamed_addr #0 { entry: %0 = load atomic i32, ptr addrspace(4) %i syncscope("workgroup") seq_cst, align 4 %1 = load atomic i32, ptr addrspace(4) %i syncscope("device") seq_cst, align 4 @@ -67,7 +67,7 @@ entry: ; CHECK-LLVM: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 4) ; CHECK-LLVM: call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 5, i32 1) -define dso_local void @test_addr(ptr addrspace(1) nocapture noundef writeonly %ig, ptr addrspace(3) nocapture noundef writeonly %il) local_unnamed_addr #0 { +define dso_local void @test_addr(ptr addrspace(1) captures(none) noundef writeonly %ig, ptr addrspace(3) captures(none) noundef writeonly %il) local_unnamed_addr #0 { entry: store atomic i32 1, ptr addrspace(1) %ig syncscope("subgroup") seq_cst, align 4 store atomic i32 1, ptr addrspace(3) %il syncscope("workgroup") seq_cst, align 4 @@ -87,7 +87,7 @@ entry: ; CHECK-LLVM: call spir_func i32 @_Z25atomic_fetch_min_explicitPU3AS4VU7_Atomicjj12memory_order12memory_scope(ptr{{.*}}, i32 1, i32 5, i32 1) ; CHECK-LLVM: call spir_func i32 @_Z25atomic_fetch_max_explicitPU3AS4VU7_Atomicjj12memory_order12memory_scope(ptr{{.*}}, i32 1, i32 5, i32 1) -define dso_local void @fi3(ptr nocapture noundef %i, ptr nocapture noundef %ui) local_unnamed_addr #0 { +define dso_local void @fi3(ptr captures(none) noundef %i, ptr captures(none) noundef %ui) local_unnamed_addr #0 { entry: %0 = atomicrmw and ptr %i, i32 1 syncscope("work_item") seq_cst, align 4 %1 = atomicrmw min ptr %i, i32 1 syncscope("all_svm_devices") seq_cst, align 4 @@ -101,7 +101,7 @@ entry: ; CHECK-SPIRV: AtomicCompareExchange [[#]] [[#]] [[#]] [[#ConstInt2]] [[#ConstInt2]] [[#ConstInt2]] [[#ConstInt1]] [[#ConstInt0]] ; CHECK-LLVM: call spir_func i1 @_Z39atomic_compare_exchange_strong_explicitPU3AS4VU7_AtomiciPU3AS4ii12memory_orderS4_12memory_scope(ptr{{.*}}, ptr{{.*}}, i32 1, i32 2, i32 2, i32 1) -define dso_local zeroext i1 @fi4(ptr nocapture noundef %i) local_unnamed_addr #0 { +define dso_local zeroext i1 @fi4(ptr captures(none) noundef %i) local_unnamed_addr #0 { entry: %0 = cmpxchg ptr %i, i32 0, i32 1 syncscope("workgroup") acquire acquire, align 4 %1 = extractvalue { i32, i1 } %0, 1 @@ -112,7 +112,7 @@ entry: ; CHECK-SPIRV: AtomicExchange [[#]] [[#]] [[#]] [[#ConstInt2]] [[#SequentiallyConsistent]] [[#Const2Power30]] ; CHECK-LLVM: call spir_func i32 @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicii12memory_order12memory_scope(ptr{{.*}}, i32 1073741824, i32 5, i32 1) -define dso_local float @ff3(ptr nocapture noundef %d) local_unnamed_addr #0 { +define dso_local float @ff3(ptr captures(none) noundef %d) local_unnamed_addr #0 { entry: %0 = atomicrmw xchg ptr %d, i32 1073741824 syncscope("workgroup") seq_cst, align 4 %1 = bitcast i32 %0 to float @@ -123,7 +123,7 @@ entry: ; CHECK-SPIRV: AtomicFAddEXT [[#]] [[#]] [[#]] [[#ConstInt2]] [[#ConstInt0]] [[#]] ; CHECK-LLVM: call spir_func float @_Z25atomic_fetch_add_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(ptr{{.*}}, i32 0, i32 1) -define dso_local float @ff4(ptr addrspace(1) nocapture noundef %d, float noundef %a) local_unnamed_addr #0 { +define dso_local float @ff4(ptr addrspace(1) captures(none) noundef %d, float noundef %a) local_unnamed_addr #0 { entry: %0 = atomicrmw fadd ptr addrspace(1) %d, float %a syncscope("workgroup") monotonic, align 4 ret float %0 diff --git a/test/transcoding/SampledImage.cl b/test/transcoding/SampledImage.cl index c37de46333..6a8d60a5b6 100644 --- a/test/transcoding/SampledImage.cl +++ b/test/transcoding/SampledImage.cl @@ -40,7 +40,7 @@ void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sam // CHECK-SPIRV: Function {{.*}} [[sample_kernel_float]] // CHECK-SPIRV: FunctionParameter {{.*}} [[InputImage:[0-9]+]] // CHECK-SPIRV: FunctionParameter [[TypeSampler]] [[argSampl:[0-9]+]] -// CHECK-LLVM: define spir_kernel void @sample_kernel_float(ptr addrspace(1) %input, <2 x float> %coords, ptr addrspace(1) nocapture align 16 %results, ptr addrspace(2) %argSampl) +// CHECK-LLVM: define spir_kernel void @sample_kernel_float(ptr addrspace(1) %input, <2 x float> %coords, ptr addrspace(1) align 16 captures(none) %results, ptr addrspace(2) %argSampl) // CHECK-SPIRV: SampledImage [[SampledImageTy]] [[SampledImage1:[0-9]+]] [[InputImage]] [[ConstSampler1]] // CHECK-SPIRV: ImageSampleExplicitLod {{[0-9]+}} {{[0-9]+}} [[SampledImage1]] @@ -63,7 +63,7 @@ void sample_kernel_int(image2d_t input, float2 coords, global int4 *results, sam // CHECK-SPIRV: Function {{.*}} [[sample_kernel_int]] // CHECK-SPIRV: FunctionParameter {{.*}} [[InputImage:[0-9]+]] // CHECK-SPIRV: FunctionParameter [[TypeSampler]] [[argSampl:[0-9]+]] -// CHECK-LLVM: define spir_kernel void @sample_kernel_int(ptr addrspace(1) %input, <2 x float> %coords, ptr addrspace(1) nocapture align 16 %results, ptr addrspace(2) %argSampl) +// CHECK-LLVM: define spir_kernel void @sample_kernel_int(ptr addrspace(1) %input, <2 x float> %coords, ptr addrspace(1) align 16 captures(none) %results, ptr addrspace(2) %argSampl) // CHECK-SPIRV: SampledImage [[SampledImageTy]] [[SampledImage4:[0-9]+]] [[InputImage]] [[ConstSampler3]] // CHECK-SPIRV: ImageSampleExplicitLod {{[0-9]+}} {{[0-9]+}} [[SampledImage4]] diff --git a/test/transcoding/SpecConstantComposite.ll b/test/transcoding/SpecConstantComposite.ll index c2226dec11..ee51dcecd9 100644 --- a/test/transcoding/SpecConstantComposite.ll +++ b/test/transcoding/SpecConstantComposite.ll @@ -90,13 +90,13 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind willreturn -declare void @llvm.memcpy.p4.p0.i64(ptr addrspace(4) noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #1 +declare void @llvm.memcpy.p4.p0.i64(ptr addrspace(4) noalias captures(none) writeonly, ptr noalias captures(none) readonly, i64, i1 immarg) #1 declare i32 @_Z20__spirv_SpecConstantii(i32, i32) diff --git a/test/transcoding/annotate_attribute.ll b/test/transcoding/annotate_attribute.ll index acf50e32a3..e4a620b02d 100644 --- a/test/transcoding/annotate_attribute.ll +++ b/test/transcoding/annotate_attribute.ll @@ -85,7 +85,7 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0(i64, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64, ptr captures(none)) #1 ; Function Attrs: inlinehint nounwind define internal spir_func void @"_ZZ4mainENK3$_0clEv"(ptr %this) #2 align 2 { @@ -99,7 +99,7 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0(i64, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64, ptr captures(none)) #1 ; Function Attrs: nounwind define spir_func void @_Z3foov() #3 { diff --git a/test/transcoding/annotation_dbg_info_drop.ll b/test/transcoding/annotation_dbg_info_drop.ll index 19ba52269f..2390875cd9 100644 --- a/test/transcoding/annotation_dbg_info_drop.ll +++ b/test/transcoding/annotation_dbg_info_drop.ll @@ -156,10 +156,10 @@ entry: declare void @llvm.dbg.declare(metadata, metadata, metadata) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #2 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #2 ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #2 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #2 ; Function Attrs: nounwind declare void @llvm.var.annotation(ptr, ptr, ptr, i32, ptr) #5 @@ -174,7 +174,7 @@ declare i64 @llvm.annotation.i64(i64, ptr, ptr, i32) #5 declare ptr @llvm.ptr.annotation.p0(ptr, ptr, ptr, i32, ptr) #5 ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p0.p0.i64(ptr nocapture writeonly, ptr nocapture readonly, i64, i1 immarg) #2 +declare void @llvm.memcpy.p0.p0.i64(ptr captures(none) writeonly, ptr captures(none) readonly, i64, i1 immarg) #2 attributes #0 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { nounwind readnone speculatable } diff --git a/test/transcoding/bitcast.ll b/test/transcoding/bitcast.ll index 3161852abe..3660e92310 100644 --- a/test/transcoding/bitcast.ll +++ b/test/transcoding/bitcast.ll @@ -13,7 +13,7 @@ target triple = "spir64-unknown-unknown" ; CHECK: bitcast ; Function Attrs: nounwind -define spir_kernel void @test_fn(ptr addrspace(1) nocapture readonly %src, ptr addrspace(1) nocapture %dst) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @test_fn(ptr addrspace(1) captures(none) readonly %src, ptr addrspace(1) captures(none) %dst) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { entry: %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 %sext = shl i64 %call, 32 diff --git a/test/transcoding/builtin_function_readnone_attr.ll b/test/transcoding/builtin_function_readnone_attr.ll index a8bedda5d2..f7cd4af46f 100644 --- a/test/transcoding/builtin_function_readnone_attr.ll +++ b/test/transcoding/builtin_function_readnone_attr.ll @@ -11,13 +11,13 @@ ; CHECK-SPIRV: Decorate [[#A]] FuncParamAttr 6 ; CHECK-SPIRV: Decorate [[#B]] FuncParamAttr 7 -; CHECK-LLVM: {{.*}}void @test_builtin_readnone(ptr nocapture readonly %{{.*}}, ptr nocapture readnone %{{.*}}) +; CHECK-LLVM: {{.*}}void @test_builtin_readnone(ptr readonly captures(none) %{{.*}}, ptr readnone captures(none) %{{.*}}) target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128" target triple = "spir-unknown-unknown" ; Function Attrs: convergent nofree norecurse nounwind uwtable -define dso_local spir_kernel void @test_builtin_readnone(ptr nocapture readonly %a, ptr nocapture readnone %b) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @test_builtin_readnone(ptr captures(none) readonly %a, ptr captures(none) readnone %b) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { entry: %0 = load double, ptr %a, align 8, !tbaa !7 %call = tail call double @_Z3expd(double %0) #2 diff --git a/test/transcoding/fclamp.ll b/test/transcoding/fclamp.ll index 3b245662ea..f5ca8ec34b 100644 --- a/test/transcoding/fclamp.ll +++ b/test/transcoding/fclamp.ll @@ -16,7 +16,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir-unknown-unknown" ; Function Attrs: nounwind -define spir_kernel void @test_scalar(ptr addrspace(1) nocapture readonly %f) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @test_scalar(ptr addrspace(1) captures(none) readonly %f) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: %0 = load float, ptr addrspace(1) %f, align 4 %call = tail call spir_func float @_Z5clampfff(float %0, float 0.000000e+00, float 1.000000e+00) #2 diff --git a/test/transcoding/image_builtins.ll b/test/transcoding/image_builtins.ll index 9eda26e658..d763d7ce20 100644 --- a/test/transcoding/image_builtins.ll +++ b/test/transcoding/image_builtins.ll @@ -22,7 +22,7 @@ target triple = "spir-unknown-unknown" ; CHECK-LLVM: call spir_func void @_Z12write_imageh14ocl_image2d_woDv2_iDv4_Dh( ; Function Attrs: convergent nounwind -define spir_kernel void @nosamp(ptr addrspace(1) %im, <2 x i32> %coord, ptr addrspace(1) nocapture %res) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define spir_kernel void @nosamp(ptr addrspace(1) %im, <2 x i32> %coord, ptr addrspace(1) captures(none) %res) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { entry: %call = tail call spir_func <4 x half> @_Z11read_imageh14ocl_image2d_roDv2_i(ptr addrspace(1) %im, <2 x i32> %coord) #3 store <4 x half> %call, ptr addrspace(1) %res, align 8, !tbaa !8 @@ -33,7 +33,7 @@ entry: declare spir_func <4 x half> @_Z11read_imageh14ocl_image2d_roDv2_i(ptr addrspace(1), <2 x i32>) local_unnamed_addr #1 ; Function Attrs: convergent nounwind -define spir_kernel void @withsamp(ptr addrspace(1) %im, ptr addrspace(2) %smp, <2 x i32> %coord, ptr addrspace(1) nocapture %res) local_unnamed_addr #0 !kernel_arg_addr_space !11 !kernel_arg_access_qual !12 !kernel_arg_type !13 !kernel_arg_base_type !14 !kernel_arg_type_qual !15 { +define spir_kernel void @withsamp(ptr addrspace(1) %im, ptr addrspace(2) %smp, <2 x i32> %coord, ptr addrspace(1) captures(none) %res) local_unnamed_addr #0 !kernel_arg_addr_space !11 !kernel_arg_access_qual !12 !kernel_arg_type !13 !kernel_arg_base_type !14 !kernel_arg_type_qual !15 { entry: %call = tail call spir_func <4 x half> @_Z11read_imageh14ocl_image2d_ro11ocl_samplerDv2_i(ptr addrspace(1) %im, ptr addrspace(2) %smp, <2 x i32> %coord) #3 store <4 x half> %call, ptr addrspace(1) %res, align 8, !tbaa !8 @@ -44,7 +44,7 @@ entry: declare spir_func <4 x half> @_Z11read_imageh14ocl_image2d_ro11ocl_samplerDv2_i(ptr addrspace(1), ptr addrspace(2), <2 x i32>) local_unnamed_addr #1 ; Function Attrs: convergent nounwind -define spir_kernel void @writehalf(ptr addrspace(1) %im, <2 x i32> %coord, ptr addrspace(1) nocapture readonly %val) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !16 !kernel_arg_type !5 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define spir_kernel void @writehalf(ptr addrspace(1) %im, <2 x i32> %coord, ptr addrspace(1) captures(none) readonly %val) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !16 !kernel_arg_type !5 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { entry: %0 = load <4 x half>, ptr addrspace(1) %val, align 8, !tbaa !8 tail call spir_func void @_Z12write_imageh14ocl_image2d_woDv2_iDv4_Dh(ptr addrspace(1) %im, <2 x i32> %coord, <4 x half> %0) #4 diff --git a/test/transcoding/image_channel.ll b/test/transcoding/image_channel.ll index d6d49a6b0b..bacba71a02 100644 --- a/test/transcoding/image_channel.ll +++ b/test/transcoding/image_channel.ll @@ -19,7 +19,7 @@ target triple = "spir-unknown-unknown" ; CHECK-SPIRV-DAG: 4 Constant {{[0-9]+}} [[OrderOffsetId:[0-9]+]] 4272 ; Function Attrs: nounwind -define spir_kernel void @f(ptr addrspace(1) %img, ptr addrspace(1) nocapture %type, ptr addrspace(1) nocapture %order) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define spir_kernel void @f(ptr addrspace(1) %img, ptr addrspace(1) captures(none) %type, ptr addrspace(1) captures(none) %order) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { ; CHECK-LLVM-DAG: [[DTCALL:%.+]] ={{.*}} call spir_func i32 @_Z27get_image_channel_data_type14ocl_image2d_ro( ; CHECK-LLVM: [[DTSUB:%.+]] = sub i32 [[DTCALL]], 4304 ; CHECK-LLVM: [[DTADD:%.+]] = add i32 [[DTSUB]], 4304 diff --git a/test/transcoding/image_signedness.ll b/test/transcoding/image_signedness.ll index 4caa213cf0..8298596981 100644 --- a/test/transcoding/image_signedness.ll +++ b/test/transcoding/image_signedness.ll @@ -30,7 +30,7 @@ target triple = "spir-unknown-unknown" ; CHECK-SPV-IR: call spir_func <4 x i32> @_Z23__spirv_ImageRead_Rint4PU3AS133__spirv_Image__void_0_0_0_0_0_0_0ii( ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @imagereads(ptr addrspace(1) %im, ptr addrspace(1) %ima, ptr addrspace(1) nocapture %res, ptr addrspace(1) nocapture %resu) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !7 !kernel_arg_type_qual !8 { +define dso_local spir_kernel void @imagereads(ptr addrspace(1) %im, ptr addrspace(1) %ima, ptr addrspace(1) captures(none) %res, ptr addrspace(1) captures(none) %resu) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !7 !kernel_arg_type_qual !8 { entry: %0 = tail call ptr addrspace(2) @__translate_sampler_initializer(i32 19) #2 %call = tail call spir_func <4 x i32> @_Z12read_imageui14ocl_image1d_ro11ocl_sampleri(ptr addrspace(1) %im, ptr addrspace(2) %0, i32 42) #3 @@ -53,7 +53,7 @@ entry: ; CHECK-SPV-IR: call spir_func void @_Z18__spirv_ImageWritePU3AS133__spirv_Image__void_1_0_0_0_0_0_1Dv2_iDv4_ji( ; Function Attrs: alwaysinline convergent nounwind -define spir_kernel void @imagewrites(i32 %offset, ptr addrspace(1) nocapture readonly %input, ptr addrspace(1) nocapture readonly %inputu, ptr addrspace(1) %output) local_unnamed_addr #0 !kernel_arg_addr_space !14 !kernel_arg_access_qual !15 !kernel_arg_type !16 !kernel_arg_base_type !17 !kernel_arg_type_qual !18 !kernel_arg_name !19 !kernel_attributes !20 { +define spir_kernel void @imagewrites(i32 %offset, ptr addrspace(1) captures(none) readonly %input, ptr addrspace(1) captures(none) readonly %inputu, ptr addrspace(1) %output) local_unnamed_addr #0 !kernel_arg_addr_space !14 !kernel_arg_access_qual !15 !kernel_arg_type !16 !kernel_arg_base_type !17 !kernel_arg_type_qual !18 !kernel_arg_name !19 !kernel_attributes !20 { entry: %idxprom = sext i32 %offset to i64 %arrayidx = getelementptr inbounds <4 x i32>, ptr addrspace(1) %input, i64 %idxprom diff --git a/test/transcoding/image_with_suffix.ll b/test/transcoding/image_with_suffix.ll index 91f5bdf31c..c847698dd1 100644 --- a/test/transcoding/image_with_suffix.ll +++ b/test/transcoding/image_with_suffix.ll @@ -16,7 +16,7 @@ target triple = "spir-unknown-unknown" %opencl.image1d_ro_t.2 = type opaque ; Function Attrs: convergent nounwind writeonly -define spir_kernel void @foo(ptr addrspace(1) %im, ptr addrspace(1) nocapture %res) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define spir_kernel void @foo(ptr addrspace(1) %im, ptr addrspace(1) captures(none) %res) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { entry: %call = tail call spir_func i32 @_Z15get_image_width14ocl_image1d_ro(ptr addrspace(1) %im) #2 store i32 %call, ptr addrspace(1) %res, align 4, !tbaa !8 diff --git a/test/transcoding/intrinsic_result_store.ll b/test/transcoding/intrinsic_result_store.ll index f542e3ecb5..c27aae6151 100644 --- a/test/transcoding/intrinsic_result_store.ll +++ b/test/transcoding/intrinsic_result_store.ll @@ -25,7 +25,7 @@ entry: } ; Function Attrs: argmemonly nounwind willreturn writeonly -declare void @llvm.memset.p1.i64(ptr addrspace(1) nocapture writeonly, i8, i64, i1 immarg) #1 +declare void @llvm.memset.p1.i64(ptr addrspace(1) captures(none) writeonly, i8, i64, i1 immarg) #1 attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { argmemonly nounwind willreturn writeonly } diff --git a/test/transcoding/isequal.ll b/test/transcoding/isequal.ll index df642825c3..aace1cf2b7 100644 --- a/test/transcoding/isequal.ll +++ b/test/transcoding/isequal.ll @@ -15,7 +15,7 @@ target triple = "spir64-unknown-unknown" ; CHECK-LLVM-NOT: call {{.*}} @_Z{{[0-9]*}}convert ; Function Attrs: nounwind -define spir_kernel void @math_kernel8(ptr addrspace(1) nocapture %out, ptr addrspace(1) nocapture readonly %in1, ptr addrspace(1) nocapture readonly %in2) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @math_kernel8(ptr addrspace(1) captures(none) %out, ptr addrspace(1) captures(none) readonly %in1, ptr addrspace(1) captures(none) readonly %in2) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { entry: %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2 %sext = shl i64 %call, 32 diff --git a/test/transcoding/relationals_double.ll b/test/transcoding/relationals_double.ll index e7da9c7588..507210c75f 100644 --- a/test/transcoding/relationals_double.ll +++ b/test/transcoding/relationals_double.ll @@ -110,7 +110,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir" ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_scalar(ptr addrspace(4) nocapture writeonly %out, double %d) local_unnamed_addr #0 { +define dso_local spir_func void @test_scalar(ptr addrspace(4) captures(none) writeonly %out, double %d) local_unnamed_addr #0 { entry: %call = tail call spir_func i32 @_Z8isfinited(double %d) #3 %call1 = tail call spir_func i32 @_Z5isinfd(double %d) #3 @@ -186,7 +186,7 @@ declare spir_func i32 @_Z9isordereddd(double, double) local_unnamed_addr #1 declare spir_func i32 @_Z11isunordereddd(double, double) local_unnamed_addr #1 ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_vector(ptr addrspace(4) nocapture writeonly %out, <2 x double> %d) local_unnamed_addr #2 { +define dso_local spir_func void @test_vector(ptr addrspace(4) captures(none) writeonly %out, <2 x double> %d) local_unnamed_addr #2 { entry: %call = tail call spir_func <2 x i64> @_Z8isfiniteDv2_d(<2 x double> %d) #3 %call1 = tail call spir_func <2 x i64> @_Z5isinfDv2_d(<2 x double> %d) #3 diff --git a/test/transcoding/relationals_float.ll b/test/transcoding/relationals_float.ll index 067ef865fa..f62b6c1b27 100644 --- a/test/transcoding/relationals_float.ll +++ b/test/transcoding/relationals_float.ll @@ -110,7 +110,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir" ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_scalar(ptr addrspace(4) nocapture writeonly %out, float %f) local_unnamed_addr #0 { +define dso_local spir_func void @test_scalar(ptr addrspace(4) captures(none) writeonly %out, float %f) local_unnamed_addr #0 { entry: %call = tail call spir_func i32 @_Z8isfinitef(float %f) #3 %call1 = tail call spir_func i32 @_Z5isinff(float %f) #3 @@ -186,7 +186,7 @@ declare spir_func i32 @_Z9isorderedff(float, float) local_unnamed_addr #1 declare spir_func i32 @_Z11isunorderedff(float, float) local_unnamed_addr #1 ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_vector(ptr addrspace(4) nocapture writeonly %out, <2 x float> %f) local_unnamed_addr #2 { +define dso_local spir_func void @test_vector(ptr addrspace(4) captures(none) writeonly %out, <2 x float> %f) local_unnamed_addr #2 { entry: %call = tail call spir_func <2 x i32> @_Z8isfiniteDv2_f(<2 x float> %f) #3 %call1 = tail call spir_func <2 x i32> @_Z5isinfDv2_f(<2 x float> %f) #3 diff --git a/test/transcoding/relationals_half.ll b/test/transcoding/relationals_half.ll index 33cdf46c67..5b0ce7ba44 100644 --- a/test/transcoding/relationals_half.ll +++ b/test/transcoding/relationals_half.ll @@ -109,7 +109,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:2 target triple = "spir" ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_scalar(ptr addrspace(4) nocapture writeonly %out, half %h) local_unnamed_addr #0 { +define dso_local spir_func void @test_scalar(ptr addrspace(4) captures(none) writeonly %out, half %h) local_unnamed_addr #0 { entry: %call = tail call spir_func i32 @_Z8isfiniteDh(half %h) #3 %call1 = tail call spir_func i32 @_Z5isinfDh(half %h) #3 @@ -185,7 +185,7 @@ declare spir_func i32 @_Z9isorderedDhDh(half, half) local_unnamed_addr #1 declare spir_func i32 @_Z11isunorderedDhDh(half, half) local_unnamed_addr #1 ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn writeonly -define dso_local spir_func void @test_vector(ptr addrspace(4) nocapture writeonly %out, <2 x half> %h) local_unnamed_addr #2 { +define dso_local spir_func void @test_vector(ptr addrspace(4) captures(none) writeonly %out, <2 x half> %h) local_unnamed_addr #2 { entry: %call = tail call spir_func <2 x i16> @_Z8isfiniteDv2_Dh(<2 x half> %h) #3 %call1 = tail call spir_func <2 x i16> @_Z5isinfDv2_Dh(<2 x half> %h) #3 diff --git a/test/transcoding/relationals_select.ll b/test/transcoding/relationals_select.ll index b7a4025f3a..bc40dd1eb8 100644 --- a/test/transcoding/relationals_select.ll +++ b/test/transcoding/relationals_select.ll @@ -39,7 +39,7 @@ declare spir_func <4 x i32> @_Z5isinfDv4_f(<4 x float> noundef) local_unnamed_ad declare spir_func <4 x i32> @_Z7signbitDv4_f(<4 x float> noundef) local_unnamed_addr #1 ; Function Attrs: convergent mustprogress nofree norecurse nounwind willreturn memory(argmem: readwrite) -define dso_local spir_kernel void @math_kernel_scalar(ptr addrspace(4) nocapture writeonly %out, float %f) local_unnamed_addr #0 { +define dso_local spir_kernel void @math_kernel_scalar(ptr addrspace(4) captures(none) writeonly %out, float %f) local_unnamed_addr #0 { entry: ; CHECK: [[DATA0:%.*]] = call spir_func i32 @_Z8isfinitef(float [[ARG0:%.*]]) ; CHECK-NEXT: [[DATA1:%.*]] = trunc i32 [[DATA0]] to i1 diff --git a/test/transcoding/spec_const.ll b/test/transcoding/spec_const.ll index d167f3f7e6..6c0c9f1f13 100644 --- a/test/transcoding/spec_const.ll +++ b/test/transcoding/spec_const.ll @@ -32,7 +32,7 @@ target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" target triple = "spir" ; Function Attrs: nofree norecurse nounwind writeonly - define spir_kernel void @foo(ptr addrspace(1) nocapture %b, ptr addrspace(1) nocapture %c, ptr addrspace(1) nocapture %s, ptr addrspace(1) nocapture %i, ptr addrspace(1) nocapture %l, ptr addrspace(1) nocapture %h, ptr addrspace(1) nocapture %f, ptr addrspace(1) nocapture %d) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { + define spir_kernel void @foo(ptr addrspace(1) captures(none) %b, ptr addrspace(1) captures(none) %c, ptr addrspace(1) captures(none) %s, ptr addrspace(1) captures(none) %i, ptr addrspace(1) captures(none) %l, ptr addrspace(1) captures(none) %h, ptr addrspace(1) captures(none) %f, ptr addrspace(1) captures(none) %d) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { entry: ; CHECK-LLVM: store i8 0, ptr addrspace(1) %b, align 1 ; CHECK-LLVM-SPEC: store i8 1, ptr addrspace(1) %b, align 1 diff --git a/test/transcoding/spirv-private-array-initialization.ll b/test/transcoding/spirv-private-array-initialization.ll index 6de1ea4f3b..c06059d900 100644 --- a/test/transcoding/spirv-private-array-initialization.ll +++ b/test/transcoding/spirv-private-array-initialization.ll @@ -51,7 +51,7 @@ entry: } ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p0.p2.i32(ptr nocapture writeonly, ptr addrspace(2) nocapture readonly, i32, i1) #1 +declare void @llvm.memcpy.p0.p2.i32(ptr captures(none) writeonly, ptr addrspace(2) captures(none) readonly, i32, i1) #1 attributes #0 = { convergent noinline nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { argmemonly nounwind } diff --git a/test/transcoding/sub_group_ballot.ll b/test/transcoding/sub_group_ballot.ll index c5a1f9a98d..4b6ee8e0b0 100644 --- a/test/transcoding/sub_group_ballot.ll +++ b/test/transcoding/sub_group_ballot.ll @@ -1176,7 +1176,7 @@ declare dso_local spir_func double @_Z25sub_group_broadcast_firstd(double) local ; CHECK-SPV-IR: call spir_func i32 @_Z36__spirv_GroupNonUniformBallotFindMSBiDv4_j(i32 3, <4 x i32> %[[ballot]]) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testBallotOperations(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @testBallotOperations(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { %2 = tail call spir_func <4 x i32> @_Z16sub_group_balloti(i32 0) #7 %3 = tail call spir_func i32 @_Z24sub_group_inverse_ballotDv4_j(<4 x i32> %2) #8 store i32 %3, ptr addrspace(1) %0, align 4, !tbaa !8 @@ -1248,7 +1248,7 @@ declare dso_local spir_func i32 @_Z25sub_group_ballot_find_msbDv4_j(<4 x i32>) l ; CHECK-SPV-IR: call spir_func <4 x i32> @_Z32__spirv_BuiltInSubgroupLtMaskKHRv() ; Function Attrs: convergent nofree nounwind writeonly -define dso_local spir_kernel void @testSubgroupMasks(ptr addrspace(1) nocapture) local_unnamed_addr #6 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !12 !kernel_arg_base_type !13 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @testSubgroupMasks(ptr addrspace(1) captures(none)) local_unnamed_addr #6 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !12 !kernel_arg_base_type !13 !kernel_arg_type_qual !7 { %2 = tail call spir_func <4 x i32> @_Z21get_sub_group_eq_maskv() #8 store <4 x i32> %2, ptr addrspace(1) %0, align 16, !tbaa !14 %3 = tail call spir_func <4 x i32> @_Z21get_sub_group_ge_maskv() #8 diff --git a/test/transcoding/sub_group_clustered_reduce.ll b/test/transcoding/sub_group_clustered_reduce.ll index dee43e6132..c4db0e0282 100644 --- a/test/transcoding/sub_group_clustered_reduce.ll +++ b/test/transcoding/sub_group_clustered_reduce.ll @@ -230,7 +230,7 @@ target triple = "spir64" ; CHECK-SPV-IR: call spir_func i8 @_Z27__spirv_GroupNonUniformSMaxiicj(i32 3, i32 3, i8 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i8 @_Z30sub_group_clustered_reduce_addcj(i8 signext 0, i32 2) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func signext i8 @_Z30sub_group_clustered_reduce_mulcj(i8 signext 0, i32 2) #2 @@ -277,7 +277,7 @@ declare dso_local spir_func signext i8 @_Z30sub_group_clustered_reduce_maxcj(i8 ; CHECK-SPV-IR: call spir_func i8 @_Z27__spirv_GroupNonUniformUMaxiihj(i32 3, i32 3, i8 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticUChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i8 @_Z30sub_group_clustered_reduce_addhj(i8 zeroext 0, i32 2) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func zeroext i8 @_Z30sub_group_clustered_reduce_mulhj(i8 zeroext 0, i32 2) #2 @@ -324,7 +324,7 @@ declare dso_local spir_func zeroext i8 @_Z30sub_group_clustered_reduce_maxhj(i8 ; CHECK-SPV-IR: call spir_func i16 @_Z27__spirv_GroupNonUniformSMaxiisj(i32 3, i32 3, i16 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i16 @_Z30sub_group_clustered_reduce_addsj(i16 signext 0, i32 2) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func signext i16 @_Z30sub_group_clustered_reduce_mulsj(i16 signext 0, i32 2) #2 @@ -371,7 +371,7 @@ declare dso_local spir_func signext i16 @_Z30sub_group_clustered_reduce_maxsj(i1 ; CHECK-SPV-IR: call spir_func i16 @_Z27__spirv_GroupNonUniformUMaxiitj(i32 3, i32 3, i16 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticUShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i16 @_Z30sub_group_clustered_reduce_addtj(i16 zeroext 0, i32 2) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func zeroext i16 @_Z30sub_group_clustered_reduce_multj(i16 zeroext 0, i32 2) #2 @@ -418,7 +418,7 @@ declare dso_local spir_func zeroext i16 @_Z30sub_group_clustered_reduce_maxtj(i1 ; CHECK-SPV-IR: call spir_func i32 @_Z27__spirv_GroupNonUniformSMaxiiij(i32 3, i32 3, i32 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z30sub_group_clustered_reduce_addij(i32 0, i32 2) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z30sub_group_clustered_reduce_mulij(i32 0, i32 2) #2 @@ -465,7 +465,7 @@ declare dso_local spir_func i32 @_Z30sub_group_clustered_reduce_maxij(i32, i32) ; CHECK-SPV-IR: call spir_func i32 @_Z27__spirv_GroupNonUniformUMaxiijj(i32 3, i32 3, i32 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticUInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticUInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z30sub_group_clustered_reduce_addjj(i32 0, i32 2) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z30sub_group_clustered_reduce_muljj(i32 0, i32 2) #2 @@ -512,7 +512,7 @@ declare dso_local spir_func i32 @_Z30sub_group_clustered_reduce_maxjj(i32, i32) ; CHECK-SPV-IR: call spir_func i64 @_Z27__spirv_GroupNonUniformSMaxiilj(i32 3, i32 3, i64 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticLong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticLong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z30sub_group_clustered_reduce_addlj(i64 0, i32 2) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z30sub_group_clustered_reduce_mullj(i64 0, i32 2) #2 @@ -559,7 +559,7 @@ declare dso_local spir_func i64 @_Z30sub_group_clustered_reduce_maxlj(i64, i32) ; CHECK-SPV-IR: call spir_func i64 @_Z27__spirv_GroupNonUniformUMaxiimj(i32 3, i32 3, i64 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticULong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticULong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z30sub_group_clustered_reduce_addmj(i64 0, i32 2) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z30sub_group_clustered_reduce_mulmj(i64 0, i32 2) #2 @@ -606,7 +606,7 @@ declare dso_local spir_func i64 @_Z30sub_group_clustered_reduce_maxmj(i64, i32) ; CHECK-SPV-IR: call spir_func float @_Z27__spirv_GroupNonUniformFMaxiifj(i32 3, i32 3, float 0.000000e+00, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticFloat(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticFloat(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { %2 = tail call spir_func float @_Z30sub_group_clustered_reduce_addfj(float 0.000000e+00, i32 2) #2 store float %2, ptr addrspace(1) %0, align 4, !tbaa !24 %3 = tail call spir_func float @_Z30sub_group_clustered_reduce_mulfj(float 0.000000e+00, i32 2) #2 @@ -653,7 +653,7 @@ declare dso_local spir_func float @_Z30sub_group_clustered_reduce_maxfj(float, i ; CHECK-SPV-IR: call spir_func half @_Z27__spirv_GroupNonUniformFMaxiiDhj(i32 3, i32 3, half 0xH0000, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticHalf(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticHalf(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { %2 = tail call spir_func half @_Z30sub_group_clustered_reduce_addDhj(half 0xH0000, i32 2) #2 store half %2, ptr addrspace(1) %0, align 2, !tbaa !27 %3 = tail call spir_func half @_Z30sub_group_clustered_reduce_mulDhj(half 0xH0000, i32 2) #2 @@ -700,7 +700,7 @@ declare dso_local spir_func half @_Z30sub_group_clustered_reduce_maxDhj(half, i3 ; CHECK-SPV-IR: call spir_func double @_Z27__spirv_GroupNonUniformFMaxiidj(i32 3, i32 3, double 0.000000e+00, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredArithmeticDouble(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredArithmeticDouble(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { %2 = tail call spir_func double @_Z30sub_group_clustered_reduce_adddj(double 0.000000e+00, i32 2) #2 store double %2, ptr addrspace(1) %0, align 8, !tbaa !30 %3 = tail call spir_func double @_Z30sub_group_clustered_reduce_muldj(double 0.000000e+00, i32 2) #2 @@ -744,7 +744,7 @@ declare dso_local spir_func double @_Z30sub_group_clustered_reduce_maxdj(double, ; CHECK-SPV-IR: call spir_func i8 @_Z33__spirv_GroupNonUniformBitwiseXoriicj(i32 3, i32 3, i8 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i8 @_Z30sub_group_clustered_reduce_andcj(i8 signext 0, i32 2) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func signext i8 @_Z29sub_group_clustered_reduce_orcj(i8 signext 0, i32 2) #2 @@ -782,7 +782,7 @@ declare dso_local spir_func signext i8 @_Z30sub_group_clustered_reduce_xorcj(i8 ; CHECK-SPV-IR: call spir_func i8 @_Z33__spirv_GroupNonUniformBitwiseXoriicj(i32 3, i32 3, i8 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseUChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i8 @_Z30sub_group_clustered_reduce_andhj(i8 zeroext 0, i32 2) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func zeroext i8 @_Z29sub_group_clustered_reduce_orhj(i8 zeroext 0, i32 2) #2 @@ -820,7 +820,7 @@ declare dso_local spir_func zeroext i8 @_Z30sub_group_clustered_reduce_xorhj(i8 ; CHECK-SPV-IR: call spir_func i16 @_Z33__spirv_GroupNonUniformBitwiseXoriisj(i32 3, i32 3, i16 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i16 @_Z30sub_group_clustered_reduce_andsj(i16 signext 0, i32 2) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func signext i16 @_Z29sub_group_clustered_reduce_orsj(i16 signext 0, i32 2) #2 @@ -858,7 +858,7 @@ declare dso_local spir_func signext i16 @_Z30sub_group_clustered_reduce_xorsj(i1 ; CHECK-SPV-IR: call spir_func i16 @_Z33__spirv_GroupNonUniformBitwiseXoriisj(i32 3, i32 3, i16 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseUShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i16 @_Z30sub_group_clustered_reduce_andtj(i16 zeroext 0, i32 2) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func zeroext i16 @_Z29sub_group_clustered_reduce_ortj(i16 zeroext 0, i32 2) #2 @@ -896,7 +896,7 @@ declare dso_local spir_func zeroext i16 @_Z30sub_group_clustered_reduce_xortj(i1 ; CHECK-SPV-IR: call spir_func i32 @_Z33__spirv_GroupNonUniformBitwiseXoriiij(i32 3, i32 3, i32 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z30sub_group_clustered_reduce_andij(i32 0, i32 2) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z29sub_group_clustered_reduce_orij(i32 0, i32 2) #2 @@ -934,7 +934,7 @@ declare dso_local spir_func i32 @_Z30sub_group_clustered_reduce_xorij(i32, i32) ; CHECK-SPV-IR: call spir_func i32 @_Z33__spirv_GroupNonUniformBitwiseXoriiij(i32 3, i32 3, i32 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseUInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseUInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z30sub_group_clustered_reduce_andjj(i32 0, i32 2) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z29sub_group_clustered_reduce_orjj(i32 0, i32 2) #2 @@ -972,7 +972,7 @@ declare dso_local spir_func i32 @_Z30sub_group_clustered_reduce_xorjj(i32, i32) ; CHECK-SPV-IR: call spir_func i64 @_Z33__spirv_GroupNonUniformBitwiseXoriilj(i32 3, i32 3, i64 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseLong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseLong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z30sub_group_clustered_reduce_andlj(i64 0, i32 2) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z29sub_group_clustered_reduce_orlj(i64 0, i32 2) #2 @@ -1010,7 +1010,7 @@ declare dso_local spir_func i64 @_Z30sub_group_clustered_reduce_xorlj(i64, i32) ; CHECK-SPV-IR: call spir_func i64 @_Z33__spirv_GroupNonUniformBitwiseXoriilj(i32 3, i32 3, i64 0, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredBitwiseULong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredBitwiseULong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z30sub_group_clustered_reduce_andmj(i64 0, i32 2) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z29sub_group_clustered_reduce_ormj(i64 0, i32 2) #2 @@ -1048,7 +1048,7 @@ declare dso_local spir_func i64 @_Z30sub_group_clustered_reduce_xormj(i64, i32) ; CHECK-SPV-IR: call spir_func i1 @_Z33__spirv_GroupNonUniformLogicalXoriibj(i32 3, i32 3, i1 false, i32 2) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testClusteredLogical(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testClusteredLogical(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z38sub_group_clustered_reduce_logical_andij(i32 0, i32 2) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z37sub_group_clustered_reduce_logical_orij(i32 0, i32 2) #2 diff --git a/test/transcoding/sub_group_extended_types.ll b/test/transcoding/sub_group_extended_types.ll index e6f65d9a34..86bf3df17e 100644 --- a/test/transcoding/sub_group_extended_types.ll +++ b/test/transcoding/sub_group_extended_types.ll @@ -990,7 +990,7 @@ declare dso_local spir_func <16 x double> @_Z19sub_group_broadcastDv16_dj(<16 x ; CHECK-LLVM: call spir_func i8 @_Z28sub_group_scan_exclusive_maxc(i8 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testReduceScanChar(ptr addrspace(1) nocapture) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @testReduceScanChar(ptr addrspace(1) captures(none)) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { %2 = tail call spir_func signext i8 @_Z20sub_group_reduce_addc(i8 signext 0) #6 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !8 %3 = tail call spir_func signext i8 @_Z20sub_group_reduce_minc(i8 signext 0) #6 @@ -1071,7 +1071,7 @@ declare dso_local spir_func signext i8 @_Z28sub_group_scan_exclusive_maxc(i8 sig ; CHECK-LLVM: call spir_func i8 @_Z28sub_group_scan_exclusive_maxh(i8 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testReduceScanUChar(ptr addrspace(1) nocapture) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @testReduceScanUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !7 { %2 = tail call spir_func zeroext i8 @_Z20sub_group_reduce_addh(i8 zeroext 0) #6 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !8 %3 = tail call spir_func zeroext i8 @_Z20sub_group_reduce_minh(i8 zeroext 0) #6 @@ -1152,7 +1152,7 @@ declare dso_local spir_func zeroext i8 @_Z28sub_group_scan_exclusive_maxh(i8 zer ; CHECK-LLVM: call spir_func i16 @_Z28sub_group_scan_exclusive_maxs(i16 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testReduceScanShort(ptr addrspace(1) nocapture) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !12 !kernel_arg_base_type !12 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @testReduceScanShort(ptr addrspace(1) captures(none)) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !12 !kernel_arg_base_type !12 !kernel_arg_type_qual !7 { %2 = tail call spir_func signext i16 @_Z20sub_group_reduce_adds(i16 signext 0) #6 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !13 %3 = tail call spir_func signext i16 @_Z20sub_group_reduce_mins(i16 signext 0) #6 @@ -1233,7 +1233,7 @@ declare dso_local spir_func signext i16 @_Z28sub_group_scan_exclusive_maxs(i16 s ; CHECK-LLVM: call spir_func i16 @_Z28sub_group_scan_exclusive_maxt(i16 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testReduceScanUShort(ptr addrspace(1) nocapture) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !7 { +define dso_local spir_kernel void @testReduceScanUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #5 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !7 { %2 = tail call spir_func zeroext i16 @_Z20sub_group_reduce_addt(i16 zeroext 0) #6 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !13 %3 = tail call spir_func zeroext i16 @_Z20sub_group_reduce_mint(i16 zeroext 0) #6 diff --git a/test/transcoding/sub_group_non_uniform_arithmetic.ll b/test/transcoding/sub_group_non_uniform_arithmetic.ll index c4c3e8e9c1..c3ff5d8b6b 100644 --- a/test/transcoding/sub_group_non_uniform_arithmetic.ll +++ b/test/transcoding/sub_group_non_uniform_arithmetic.ll @@ -388,7 +388,7 @@ target triple = "spir64" ; CHECK-SPV-IR: call spir_func i8 @_Z27__spirv_GroupNonUniformSMaxiic(i32 3, i32 2, i8 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i8 @_Z32sub_group_non_uniform_reduce_addc(i8 signext 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func signext i8 @_Z32sub_group_non_uniform_reduce_mulc(i8 signext 0) #2 @@ -507,7 +507,7 @@ declare dso_local spir_func signext i8 @_Z40sub_group_non_uniform_scan_exclusive ; CHECK-SPV-IR: call spir_func i8 @_Z27__spirv_GroupNonUniformUMaxiih(i32 3, i32 2, i8 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticUChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i8 @_Z32sub_group_non_uniform_reduce_addh(i8 zeroext 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func zeroext i8 @_Z32sub_group_non_uniform_reduce_mulh(i8 zeroext 0) #2 @@ -626,7 +626,7 @@ declare dso_local spir_func zeroext i8 @_Z40sub_group_non_uniform_scan_exclusive ; CHECK-SPV-IR: call spir_func i16 @_Z27__spirv_GroupNonUniformSMaxiis(i32 3, i32 2, i16 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i16 @_Z32sub_group_non_uniform_reduce_adds(i16 signext 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func signext i16 @_Z32sub_group_non_uniform_reduce_muls(i16 signext 0) #2 @@ -745,7 +745,7 @@ declare dso_local spir_func signext i16 @_Z40sub_group_non_uniform_scan_exclusiv ; CHECK-SPV-IR: call spir_func i16 @_Z27__spirv_GroupNonUniformUMaxiit(i32 3, i32 2, i16 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticUShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i16 @_Z32sub_group_non_uniform_reduce_addt(i16 zeroext 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func zeroext i16 @_Z32sub_group_non_uniform_reduce_mult(i16 zeroext 0) #2 @@ -864,7 +864,7 @@ declare dso_local spir_func zeroext i16 @_Z40sub_group_non_uniform_scan_exclusiv ; CHECK-SPV-IR: call spir_func i32 @_Z27__spirv_GroupNonUniformSMaxiii(i32 3, i32 2, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z32sub_group_non_uniform_reduce_addi(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z32sub_group_non_uniform_reduce_muli(i32 0) #2 @@ -983,7 +983,7 @@ declare dso_local spir_func i32 @_Z40sub_group_non_uniform_scan_exclusive_maxi(i ; CHECK-SPV-IR: call spir_func i32 @_Z27__spirv_GroupNonUniformUMaxiij(i32 3, i32 2, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticUInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticUInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z32sub_group_non_uniform_reduce_addj(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z32sub_group_non_uniform_reduce_mulj(i32 0) #2 @@ -1102,7 +1102,7 @@ declare dso_local spir_func i32 @_Z40sub_group_non_uniform_scan_exclusive_maxj(i ; CHECK-SPV-IR: call spir_func i64 @_Z27__spirv_GroupNonUniformSMaxiil(i32 3, i32 2, i64 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticLong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticLong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z32sub_group_non_uniform_reduce_addl(i64 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z32sub_group_non_uniform_reduce_mull(i64 0) #2 @@ -1221,7 +1221,7 @@ declare dso_local spir_func i64 @_Z40sub_group_non_uniform_scan_exclusive_maxl(i ; CHECK-SPV-IR: call spir_func i64 @_Z27__spirv_GroupNonUniformUMaxiim(i32 3, i32 2, i64 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticULong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticULong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z32sub_group_non_uniform_reduce_addm(i64 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z32sub_group_non_uniform_reduce_mulm(i64 0) #2 @@ -1340,7 +1340,7 @@ declare dso_local spir_func i64 @_Z40sub_group_non_uniform_scan_exclusive_maxm(i ; CHECK-SPV-IR: call spir_func float @_Z27__spirv_GroupNonUniformFMaxiif(i32 3, i32 2, float 0.000000e+00) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticFloat(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticFloat(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { %2 = tail call spir_func float @_Z32sub_group_non_uniform_reduce_addf(float 0.000000e+00) #2 store float %2, ptr addrspace(1) %0, align 4, !tbaa !24 %3 = tail call spir_func float @_Z32sub_group_non_uniform_reduce_mulf(float 0.000000e+00) #2 @@ -1459,7 +1459,7 @@ declare dso_local spir_func float @_Z40sub_group_non_uniform_scan_exclusive_maxf ; CHECK-SPV-IR: call spir_func half @_Z27__spirv_GroupNonUniformFMaxiiDh(i32 3, i32 2, half 0xH0000) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticHalf(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticHalf(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { %2 = tail call spir_func half @_Z32sub_group_non_uniform_reduce_addDh(half 0xH0000) #2 store half %2, ptr addrspace(1) %0, align 2, !tbaa !27 %3 = tail call spir_func half @_Z32sub_group_non_uniform_reduce_mulDh(half 0xH0000) #2 @@ -1578,7 +1578,7 @@ declare dso_local spir_func half @_Z40sub_group_non_uniform_scan_exclusive_maxDh ; CHECK-SPV-IR: call spir_func double @_Z27__spirv_GroupNonUniformFMaxiid(i32 3, i32 2, double 0.000000e+00) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformArithmeticDouble(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformArithmeticDouble(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { %2 = tail call spir_func double @_Z32sub_group_non_uniform_reduce_addd(double 0.000000e+00) #2 store double %2, ptr addrspace(1) %0, align 8, !tbaa !30 %3 = tail call spir_func double @_Z32sub_group_non_uniform_reduce_muld(double 0.000000e+00) #2 @@ -1688,7 +1688,7 @@ declare dso_local spir_func double @_Z40sub_group_non_uniform_scan_exclusive_max ; CHECK-SPV-IR: call spir_func i8 @_Z33__spirv_GroupNonUniformBitwiseXoriic(i32 3, i32 2, i8 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i8 @_Z32sub_group_non_uniform_reduce_andc(i8 signext 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func signext i8 @_Z31sub_group_non_uniform_reduce_orc(i8 signext 0) #2 @@ -1780,7 +1780,7 @@ declare dso_local spir_func signext i8 @_Z40sub_group_non_uniform_scan_exclusive ; CHECK-SPV-IR: call spir_func i8 @_Z33__spirv_GroupNonUniformBitwiseXoriic(i32 3, i32 2, i8 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseUChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i8 @_Z32sub_group_non_uniform_reduce_andh(i8 zeroext 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func zeroext i8 @_Z31sub_group_non_uniform_reduce_orh(i8 zeroext 0) #2 @@ -1872,7 +1872,7 @@ declare dso_local spir_func zeroext i8 @_Z40sub_group_non_uniform_scan_exclusive ; CHECK-SPV-IR: call spir_func i16 @_Z33__spirv_GroupNonUniformBitwiseXoriis(i32 3, i32 2, i16 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i16 @_Z32sub_group_non_uniform_reduce_ands(i16 signext 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func signext i16 @_Z31sub_group_non_uniform_reduce_ors(i16 signext 0) #2 @@ -1964,7 +1964,7 @@ declare dso_local spir_func signext i16 @_Z40sub_group_non_uniform_scan_exclusiv ; CHECK-SPV-IR: call spir_func i16 @_Z33__spirv_GroupNonUniformBitwiseXoriis(i32 3, i32 2, i16 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseUShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i16 @_Z32sub_group_non_uniform_reduce_andt(i16 zeroext 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func zeroext i16 @_Z31sub_group_non_uniform_reduce_ort(i16 zeroext 0) #2 @@ -2057,7 +2057,7 @@ declare dso_local spir_func zeroext i16 @_Z40sub_group_non_uniform_scan_exclusiv ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z32sub_group_non_uniform_reduce_andi(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z31sub_group_non_uniform_reduce_ori(i32 0) #2 @@ -2149,7 +2149,7 @@ declare dso_local spir_func i32 @_Z40sub_group_non_uniform_scan_exclusive_xori(i ; CHECK-SPV-IR: call spir_func i32 @_Z33__spirv_GroupNonUniformBitwiseXoriii(i32 3, i32 2, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseUInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseUInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z32sub_group_non_uniform_reduce_andj(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z31sub_group_non_uniform_reduce_orj(i32 0) #2 @@ -2241,7 +2241,7 @@ declare dso_local spir_func i32 @_Z40sub_group_non_uniform_scan_exclusive_xorj(i ; CHECK-SPV-IR: call spir_func i64 @_Z33__spirv_GroupNonUniformBitwiseXoriil(i32 3, i32 2, i64 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseLong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseLong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z32sub_group_non_uniform_reduce_andl(i64 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z31sub_group_non_uniform_reduce_orl(i64 0) #2 @@ -2333,7 +2333,7 @@ declare dso_local spir_func i64 @_Z40sub_group_non_uniform_scan_exclusive_xorl(i ; CHECK-SPV-IR: call spir_func i64 @_Z33__spirv_GroupNonUniformBitwiseXoriil(i32 3, i32 2, i64 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformBitwiseULong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformBitwiseULong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z32sub_group_non_uniform_reduce_andm(i64 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z31sub_group_non_uniform_reduce_orm(i64 0) #2 @@ -2425,7 +2425,7 @@ declare dso_local spir_func i64 @_Z40sub_group_non_uniform_scan_exclusive_xorm(i ; CHECK-SPV-IR: call spir_func i1 @_Z33__spirv_GroupNonUniformLogicalXoriib(i32 3, i32 2, i1 false) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testNonUniformLogical(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testNonUniformLogical(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z40sub_group_non_uniform_reduce_logical_andi(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z39sub_group_non_uniform_reduce_logical_ori(i32 0) #2 diff --git a/test/transcoding/sub_group_non_uniform_vote.ll b/test/transcoding/sub_group_non_uniform_vote.ll index 4c8347e19f..541f69e9d3 100644 --- a/test/transcoding/sub_group_non_uniform_vote.ll +++ b/test/transcoding/sub_group_non_uniform_vote.ll @@ -104,7 +104,7 @@ target triple = "spir64" ; CHECK-LLVM: call spir_func i32 @_Z15sub_group_electv() ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testSubGroupElect(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testSubGroupElect(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z15sub_group_electv() #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !7 ret void @@ -121,7 +121,7 @@ declare dso_local spir_func i32 @_Z15sub_group_electv() local_unnamed_addr #1 ; CHECK-LLVM: call spir_func i32 @_Z25sub_group_non_uniform_alli(i32 {{.*}}) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testSubGroupNonUniformAll(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testSubGroupNonUniformAll(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z25sub_group_non_uniform_alli(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !7 ret void @@ -138,7 +138,7 @@ declare dso_local spir_func i32 @_Z25sub_group_non_uniform_alli(i32) local_unnam ; CHECK-LLVM: call spir_func i32 @_Z25sub_group_non_uniform_anyi(i32 {{.*}}) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testSubGroupNonUniformAny(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testSubGroupNonUniformAny(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z25sub_group_non_uniform_anyi(i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !7 ret void @@ -175,7 +175,7 @@ declare dso_local spir_func i32 @_Z25sub_group_non_uniform_anyi(i32) local_unnam ; CHECK-LLVM: call spir_func i32 @_Z31sub_group_non_uniform_all_equald(double {{.*}}) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testSubGroupNonUniformAllEqual(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testSubGroupNonUniformAllEqual(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z31sub_group_non_uniform_all_equalc(i8 signext 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !7 %3 = tail call spir_func i32 @_Z31sub_group_non_uniform_all_equalh(i8 zeroext 0) #2 diff --git a/test/transcoding/sub_group_shuffle.ll b/test/transcoding/sub_group_shuffle.ll index 611dc71dfa..6d2f93b04d 100644 --- a/test/transcoding/sub_group_shuffle.ll +++ b/test/transcoding/sub_group_shuffle.ll @@ -127,7 +127,7 @@ target triple = "spir64" ; CHECK-SPV-IR: call spir_func i8 @_Z33__spirv_GroupNonUniformShuffleXoricj(i32 3, i8 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i8 @_Z17sub_group_shufflecj(i8 signext 0, i32 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func signext i8 @_Z21sub_group_shuffle_xorcj(i8 signext 0, i32 0) #2 @@ -156,7 +156,7 @@ declare dso_local spir_func signext i8 @_Z21sub_group_shuffle_xorcj(i8 signext, ; CHECK-SPV-IR: call spir_func i8 @_Z33__spirv_GroupNonUniformShuffleXoricj(i32 3, i8 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleUChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i8 @_Z17sub_group_shufflehj(i8 zeroext 0, i32 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func zeroext i8 @_Z21sub_group_shuffle_xorhj(i8 zeroext 0, i32 0) #2 @@ -185,7 +185,7 @@ declare dso_local spir_func zeroext i8 @_Z21sub_group_shuffle_xorhj(i8 zeroext, ; CHECK-SPV-IR: call spir_func i16 @_Z33__spirv_GroupNonUniformShuffleXorisj(i32 3, i16 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i16 @_Z17sub_group_shufflesj(i16 signext 0, i32 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func signext i16 @_Z21sub_group_shuffle_xorsj(i16 signext 0, i32 0) #2 @@ -214,7 +214,7 @@ declare dso_local spir_func signext i16 @_Z21sub_group_shuffle_xorsj(i16 signext ; CHECK-SPV-IR: call spir_func i16 @_Z33__spirv_GroupNonUniformShuffleXorisj(i32 3, i16 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleUShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i16 @_Z17sub_group_shuffletj(i16 zeroext 0, i32 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func zeroext i16 @_Z21sub_group_shuffle_xortj(i16 zeroext 0, i32 0) #2 @@ -243,7 +243,7 @@ declare dso_local spir_func zeroext i16 @_Z21sub_group_shuffle_xortj(i16 zeroext ; CHECK-SPV-IR: call spir_func i32 @_Z33__spirv_GroupNonUniformShuffleXoriij(i32 3, i32 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z17sub_group_shuffleij(i32 0, i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z21sub_group_shuffle_xorij(i32 0, i32 0) #2 @@ -272,7 +272,7 @@ declare dso_local spir_func i32 @_Z21sub_group_shuffle_xorij(i32, i32) local_unn ; CHECK-SPV-IR: call spir_func i32 @_Z33__spirv_GroupNonUniformShuffleXoriij(i32 3, i32 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleUInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleUInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z17sub_group_shufflejj(i32 0, i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z21sub_group_shuffle_xorjj(i32 0, i32 0) #2 @@ -301,7 +301,7 @@ declare dso_local spir_func i32 @_Z21sub_group_shuffle_xorjj(i32, i32) local_unn ; CHECK-SPV-IR: call spir_func i64 @_Z33__spirv_GroupNonUniformShuffleXorilj(i32 3, i64 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleLong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleLong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z17sub_group_shufflelj(i64 0, i32 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z21sub_group_shuffle_xorlj(i64 0, i32 0) #2 @@ -330,7 +330,7 @@ declare dso_local spir_func i64 @_Z21sub_group_shuffle_xorlj(i64, i32) local_unn ; CHECK-SPV-IR: call spir_func i64 @_Z33__spirv_GroupNonUniformShuffleXorilj(i32 3, i64 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleULong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleULong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z17sub_group_shufflemj(i64 0, i32 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z21sub_group_shuffle_xormj(i64 0, i32 0) #2 @@ -359,7 +359,7 @@ declare dso_local spir_func i64 @_Z21sub_group_shuffle_xormj(i64, i32) local_unn ; CHECK-SPV-IR: call spir_func float @_Z33__spirv_GroupNonUniformShuffleXorifj(i32 3, float 0.000000e+00, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleFloat(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleFloat(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { %2 = tail call spir_func float @_Z17sub_group_shufflefj(float 0.000000e+00, i32 0) #2 store float %2, ptr addrspace(1) %0, align 4, !tbaa !24 %3 = tail call spir_func float @_Z21sub_group_shuffle_xorfj(float 0.000000e+00, i32 0) #2 @@ -388,7 +388,7 @@ declare dso_local spir_func float @_Z21sub_group_shuffle_xorfj(float, i32) local ; CHECK-SPV-IR: call spir_func half @_Z33__spirv_GroupNonUniformShuffleXoriDhj(i32 3, half 0xH0000, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleHalf(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleHalf(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { %2 = tail call spir_func half @_Z17sub_group_shuffleDhj(half 0xH0000, i32 0) #2 store half %2, ptr addrspace(1) %0, align 2, !tbaa !27 %3 = tail call spir_func half @_Z21sub_group_shuffle_xorDhj(half 0xH0000, i32 0) #2 @@ -417,7 +417,7 @@ declare dso_local spir_func half @_Z21sub_group_shuffle_xorDhj(half, i32) local_ ; CHECK-SPV-IR: call spir_func double @_Z33__spirv_GroupNonUniformShuffleXoridj(i32 3, double 0.000000e+00, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleDouble(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleDouble(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { %2 = tail call spir_func double @_Z17sub_group_shuffledj(double 0.000000e+00, i32 0) #2 store double %2, ptr addrspace(1) %0, align 8, !tbaa !30 %3 = tail call spir_func double @_Z21sub_group_shuffle_xordj(double 0.000000e+00, i32 0) #2 diff --git a/test/transcoding/sub_group_shuffle_relative.ll b/test/transcoding/sub_group_shuffle_relative.ll index 8aa7f775b4..4d90d90f48 100644 --- a/test/transcoding/sub_group_shuffle_relative.ll +++ b/test/transcoding/sub_group_shuffle_relative.ll @@ -121,7 +121,7 @@ target triple = "spir64" ; CHECK-LLVM: call spir_func i8 @_Z22sub_group_shuffle_downcj(i8 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i8 @_Z20sub_group_shuffle_upcj(i8 signext 0, i32 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func signext i8 @_Z22sub_group_shuffle_downcj(i8 signext 0, i32 0) #2 @@ -146,7 +146,7 @@ declare dso_local spir_func signext i8 @_Z22sub_group_shuffle_downcj(i8 signext, ; CHECK-LLVM: call spir_func i8 @_Z22sub_group_shuffle_downcj(i8 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeUChar(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeUChar(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !10 !kernel_arg_base_type !10 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i8 @_Z20sub_group_shuffle_uphj(i8 zeroext 0, i32 0) #2 store i8 %2, ptr addrspace(1) %0, align 1, !tbaa !7 %3 = tail call spir_func zeroext i8 @_Z22sub_group_shuffle_downhj(i8 zeroext 0, i32 0) #2 @@ -171,7 +171,7 @@ declare dso_local spir_func zeroext i8 @_Z22sub_group_shuffle_downhj(i8 zeroext, ; CHECK-LLVM: call spir_func i16 @_Z22sub_group_shuffle_downsj(i16 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !11 !kernel_arg_base_type !11 !kernel_arg_type_qual !6 { %2 = tail call spir_func signext i16 @_Z20sub_group_shuffle_upsj(i16 signext 0, i32 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func signext i16 @_Z22sub_group_shuffle_downsj(i16 signext 0, i32 0) #2 @@ -196,7 +196,7 @@ declare dso_local spir_func signext i16 @_Z22sub_group_shuffle_downsj(i16 signex ; CHECK-LLVM: call spir_func i16 @_Z22sub_group_shuffle_downsj(i16 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeUShort(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeUShort(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !14 !kernel_arg_base_type !14 !kernel_arg_type_qual !6 { %2 = tail call spir_func zeroext i16 @_Z20sub_group_shuffle_uptj(i16 zeroext 0, i32 0) #2 store i16 %2, ptr addrspace(1) %0, align 2, !tbaa !12 %3 = tail call spir_func zeroext i16 @_Z22sub_group_shuffle_downtj(i16 zeroext 0, i32 0) #2 @@ -221,7 +221,7 @@ declare dso_local spir_func zeroext i16 @_Z22sub_group_shuffle_downtj(i16 zeroex ; CHECK-LLVM: call spir_func i32 @_Z22sub_group_shuffle_downij(i32 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !15 !kernel_arg_base_type !15 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z20sub_group_shuffle_upij(i32 0, i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z22sub_group_shuffle_downij(i32 0, i32 0) #2 @@ -246,7 +246,7 @@ declare dso_local spir_func i32 @_Z22sub_group_shuffle_downij(i32, i32) local_un ; CHECK-LLVM: call spir_func i32 @_Z22sub_group_shuffle_downij(i32 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeUInt(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeUInt(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !18 !kernel_arg_base_type !18 !kernel_arg_type_qual !6 { %2 = tail call spir_func i32 @_Z20sub_group_shuffle_upjj(i32 0, i32 0) #2 store i32 %2, ptr addrspace(1) %0, align 4, !tbaa !16 %3 = tail call spir_func i32 @_Z22sub_group_shuffle_downjj(i32 0, i32 0) #2 @@ -271,7 +271,7 @@ declare dso_local spir_func i32 @_Z22sub_group_shuffle_downjj(i32, i32) local_un ; CHECK-LLVM: call spir_func i64 @_Z22sub_group_shuffle_downlj(i64 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeLong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeLong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !19 !kernel_arg_base_type !19 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z20sub_group_shuffle_uplj(i64 0, i32 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z22sub_group_shuffle_downlj(i64 0, i32 0) #2 @@ -296,7 +296,7 @@ declare dso_local spir_func i64 @_Z22sub_group_shuffle_downlj(i64, i32) local_un ; CHECK-LLVM: call spir_func i64 @_Z22sub_group_shuffle_downlj(i64 0, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeULong(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeULong(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !22 !kernel_arg_base_type !22 !kernel_arg_type_qual !6 { %2 = tail call spir_func i64 @_Z20sub_group_shuffle_upmj(i64 0, i32 0) #2 store i64 %2, ptr addrspace(1) %0, align 8, !tbaa !20 %3 = tail call spir_func i64 @_Z22sub_group_shuffle_downmj(i64 0, i32 0) #2 @@ -321,7 +321,7 @@ declare dso_local spir_func i64 @_Z22sub_group_shuffle_downmj(i64, i32) local_un ; CHECK-LLVM: call spir_func float @_Z22sub_group_shuffle_downfj(float 0.000000e+00, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeFloat(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeFloat(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !23 !kernel_arg_base_type !23 !kernel_arg_type_qual !6 { %2 = tail call spir_func float @_Z20sub_group_shuffle_upfj(float 0.000000e+00, i32 0) #2 store float %2, ptr addrspace(1) %0, align 4, !tbaa !24 %3 = tail call spir_func float @_Z22sub_group_shuffle_downfj(float 0.000000e+00, i32 0) #2 @@ -346,7 +346,7 @@ declare dso_local spir_func float @_Z22sub_group_shuffle_downfj(float, i32) loca ; CHECK-LLVM: call spir_func half @_Z22sub_group_shuffle_downDhj(half 0xH0000, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeHalf(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeHalf(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !26 !kernel_arg_base_type !26 !kernel_arg_type_qual !6 { %2 = tail call spir_func half @_Z20sub_group_shuffle_upDhj(half 0xH0000, i32 0) #2 store half %2, ptr addrspace(1) %0, align 2, !tbaa !27 %3 = tail call spir_func half @_Z22sub_group_shuffle_downDhj(half 0xH0000, i32 0) #2 @@ -371,7 +371,7 @@ declare dso_local spir_func half @_Z22sub_group_shuffle_downDhj(half, i32) local ; CHECK-LLVM: call spir_func double @_Z22sub_group_shuffle_downdj(double 0.000000e+00, i32 0) ; Function Attrs: convergent nounwind -define dso_local spir_kernel void @testShuffleRelativeDouble(ptr addrspace(1) nocapture) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { +define dso_local spir_kernel void @testShuffleRelativeDouble(ptr addrspace(1) captures(none)) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !29 !kernel_arg_base_type !29 !kernel_arg_type_qual !6 { %2 = tail call spir_func double @_Z20sub_group_shuffle_updj(double 0.000000e+00, i32 0) #2 store double %2, ptr addrspace(1) %0, align 8, !tbaa !30 %3 = tail call spir_func double @_Z22sub_group_shuffle_downdj(double 0.000000e+00, i32 0) #2 diff --git a/test/transcoding/sycl_array_zero_init.ll b/test/transcoding/sycl_array_zero_init.ll index 31af481090..0ee6367268 100644 --- a/test/transcoding/sycl_array_zero_init.ll +++ b/test/transcoding/sycl_array_zero_init.ll @@ -25,13 +25,13 @@ define weak_odr dso_local spir_kernel void @main() #0 comdat !kernel_arg_addr_sp } ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) #1 +declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #1 ; Function Attrs: argmemonly nounwind -declare void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) nocapture writeonly, ptr addrspace(4) nocapture readonly, i64, i1 immarg) #1 +declare void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) captures(none) writeonly, ptr addrspace(4) captures(none) readonly, i64, i1 immarg) #1 attributes #0 = { "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } attributes #1 = { argmemonly nounwind } diff --git a/test/transcoding/unreachable.ll b/test/transcoding/unreachable.ll index 3879c26e1c..7dae3e5bb0 100644 --- a/test/transcoding/unreachable.ll +++ b/test/transcoding/unreachable.ll @@ -14,7 +14,7 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256: target triple = "spir64-unknown-unknown" ; Function Attrs: nounwind -define spir_kernel void @unreachable_simple(ptr addrspace(1) nocapture %in, ptr addrspace(1) %out) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @unreachable_simple(ptr addrspace(1) captures(none) %in, ptr addrspace(1) %out) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { %1 = call spir_func i64 @_Z13get_global_idj(i32 0) #1 %2 = shl i64 %1, 32 %3 = ashr exact i64 %2, 32 diff --git a/test/transcoding/vector_casts.ll b/test/transcoding/vector_casts.ll index ec0c1b4c70..0f61e526ca 100644 --- a/test/transcoding/vector_casts.ll +++ b/test/transcoding/vector_casts.ll @@ -22,7 +22,7 @@ target triple = "spir64-unknown-unknown" ; CHECK: call spir_func <8 x i32> @_Z12convert_int8Dv8_f(<8 x float> ; Function Attrs: nounwind -define spir_kernel void @test_default_conversions(ptr addrspace(1) nocapture %out, <8 x i8> %in) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { +define spir_kernel void @test_default_conversions(ptr addrspace(1) captures(none) %out, <8 x i8> %in) #0 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_base_type !5 !kernel_arg_type_qual !4 { %1 = tail call spir_func <8 x i16> @_Z15convert_ushort8Dv8_c(<8 x i8> %in) #1 %2 = tail call spir_func <8 x i32> @_Z12convert_int8Dv8_t(<8 x i16> %1) #1 %3 = tail call spir_func <8 x i8> @_Z13convert_char8Dv8_i(<8 x i32> %2) #1 diff --git a/test/uitofp-with-bool.ll b/test/uitofp-with-bool.ll index 16052d6f2f..729982f61b 100644 --- a/test/uitofp-with-bool.ll +++ b/test/uitofp-with-bool.ll @@ -85,7 +85,7 @@ target triple = "spir64" ; SPV-DAG: FunctionParameter {{[0-9]+}} [[i1v:[0-9]+]] ; Function Attrs: nofree norecurse nounwind writeonly -define dso_local spir_kernel void @K(ptr addrspace(1) nocapture %A, i32 %B, i1 %i1s, <2 x i1> %i1v) local_unnamed_addr #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { +define dso_local spir_kernel void @K(ptr addrspace(1) captures(none) %A, i32 %B, i1 %i1s, <2 x i1> %i1v) local_unnamed_addr #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 { entry: