Skip to content

Commit 14509a1

Browse files
wenju-hewenju-he
committed
fix function and call in SpecConstantsPass
1 parent c342e0a commit 14509a1

File tree

10 files changed

+90
-79
lines changed

10 files changed

+90
-79
lines changed

‎llvm/lib/SYCLLowerIR/SpecConstants.cpp‎

Lines changed: 12 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -475,6 +475,7 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
475475
auto *FT = FunctionType::get(RetTy, ArgTys, false /*isVarArg*/);
476476
std::string FunctionName = mangleFuncItanium(BaseFunctionName, FT);
477477
Module *M = InsertBefore->getFunction()->getParent();
478+
bool IsSPIROrSPIRV = llvm::Triple(M->getTargetTriple()).isSPIROrSPIRV();
478479

479480
if (RetTy->isIntegerTy(1)) {
480481
assert(ArgTys.size() == 2 && "Expected a scalar spec constant");
@@ -500,6 +501,11 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
500501

501502
auto *Call =
502503
CallInst::Create(NewFT, NewFC.getCallee(), Args, "", InsertBefore);
504+
if (IsSPIROrSPIRV) {
505+
cast<Function>(NewFC.getCallee())
506+
->setCallingConv(CallingConv::SPIR_FUNC);
507+
Call->setCallingConv(CallingConv::SPIR_FUNC);
508+
}
503509
return CastInst::CreateTruncOrBitCast(Call, RetTy, "tobool",
504510
InsertBefore);
505511
}
@@ -520,7 +526,12 @@ Instruction *emitCall(Type *RetTy, StringRef BaseFunctionName,
520526
// types? Is it necessary?
521527

522528
FunctionCallee FC = M->getOrInsertFunction(FunctionName, FT);
523-
return CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
529+
auto *Call = CallInst::Create(FT, FC.getCallee(), Args, "", InsertBefore);
530+
if (IsSPIROrSPIRV) {
531+
cast<Function>(FC.getCallee())->setCallingConv(CallingConv::SPIR_FUNC);
532+
Call->setCallingConv(CallingConv::SPIR_FUNC);
533+
}
534+
return Call;
524535
}
525536

526537
Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,

‎llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer-array-of-arrays.ll‎

Lines changed: 34 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -39,44 +39,44 @@ define weak_odr spir_kernel void @_ZN2cl4sycl14kernel_handler33getSpecialization
3939
%9 = getelementptr inbounds %"class.cl::sycl::kernel_handler", %"class.cl::sycl::kernel_handler" addrspace(4)* %7, i32 0, i32 0
4040
%10 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %9, align 8, !tbaa !8
4141
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
42-
; CHECK-IR: %[[#NS0:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00)
43-
; CHECK-IR: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00)
44-
; CHECK-IR: %[[#NS2:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
45-
; CHECK-IR: %[[#NS3:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]])
46-
; CHECK-IR: %[[#NS4:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS3]])
47-
; CHECK-IR: %[[#NS5:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID3:]], float 0.000000e+00)
48-
; CHECK-IR: %[[#NS6:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
49-
; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
50-
; CHECK-IR: %[[#NS8:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS5]], float %[[#NS6]], float %[[#NS7]])
51-
; CHECK-IR: %[[#NS9:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS8]])
52-
; CHECK-IR: %[[#NS10:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
53-
; CHECK-IR: %[[#NS11:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID7:]], float 0.000000e+00)
54-
; CHECK-IR: %[[#NS12:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID8:]], float 0.000000e+00)
55-
; CHECK-IR: %[[#NS13:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS10]], float %[[#NS11]], float %[[#NS12]])
56-
; CHECK-IR: %[[#NS14:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS13]])
57-
; 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]])
58-
; CHECK-IR: %[[#NS16:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS15]])
42+
; CHECK-IR: %[[#NS0:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00)
43+
; CHECK-IR: %[[#NS1:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00)
44+
; CHECK-IR: %[[#NS2:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
45+
; CHECK-IR: %[[#NS3:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]])
46+
; CHECK-IR: %[[#NS4:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS3]])
47+
; CHECK-IR: %[[#NS5:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID3:]], float 0.000000e+00)
48+
; CHECK-IR: %[[#NS6:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
49+
; CHECK-IR: %[[#NS7:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
50+
; CHECK-IR: %[[#NS8:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS5]], float %[[#NS6]], float %[[#NS7]])
51+
; CHECK-IR: %[[#NS9:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS8]])
52+
; CHECK-IR: %[[#NS10:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
53+
; CHECK-IR: %[[#NS11:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID7:]], float 0.000000e+00)
54+
; CHECK-IR: %[[#NS12:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID8:]], float 0.000000e+00)
55+
; CHECK-IR: %[[#NS13:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS10]], float %[[#NS11]], float %[[#NS12]])
56+
; CHECK-IR: %[[#NS14:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS13]])
57+
; 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]])
58+
; 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]])
5959

6060
%11 = alloca %"class.std::array", align 4
6161
%12 = addrspacecast %"class.std::array"* %11 to %"class.std::array" addrspace(4)*
6262
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
63-
; CHECK-IR: %[[#NS17:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID9:]], float 0.000000e+00)
64-
; CHECK-IR: %[[#NS18:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID10:]], float 0.000000e+00)
65-
; CHECK-IR: %[[#NS19:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID11:]], float 0.000000e+00)
66-
; CHECK-IR: %[[#NS20:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS17]], float %[[#NS18]], float %[[#NS19]])
67-
; CHECK-IR: %[[#NS21:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS20]])
68-
; CHECK-IR: %[[#NS22:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID12:]], float 0.000000e+00)
69-
; CHECK-IR: %[[#NS23:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID13:]], float 1.000000e+00)
70-
; CHECK-IR: %[[#NS24:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID14:]], float 2.000000e+00)
71-
; CHECK-IR: %[[#NS25:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS22]], float %[[#NS23]], float %[[#NS24]])
72-
; CHECK-IR: %[[#NS26:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS25]])
73-
; CHECK-IR: %[[#NS27:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID15:]], float 0x4010666660000000)
74-
; CHECK-IR: %[[#NS28:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID16:]], float 0x4014666660000000)
75-
; CHECK-IR: %[[#NS29:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID17:]], float 0x4018CCCCC0000000)
76-
; CHECK-IR: %[[#NS30:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS27]], float %[[#NS28]], float %[[#NS29]])
77-
; CHECK-IR: %[[#NS31:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS30]])
78-
; 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]])
79-
; CHECK-IR: %[[#NS33:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS32]])
63+
; CHECK-IR: %[[#NS17:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID9:]], float 0.000000e+00)
64+
; CHECK-IR: %[[#NS18:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID10:]], float 0.000000e+00)
65+
; CHECK-IR: %[[#NS19:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID11:]], float 0.000000e+00)
66+
; CHECK-IR: %[[#NS20:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS17]], float %[[#NS18]], float %[[#NS19]])
67+
; CHECK-IR: %[[#NS21:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS20]])
68+
; CHECK-IR: %[[#NS22:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID12:]], float 0.000000e+00)
69+
; CHECK-IR: %[[#NS23:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID13:]], float 1.000000e+00)
70+
; CHECK-IR: %[[#NS24:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID14:]], float 2.000000e+00)
71+
; CHECK-IR: %[[#NS25:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS22]], float %[[#NS23]], float %[[#NS24]])
72+
; CHECK-IR: %[[#NS26:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS25]])
73+
; CHECK-IR: %[[#NS27:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID15:]], float 0x4010666660000000)
74+
; CHECK-IR: %[[#NS28:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID16:]], float 0x4014666660000000)
75+
; CHECK-IR: %[[#NS29:]] = call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID17:]], float 0x4018CCCCC0000000)
76+
; CHECK-IR: %[[#NS30:]] = call spir_func [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS27]], float %[[#NS28]], float %[[#NS29]])
77+
; CHECK-IR: %[[#NS31:]] = call spir_func %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS30]])
78+
; 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]])
79+
; 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]])
8080

8181
%13 = alloca %struct.coeff_str_t, align 8
8282
%14 = addrspacecast %struct.coeff_str_t* %13 to %struct.coeff_str_t addrspace(4)*

‎llvm/test/tools/sycl-post-link/spec-constants/SYCL-2020-zeroinitializer.ll‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -31,16 +31,16 @@ entry:
3131
%ref.tmp.i = alloca %struct._ZTS9composite.composite, align 8
3232
%ref.tmp.ascast.i = addrspacecast %struct._ZTS9composite.composite* %ref.tmp.i to %struct._ZTS9composite.composite addrspace(4)*
3333
%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
34-
; CHECK: call i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID0:]], i32 0)
34+
; CHECK: call spir_func i32 @_Z20__spirv_SpecConstantii(i32 [[#SCID0:]], i32 0)
3535

3636
%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
37-
; CHECK: call double @_Z20__spirv_SpecConstantid(i32 [[#SCID1:]], double 0.000000e+00)
37+
; CHECK: call spir_func double @_Z20__spirv_SpecConstantid(i32 [[#SCID1:]], double 0.000000e+00)
3838

3939
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
40-
; CHECK: call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
41-
; CHECK: call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID3:]], i8 0)
42-
; CHECK: call i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 0)
43-
; CHECK: call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID5:]], i64 0)
40+
; CHECK: call spir_func float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
41+
; CHECK: call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID3:]], i8 0)
42+
; CHECK: call spir_func i8 @_Z20__spirv_SpecConstantia(i32 [[#SCID4:]], i8 0)
43+
; CHECK: call spir_func i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID5:]], i64 0)
4444

4545
; CHECK-LOG: sycl.specialization-constants
4646
; CHECK-LOG:[[UNIQUE_PREFIX:[0-9a-zA-Z]+]]={0, 0, 4}

0 commit comments

Comments
 (0)