diff --git a/llvm/lib/SYCLLowerIR/LowerWGScope.cpp b/llvm/lib/SYCLLowerIR/LowerWGScope.cpp index 882129b369da1..b087221e5f37a 100644 --- a/llvm/lib/SYCLLowerIR/LowerWGScope.cpp +++ b/llvm/lib/SYCLLowerIR/LowerWGScope.cpp @@ -1001,6 +1001,8 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) { FunctionCallee FC = M.getOrInsertFunction(Name, Attr, RetTy, ScopeTy, ScopeTy, SemanticsTy); assert(FC.getCallee() && "spirv intrinsic creation failed"); + if (TT.isSPIROrSPIRV()) + cast(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC); IRBuilder<> Bld(Ctx); Bld.SetInsertPoint(&Before); @@ -1011,5 +1013,7 @@ Instruction *spirv::genWGBarrier(Instruction &Before, const Triple &TT) { asUInt(spirv::MemorySemantics::WorkgroupMemory)); auto BarrierCall = Bld.CreateCall(FC, {ArgExec, ArgMem, ArgSema}); BarrierCall->addFnAttr(llvm::Attribute::Convergent); + if (TT.isSPIROrSPIRV()) + BarrierCall->setCallingConv(CallingConv::SPIR_FUNC); return BarrierCall; } diff --git a/llvm/lib/SYCLLowerIR/SpecConstants.cpp b/llvm/lib/SYCLLowerIR/SpecConstants.cpp index 83bb489486911..ab31be8ec831a 100644 --- a/llvm/lib/SYCLLowerIR/SpecConstants.cpp +++ b/llvm/lib/SYCLLowerIR/SpecConstants.cpp @@ -475,6 +475,7 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, auto *FT = FunctionType::get(RetTy, ArgTys, false /*isVarArg*/); std::string FunctionName = mangleFuncItanium(BaseFunctionName, FT); Module *M = InsertBefore->getFunction()->getParent(); + bool IsSPIROrSPIRV = llvm::Triple(M->getTargetTriple()).isSPIROrSPIRV(); if (RetTy->isIntegerTy(1)) { assert(ArgTys.size() == 2 && "Expected a scalar spec constant"); @@ -500,6 +501,11 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, auto *Call = CallInst::Create(NewFT, NewFC.getCallee(), Args, "", InsertBefore); + if (IsSPIROrSPIRV) { + cast(NewFC.getCallee()) + ->setCallingConv(CallingConv::SPIR_FUNC); + Call->setCallingConv(CallingConv::SPIR_FUNC); + } return CastInst::CreateTruncOrBitCast(Call, RetTy, "tobool", InsertBefore); } @@ -520,7 +526,12 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName, // types? Is it necessary? FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT); - return CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); + auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore); + if (IsSPIROrSPIRV) { + cast(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC); + Call->setCallingConv(CallingConv::SPIR_FUNC); + } + return Call; } Instruction *emitSpecConstant(unsigned NumericID, Type *Ty, diff --git a/llvm/test/SYCLLowerIR/LowerWGScope/barrier-calling-conv.ll b/llvm/test/SYCLLowerIR/LowerWGScope/barrier-calling-conv.ll new file mode 100644 index 0000000000000..cd7a1a4ba66e5 --- /dev/null +++ b/llvm/test/SYCLLowerIR/LowerWGScope/barrier-calling-conv.ll @@ -0,0 +1,27 @@ +; RUN: opt -passes=LowerWGScope -S %s -o - | FileCheck %s + +; Check newly created barrier call has spir_func calling convention. + +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1" +target triple = "spir64-unknown-unknown" + +%"class.sycl::_V1::group" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::id" } +%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" } +%"class.sycl::_V1::detail::array" = type { [1 x i64] } +%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" } + +define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_5groupILi1EEEE_clES5_(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this, ptr noundef byval(%"class.sycl::_V1::group") align 8 %group) !work_group_scope !0 { +entry: +; CHECK: call spir_func void @_Z22__spirv_ControlBarrierjjj( + + %this.addr = alloca ptr addrspace(4), align 8 + %this.addr.ascast = addrspacecast ptr %this.addr to ptr addrspace(4) + store ptr addrspace(4) %this, ptr addrspace(4) %this.addr.ascast, align 8 + %group.ascast = addrspacecast ptr %group to ptr addrspace(4) + %this1 = load ptr addrspace(4), ptr addrspace(4) %this.addr.ascast, align 8 + ret void +} + +; CHECK: declare spir_func void @_Z22__spirv_ControlBarrierjjj( + +!0 = !{} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll index 36b8a81aaa4c3..effea3aa01e03 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll @@ -39,44 +39,44 @@ define weak_odr spir_kernel void @_ZN2cl4sycl14kernel_handler33getSpecialization %9 = getelementptr inbounds %"class.cl::sycl::kernel_handler", %"class.cl::sycl::kernel_handler" addrspace(4)* %7, i32 0, i32 0 %10 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %9, align 8, !tbaa !8 call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueISt5arrayIS0_IfLy3EELy3EEET_PKcPKvS7_(%"class.std::array" addrspace(4)* sret(%"class.std::array") align 4 %0, i8 addrspace(4)* %8, i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id" addrspace(1)* @_ZL8coeff_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* %10) #13 -; CHECK-IR: %[[#NS0:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00) -; CHECK-IR: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00) -; CHECK-IR: %[[#NS2:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00) -; CHECK-IR: %[[#NS3:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]]) -; CHECK-IR: %[[#NS4:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS3]]) -; CHECK-IR: %[[#NS5:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID3:]], float 0.000000e+00) -; CHECK-IR: %[[#NS6:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00) -; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00) -; CHECK-IR: %[[#NS8:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS5]], float %[[#NS6]], float %[[#NS7]]) -; CHECK-IR: %[[#NS9:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS8]]) -; CHECK-IR: %[[#NS10:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00) -; CHECK-IR: %[[#NS11:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID7:]], float 0.000000e+00) -; CHECK-IR: %[[#NS12:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID8:]], float 0.000000e+00) -; CHECK-IR: %[[#NS13:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS10]], float %[[#NS11]], float %[[#NS12]]) -; CHECK-IR: %[[#NS14:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS13]]) -; CHECK-IR: %[[#NS15:]] = call [3 x %"class.std::array.1"] @"_Z29__spirv_SpecConstantCompositeclass.std::array.1class.std::array.1class.std::array.1_RA3_class.std::array.1"(%"class.std::array.1" %[[#NS4]], %"class.std::array.1" %[[#NS9]], %"class.std::array.1" %[[#NS14]]) -; CHECK-IR: %[[#NS16:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS15]]) +; CHECK-IR: %[[#NS0:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00) +; CHECK-IR: %[[#NS1:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00) +; CHECK-IR: %[[#NS2:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00) +; CHECK-IR: %[[#NS3:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]]) +; CHECK-IR: %[[#NS4:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS3]]) +; CHECK-IR: %[[#NS5:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID3:]], float 0.000000e+00) +; CHECK-IR: %[[#NS6:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00) +; CHECK-IR: %[[#NS7:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00) +; CHECK-IR: %[[#NS8:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS5]], float %[[#NS6]], float %[[#NS7]]) +; CHECK-IR: %[[#NS9:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS8]]) +; CHECK-IR: %[[#NS10:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00) +; CHECK-IR: %[[#NS11:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID7:]], float 0.000000e+00) +; CHECK-IR: %[[#NS12:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID8:]], float 0.000000e+00) +; CHECK-IR: %[[#NS13:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS10]], float %[[#NS11]], float %[[#NS12]]) +; CHECK-IR: %[[#NS14:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS13]]) +; CHECK-IR: %[[#NS15:]] = call spir_func [3 x %"class.std::array.1"] @"_Z29__spirv_SpecConstantCompositeclass.std::array.1class.std::array.1class.std::array.1_RA3_class.std::array.1"(%"class.std::array.1" %[[#NS4]], %"class.std::array.1" %[[#NS9]], %"class.std::array.1" %[[#NS14]]) +; CHECK-IR: %[[#NS16:]] = call spir_func %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS15]]) %11 = alloca %"class.std::array", align 4 %12 = addrspacecast %"class.std::array"* %11 to %"class.std::array" addrspace(4)* call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueISt5arrayIS0_IfLy3EELy3EEET_PKcPKvS7_(%"class.std::array" addrspace(4)* sret(%"class.std::array") align 4 %12, i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([33 x i8], [33 x i8]* @__usid_str.2, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id" addrspace(1)* @_ZL9coeff_id2 to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #13 -; CHECK-IR: %[[#NS17:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID9:]], float 0.000000e+00) -; CHECK-IR: %[[#NS18:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID10:]], float 0.000000e+00) -; CHECK-IR: %[[#NS19:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID11:]], float 0.000000e+00) -; CHECK-IR: %[[#NS20:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS17]], float %[[#NS18]], float %[[#NS19]]) -; CHECK-IR: %[[#NS21:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS20]]) -; CHECK-IR: %[[#NS22:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID12:]], float 0.000000e+00) -; CHECK-IR: %[[#NS23:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID13:]], float 1.000000e+00) -; CHECK-IR: %[[#NS24:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID14:]], float 2.000000e+00) -; CHECK-IR: %[[#NS25:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS22]], float %[[#NS23]], float %[[#NS24]]) -; CHECK-IR: %[[#NS26:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS25]]) -; CHECK-IR: %[[#NS27:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID15:]], float 0x4010666660000000) -; CHECK-IR: %[[#NS28:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID16:]], float 0x4014666660000000) -; CHECK-IR: %[[#NS29:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID17:]], float 0x4018CCCCC0000000) -; CHECK-IR: %[[#NS30:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS27]], float %[[#NS28]], float %[[#NS29]]) -; CHECK-IR: %[[#NS31:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS30]]) -; CHECK-IR: %[[#NS32:]] = call [3 x %"class.std::array.1"] @"_Z29__spirv_SpecConstantCompositeclass.std::array.1class.std::array.1class.std::array.1_RA3_class.std::array.1"(%"class.std::array.1" %[[#NS21]], %"class.std::array.1" %[[#NS26]], %"class.std::array.1" %[[#NS31]]) -; CHECK-IR: %[[#NS33:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS32]]) +; CHECK-IR: %[[#NS17:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID9:]], float 0.000000e+00) +; CHECK-IR: %[[#NS18:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID10:]], float 0.000000e+00) +; CHECK-IR: %[[#NS19:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID11:]], float 0.000000e+00) +; CHECK-IR: %[[#NS20:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS17]], float %[[#NS18]], float %[[#NS19]]) +; CHECK-IR: %[[#NS21:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS20]]) +; CHECK-IR: %[[#NS22:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID12:]], float 0.000000e+00) +; CHECK-IR: %[[#NS23:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID13:]], float 1.000000e+00) +; CHECK-IR: %[[#NS24:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID14:]], float 2.000000e+00) +; CHECK-IR: %[[#NS25:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS22]], float %[[#NS23]], float %[[#NS24]]) +; CHECK-IR: %[[#NS26:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS25]]) +; CHECK-IR: %[[#NS27:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID15:]], float 0x4010666660000000) +; CHECK-IR: %[[#NS28:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID16:]], float 0x4014666660000000) +; CHECK-IR: %[[#NS29:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID17:]], float 0x4018CCCCC0000000) +; CHECK-IR: %[[#NS30:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS27]], float %[[#NS28]], float %[[#NS29]]) +; CHECK-IR: %[[#NS31:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS30]]) +; CHECK-IR: %[[#NS32:]] = call spir_func [3 x %"class.std::array.1"] @"_Z29__spirv_SpecConstantCompositeclass.std::array.1class.std::array.1class.std::array.1_RA3_class.std::array.1"(%"class.std::array.1" %[[#NS21]], %"class.std::array.1" %[[#NS26]], %"class.std::array.1" %[[#NS31]]) +; CHECK-IR: %[[#NS33:]] = call spir_func %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS32]]) %13 = alloca %struct.coeff_str_t, align 8 %14 = addrspacecast %struct.coeff_str_t* %13 to %struct.coeff_str_t addrspace(4)* diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll index 21fdc1c835b79..e67daf1061dc5 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll @@ -31,16 +31,16 @@ entry: %ref.tmp.i = alloca %struct._ZTS9composite.composite, align 8 %ref.tmp.ascast.i = addrspacecast %struct._ZTS9composite.composite* %ref.tmp.i to %struct._ZTS9composite.composite addrspace(4)* %call.i.i.i = tail call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)* getelementptr inbounds ([70 x i8], [70 x i8] addrspace(4)* addrspacecast ([70 x i8] addrspace(1)* @__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL6int_idEiLPv0EEET0_v to [70 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class._ZTSN2cl4sycl17specialization_idIiEE.cl::sycl::specialization_id" addrspace(1)* @_ZL6int_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #3 -; CHECK: call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID0:]], i32 0) +; CHECK: call spir_func i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID0:]], i32 0) %call.i.i23.i = tail call spir_func double @_Z37__sycl_getScalar2020SpecConstantValueIdET_PKcPKvS4_(i8 addrspace(4)* getelementptr inbounds ([73 x i8], [73 x i8] addrspace(4)* addrspacecast ([73 x i8] addrspace(1)* @__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL9double_idEdLPv0EEET0_v to [73 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class._ZTSN2cl4sycl17specialization_idIdEE.cl::sycl::specialization_id" addrspace(1)* @_ZL9double_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #3 -; CHECK: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID1:]], double 0.000000e+00) +; CHECK: call spir_func double @_Z20__spirv_SpecConstantid(i32 [[#SCID1:]], double 0.000000e+00) call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI9compositeET_PKcPKvS5_(%struct._ZTS9composite.composite addrspace(4)* sret(%struct._ZTS9composite.composite) align 8 %ref.tmp.ascast.i, i8 addrspace(4)* getelementptr inbounds ([77 x i8], [77 x i8] addrspace(4)* addrspacecast ([77 x i8] addrspace(1)* @__builtin_unique_stable_name._ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL12composite_idE9compositeLPv0EEET0_v to [77 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class._ZTSN2cl4sycl17specialization_idI9compositeEE.cl::sycl::specialization_id" addrspace(1)* @_ZL12composite_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #3 -; CHECK: call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00) -; CHECK: call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID3:]], i8 0) -; CHECK: call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 0) -; CHECK: call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID5:]], i64 0) +; CHECK: call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00) +; CHECK: call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID3:]], i8 0) +; CHECK: call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 0) +; CHECK: call spir_func i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID5:]], i64 0) ; CHECK-LOG: sycl.specialization-constants ; CHECK-LOG:[[UNIQUE_PREFIX:[0-9a-zA-Z]+]]={0, 0, 4} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll b/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll index 4eb6862b2ca44..48824faa78d75 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll @@ -29,25 +29,25 @@ define weak_odr dso_local spir_kernel void @_ZTSZ4mainEUlN2cl4sycl14kernel_handl %2 = addrspacecast %struct.coeff_str_aligned_t* %1 to %struct.coeff_str_aligned_t addrspace(4)* %3 = bitcast %struct.coeff_str_aligned_t* %1 to i8* call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff_str_aligned_tET_PKcPKvS5_(%struct.coeff_str_aligned_t addrspace(4)* sret(%struct.coeff_str_aligned_t) align 32 %2, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([32 x i8], [32 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id" addrspace(1)* @_ZL8coeff_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4 -; CHECK-IR: %[[#NS0:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00) -; CHECK-IR: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00) -; CHECK-IR: %[[#NS2:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00) -; CHECK-IR: %[[#NS3:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]]) -; CHECK-IR: %[[#NS4:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS3]]) -; CHECK-IR: %[[#NS5:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID3:]], i64 0) -; CHECK-IR: %[[#NS6:]] = call %struct.coeff_str_aligned_t @"_Z29__spirv_SpecConstantCompositeclass.std::arrayxA8_a_Rstruct.coeff_str_aligned_t"(%"class.std::array" %[[#NS4]], i64 %[[#NS5]], [8 x i8] undef) +; CHECK-IR: %[[#NS0:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00) +; CHECK-IR: %[[#NS1:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00) +; CHECK-IR: %[[#NS2:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00) +; CHECK-IR: %[[#NS3:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]]) +; CHECK-IR: %[[#NS4:]] = call spir_func %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS3]]) +; CHECK-IR: %[[#NS5:]] = call spir_func i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID3:]], i64 0) +; CHECK-IR: %[[#NS6:]] = call spir_func %struct.coeff_str_aligned_t @"_Z29__spirv_SpecConstantCompositeclass.std::arrayxA8_a_Rstruct.coeff_str_aligned_t"(%"class.std::array" %[[#NS4]], i64 %[[#NS5]], [8 x i8] undef) %4 = alloca %struct.coeff2_str_aligned_t, align 32 %5 = addrspacecast %struct.coeff2_str_aligned_t* %4 to %struct.coeff2_str_aligned_t addrspace(4)* %6 = bitcast %struct.coeff2_str_aligned_t* %4 to i8* call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff2_str_aligned_tET_PKcPKvS5_(%struct.coeff2_str_aligned_t addrspace(4)* sret(%struct.coeff2_str_aligned_t) align 32 %5, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([33 x i8], [33 x i8]* @__usid_str.0, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id.1" addrspace(1)* @_ZL8coeff_id2 to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4 -; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00) -; CHECK-IR: %[[#NS8:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00) -; CHECK-IR: %[[#NS9:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00) -; CHECK-IR: %[[#NS10:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS7]], float %[[#NS8]], float %[[#NS9]]) -; CHECK-IR: %[[#NS11:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS10]]) -; CHECK-IR: %[[#NS12:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID7:]], i64 0) -; CHECK-IR: %[[#NS13:]] = call %struct.coeff2_str_aligned_t @"_Z29__spirv_SpecConstantCompositeclass.std::arrayxA7_aa_Rstruct.coeff2_str_aligned_t"(%"class.std::array" %[[#NS11]], i64 %[[#NS12]], [7 x i8] undef, i8 undef) +; CHECK-IR: %[[#NS7:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00) +; CHECK-IR: %[[#NS8:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00) +; CHECK-IR: %[[#NS9:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00) +; CHECK-IR: %[[#NS10:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS7]], float %[[#NS8]], float %[[#NS9]]) +; CHECK-IR: %[[#NS11:]] = call spir_func %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS10]]) +; CHECK-IR: %[[#NS12:]] = call spir_func i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID7:]], i64 0) +; CHECK-IR: %[[#NS13:]] = call spir_func %struct.coeff2_str_aligned_t @"_Z29__spirv_SpecConstantCompositeclass.std::arrayxA7_aa_Rstruct.coeff2_str_aligned_t"(%"class.std::array" %[[#NS11]], i64 %[[#NS12]], [7 x i8] undef, i8 undef) ret void } diff --git a/llvm/test/tools/sycl-post-link/spec-constants/bool.ll b/llvm/test/tools/sycl-post-link/spec-constants/bool.ll index 7981450a12f00..5962f20d79ae5 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/bool.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/bool.ll @@ -7,7 +7,7 @@ ; CHECK-LABEL: void @kernel_A -; CHECK-RT: %[[CALL:.*]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#]], i8 1) +; CHECK-RT: %[[CALL:.*]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#]], i8 1) ; CHECK-RT: trunc i8 %[[CALL]] to i1 ; ; CHECK-DEF: %[[GEP:gep.*]] = getelementptr i8, ptr addrspace(4) null, i32 0 @@ -15,7 +15,7 @@ ; CHECK-DEF: %[[TOBOOL:tobool.*]] = trunc i8 %[[LOAD]] to i1 ; ; CHECK-LABEL: void @kernel_B -; CHECK-RT: call i8 @_Z20__spirv_SpecConstantia(i32 [[#]], i8 +; CHECK-RT: call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#]], i8 ; ; CHECK-DEF: %[[GEP:gep.*]] = getelementptr i8, ptr addrspace(4) null, i32 4 ; CHECK-DEF: %[[BC:bc.*]] = bitcast ptr addrspace(4) %[[GEP]] to ptr addrspace(4) diff --git a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll index 2ef1141cbd740..8da5923879b36 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/composite-O2.ll @@ -7,9 +7,9 @@ ; friendly IR operations representing those constants. ; ; CHECK-LABEL: define {{.*}} spir_kernel void @_ZTS17SpecializedKernel -; CHECK: %[[#N0:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32 -; CHECK: %[[#N1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID + 1]], float -; CHECK: %[[#CONST:]] = call %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#N0]], float %[[#N1]]) +; CHECK: %[[#N0:]] = call spir_func i32 @_Z20__spirv_SpecConstantii(i32 [[#ID:]], i32 +; CHECK: %[[#N1:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#ID + 1]], float +; CHECK: %[[#CONST:]] = call spir_func %struct._ZTS1A.A @_Z29__spirv_SpecConstantCompositeif_Rstruct._ZTS1A.A(i32 %[[#N0]], float %[[#N1]]) ; CHECK: %struct._ZTS1A.A %[[#CONST]] ; ; CHECK: !sycl.specialization-constants = !{![[#MD0:]]} diff --git a/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll b/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll index d38e2469b7e2b..c00e8fc3d8365 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll @@ -19,9 +19,9 @@ ; CHECK-PROP1: specConstsReplacedWithDefault=1|1 -; CHECK-IR0: call i32 @_Z20__spirv_SpecConstantii -; CHECK-IR0: call %struct.B @_Z29__spirv_SpecConstantCompositeiii_Rstruct.B -; CHECK-IR0: call %struct.A @_Z29__spirv_SpecConstantCompositeistruct.B_Rstruct.A +; CHECK-IR0: call spir_func i32 @_Z20__spirv_SpecConstantii +; CHECK-IR0: call spir_func %struct.B @_Z29__spirv_SpecConstantCompositeiii_Rstruct.B +; CHECK-IR0: call spir_func %struct.A @_Z29__spirv_SpecConstantCompositeistruct.B_Rstruct.A ; CHECK-IR1: store %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } }, ptr addrspace(4) %a.ascast.i, align 4 diff --git a/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll b/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll index d68bf98094103..0744742817758 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/nested-struct.ll @@ -5,13 +5,13 @@ ; RUN: FileCheck %s -input-file=%t_0.ll ; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst -spec-const=native < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} -; CHECK: %[[#SCV1:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID1:]], i8 120) -; CHECK: %[[#SCV2:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID2:]], i8 121) -; CHECK: %[[#SCV3:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID3:]], i32 122) -; CHECK: %[[#SCV4:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 97) -; CHECK: %[[#SCV5:]] = call %struct.anon @_Z29__spirv_SpecConstantCompositeaia_Rstruct.anon(i8 %[[#SCV2:]], i32 %[[#SCV3:]], i8 %[[#SCV4:]]) -; CHECK: %[[#SCV6:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID5:]], i8 98) -; CHECK: call %struct.user_defined_type3 @_Z29__spirv_SpecConstantCompositeastruct.anona_Rstruct.user_defined_type3(i8 %[[#SCV1:]], %struct.anon %[[#SCV5:]], i8 %[[#SCV6:]]) +; CHECK: %[[#SCV1:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID1:]], i8 120) +; CHECK: %[[#SCV2:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID2:]], i8 121) +; CHECK: %[[#SCV3:]] = call spir_func i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID3:]], i32 122) +; CHECK: %[[#SCV4:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 97) +; CHECK: %[[#SCV5:]] = call spir_func %struct.anon @_Z29__spirv_SpecConstantCompositeaia_Rstruct.anon(i8 %[[#SCV2:]], i32 %[[#SCV3:]], i8 %[[#SCV4:]]) +; CHECK: %[[#SCV6:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID5:]], i8 98) +; CHECK: call spir_func %struct.user_defined_type3 @_Z29__spirv_SpecConstantCompositeastruct.anona_Rstruct.user_defined_type3(i8 %[[#SCV1:]], %struct.anon %[[#SCV5:]], i8 %[[#SCV6:]]) ; CHECK: !sycl.specialization-constants = !{![[#SC:]]} ; CHECK: ![[#SC]] = !{!"uid0a28d8a0a23067ab____ZL8spec_id3", diff --git a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll index 2dc002140362f..80b03beb4f440 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-padding-in-the-middle.ll @@ -39,11 +39,11 @@ ; a situation, where spec constant default value contains less elements than ; spec constant type, due to padding inserted by a compiler. -; CHECK: %[[#SCV1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0x40091EB860000000) -; CHECK: %[[#SCV2:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID2:]], i8 97) -; CHECK: %[[#SCV3:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID3:]], i32 42) -; CHECK: %[[#SCV4:]] = call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 8) -; CHECK: call %struct.user_defined_type @_Z29__spirv_SpecConstantCompositefaA3_aiaA3_a_Rstruct.user_defined_type(float %[[#SCV1]], i8 %[[#SCV2]], [3 x i8] undef, i32 %[[#SCV3]], i8 %[[#SCV4]], [3 x i8] undef) +; CHECK: %[[#SCV1:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0x40091EB860000000) +; CHECK: %[[#SCV2:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID2:]], i8 97) +; CHECK: %[[#SCV3:]] = call spir_func i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID3:]], i32 42) +; CHECK: %[[#SCV4:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 8) +; CHECK: call spir_func %struct.user_defined_type @_Z29__spirv_SpecConstantCompositefaA3_aiaA3_a_Rstruct.user_defined_type(float %[[#SCV1]], i8 %[[#SCV2]], [3 x i8] undef, i32 %[[#SCV3]], i8 %[[#SCV4]], [3 x i8] undef) ; CHECK: !sycl.specialization-constants = !{![[#SC:]]} ; CHECK: ![[#SC]] = !{!"uidc5885cee0b80ad9d____ZL7spec_id", diff --git a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll index 460699ae68ed3..a7c73f4e52f40 100644 --- a/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll +++ b/llvm/test/tools/sycl-post-link/spec-constants/struct-with-undef-padding-2.ll @@ -38,10 +38,10 @@ ; RUN: FileCheck %s -input-file=%t.files_0.ll ; RUN: %if asserts %{ sycl-post-link -properties -debug-only=SpecConst --spec-const=native -S < %s 2>&1 | FileCheck %s --check-prefix=CHECK-LOG %} ; -; CHECK: %[[#A:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#ID:]], float 0x40091EB860000000) -; CHECK: %[[#B:]] = call i32 @_Z20__spirv_SpecConstantii(i32 [[#ID+1]], i32 42) -; CHECK: %[[#C:]] = call i8 @_Z20__spirv_SpecConstantia(i32 2, i8 8) -; CHECK: call %struct.user_defined_type @_Z29__spirv_SpecConstantCompositefiaA3_a_Rstruct.user_defined_type(float %[[#A]], i32 %[[#B]], i8 %[[#C]], [3 x i8] undef) +; CHECK: %[[#A:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#ID:]], float 0x40091EB860000000) +; CHECK: %[[#B:]] = call spir_func i32 @_Z20__spirv_SpecConstantii(i32 [[#ID+1]], i32 42) +; CHECK: %[[#C:]] = call spir_func i8 @_Z20__spirv_SpecConstantia(i32 2, i8 8) +; CHECK: call spir_func %struct.user_defined_type @_Z29__spirv_SpecConstantCompositefiaA3_a_Rstruct.user_defined_type(float %[[#A]], i32 %[[#B]], i8 %[[#C]], [3 x i8] undef) ; ; CHECK: !sycl.specialization-constants = !{![[#SC:]]} ; CHECK: ![[#SC]] = !{!"uidac684fbd602505be____ZL7spec_id",