Skip to content

[SYCL][LowerWGScope][SpecConstantsPass] fix new call's calling conv for spir/spirv #15718

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Oct 21, 2024
Merged

[SYCL][LowerWGScope][SpecConstantsPass] fix new call's calling conv for spir/spirv #15718

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
033a1f6
[SYCL][LowerWGScope] Set new call's calling conv to spirv_func for sp…
wenju-he Oct 16, 2024
c342e0a
isSPIR -> isSPIROrSPIRV
wenju-he Oct 17, 2024
14509a1
fix function and call in SpecConstantsPass
wenju-he Oct 17, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions llvm/lib/SYCLLowerIR/LowerWGScope.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<Function>(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC);

IRBuilder<> Bld(Ctx);
Bld.SetInsertPoint(&Before);
Expand All @@ -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;
}
13 changes: 12 additions & 1 deletion llvm/lib/SYCLLowerIR/SpecConstants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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");
Expand All @@ -500,6 +501,11 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,

auto *Call =
CallInst::Create(NewFT, NewFC.getCallee(), Args, "", InsertBefore);
if (IsSPIROrSPIRV) {
cast<Function>(NewFC.getCallee())
->setCallingConv(CallingConv::SPIR_FUNC);
Call->setCallingConv(CallingConv::SPIR_FUNC);
}
return CastInst::CreateTruncOrBitCast(Call, RetTy, "tobool",
InsertBefore);
}
Expand All @@ -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<Function>(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC);
Call->setCallingConv(CallingConv::SPIR_FUNC);
}
return Call;
}

Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,
Expand Down
27 changes: 27 additions & 0 deletions llvm/test/SYCLLowerIR/LowerWGScope/barrier-calling-conv.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 = !{}
68 changes: 34 additions & 34 deletions llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)*
Expand Down
12 changes: 6 additions & 6 deletions llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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}
Expand Down
Loading
Loading