Skip to content

Commit fbf6e21

Browse files
Pavel Samolysovsteffenlarsen
Pavel Samolysov
andauthored
[sycl-post-link] Fix a crash during spec-constant properties generation (#5538)
The sycl-post-link tool crashed during processing an IR where a specialization constant with a padding is present, for example a specialization constant of such type: struct alignas(32) coeff_str_aligned_t {   std::array<float, 3> coeffs;   size_t number; }; In the IR, the constant looks like the following: @_ZL8coeff_id = internal addrspace(1) constant %"class.cl::sycl::specialization_id" { %struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [8 x i8] undef } }, align 32 That '[8 x i8] undef' led to the crash. This patch fixes the issue and adds the following optimization: if the default value of a specialization constant or a part of such default value is undefined, the value (or its part) is passed into an invocation of the '__spirv_SpecConstantComposite' function as is. Co-authored-by: Steffen Larsen <steffen.larsen@intel.com> Signed-off-by: Pavel Samolysov <pavel.samolysov@intel.com>
1 parent 64e92cb commit fbf6e21

File tree

2 files changed

+166
-41
lines changed

2 files changed

+166
-41
lines changed
Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
; RUN: sycl-post-link --spec-const=rt -S %s -o %t.files.table
2+
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
3+
; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP
4+
;
5+
; This test is intended to check that SpecConstantsPass is able to handle the
6+
; situation where specialization constants with complex types such as structs
7+
; have an 'undef' value for padding in LLVM IR
8+
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
%"class.cl::sycl::specialization_id" = type { %struct.coeff_str_aligned_t }
13+
%"class.cl::sycl::specialization_id.1" = type { %struct.coeff2_str_aligned_t }
14+
%struct.coeff_str_aligned_t = type { %"class.std::array", i64, [8 x i8] }
15+
%struct.coeff2_str_aligned_t = type { %"class.std::array", i64, [7 x i8], i8 }
16+
%"class.std::array" = type { [3 x float] }
17+
18+
$_ZTSZ4mainEUlN2cl4sycl14kernel_handlerEE_ = comdat any
19+
20+
@__usid_str = private unnamed_addr constant [32 x i8] c"ef880fa09cf7a9d7____ZL8coeff_id\00", align 1
21+
@_ZL8coeff_id = internal addrspace(1) constant %"class.cl::sycl::specialization_id" { %struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [8 x i8] undef } }, align 32
22+
@__usid_str.0 = private unnamed_addr constant [33 x i8] c"df991fa0adf9bad8____ZL8coeff_id2\00", align 1
23+
@_ZL8coeff_id2 = internal addrspace(1) constant %"class.cl::sycl::specialization_id.1" { %struct.coeff2_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [7 x i8] undef, i8 undef } }, align 32
24+
25+
; Function Attrs: convergent norecurse
26+
define weak_odr dso_local spir_kernel void @_ZTSZ4mainEUlN2cl4sycl14kernel_handlerEE_() local_unnamed_addr #0 comdat !kernel_arg_buffer_location !6 !sycl_kernel_omit_args !7 {
27+
%1 = alloca %struct.coeff_str_aligned_t, align 32
28+
%2 = addrspacecast %struct.coeff_str_aligned_t* %1 to %struct.coeff_str_aligned_t addrspace(4)*
29+
%3 = bitcast %struct.coeff_str_aligned_t* %1 to i8*
30+
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
31+
; CHECK-IR: %[[#NS0:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00)
32+
; CHECK-IR: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00)
33+
; CHECK-IR: %[[#NS2:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
34+
; CHECK-IR: %[[#NS3:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]])
35+
; CHECK-IR: %[[#NS4:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS3]])
36+
; CHECK-IR: %[[#NS5:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID3:]], i64 0)
37+
; 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)
38+
39+
%4 = alloca %struct.coeff2_str_aligned_t, align 32
40+
%5 = addrspacecast %struct.coeff2_str_aligned_t* %4 to %struct.coeff2_str_aligned_t addrspace(4)*
41+
%6 = bitcast %struct.coeff2_str_aligned_t* %4 to i8*
42+
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
43+
; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
44+
; CHECK-IR: %[[#NS8:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
45+
; CHECK-IR: %[[#NS9:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
46+
; CHECK-IR: %[[#NS10:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS7]], float %[[#NS8]], float %[[#NS9]])
47+
; CHECK-IR: %[[#NS11:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS10]])
48+
; CHECK-IR: %[[#NS12:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID7:]], i64 0)
49+
; 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)
50+
51+
ret void
52+
}
53+
; Function Attrs: convergent
54+
declare dso_local 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, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef) local_unnamed_addr #2
55+
56+
declare dso_local 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, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef) local_unnamed_addr #2
57+
58+
attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="spec-constant-test.cpp" "uniform-work-group-size"="true" }
59+
attributes #1 = { argmemonly mustprogress nofree nosync nounwind willreturn }
60+
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
61+
attributes #3 = { nounwind }
62+
attributes #4 = { convergent }
63+
64+
!llvm.dependent-libraries = !{!0}
65+
!llvm.module.flags = !{!1, !2}
66+
!opencl.spir.version = !{!3}
67+
!spirv.Source = !{!4}
68+
!llvm.ident = !{!5}
69+
; CHECK-IR: !sycl.specialization-constants = !{![[#MN0:]], ![[#MN1:]]}
70+
; CHECK-IR: !sycl.specialization-constants-default-values = !{![[#MN2:]], ![[#MN3:]]}
71+
72+
!0 = !{!"libcpmt"}
73+
!1 = !{i32 1, !"wchar_size", i32 2}
74+
!2 = !{i32 7, !"frame-pointer", i32 2}
75+
!3 = !{i32 1, i32 2}
76+
!4 = !{i32 4, i32 100000}
77+
!5 = !{!"clang version 14.0.0"}
78+
!6 = !{i32 -1}
79+
!7 = !{i1 true}
80+
; CHECK-IR: ![[#MN0]] = !{!"ef880fa09cf7a9d7____ZL8coeff_id", i32 0, i32 0, i32 4, i32 1, i32 4, i32 4, i32 2, i32 8, i32 4, i32 3, i32 16, i32 8, i32 -1, i32 24, i32 8}
81+
; CHECK-IR: ![[#MN1]] = !{!"df991fa0adf9bad8____ZL8coeff_id2", i32 5, i32 0, i32 4, i32 6, i32 4, i32 4, i32 7, i32 8, i32 4, i32 8, i32 16, i32 8, i32 -1, i32 31, i32 1}
82+
; CHECK-IR: ![[#MN2]] = !{%struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [8 x i8] undef }}
83+
; CHECK-IR: ![[#MN3]] = !{%struct.coeff2_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [7 x i8] undef, i8 undef }}
84+
85+
; CHECK-PROP: [SYCL/specialization constants]
86+
; CHECK-PROP-NEXT: ef880fa09cf7a9d7____ZL8coeff_id=2|
87+
88+
; CHECK-PROP: [SYCL/specialization constants default values]
89+
; CHECK-PROP-NEXT: all=2|

llvm/tools/sycl-post-link/SpecConstants.cpp

Lines changed: 77 additions & 41 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,14 @@ constexpr char SPEC_CONST_MD_STRING[] = "sycl.specialization-constants";
5151
constexpr char SPEC_CONST_DEFAULT_VAL_MD_STRING[] =
5252
"sycl.specialization-constants-default-values";
5353

54+
/// Spec. Constant ID is a pair of Id and a flag whether this Id belongs to an
55+
/// undefined value. Undefined values ('undef' in the IR) are used to get the
56+
/// required alignment and should be handled in a special manner as padding.
57+
struct ID {
58+
unsigned ID;
59+
bool Undef;
60+
};
61+
5462
StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo,
5563
SmallVectorImpl<Instruction *> &DelInsts) {
5664
Value *V = CI->getArgOperand(ArgNo)->stripPointerCasts();
@@ -236,8 +244,13 @@ MDNode *generateSpecConstDefaultValueMetadata(StringRef SymID, Value *Default) {
236244
/// Recursively iterates over a composite type in order to collect information
237245
/// about its scalar elements.
238246
void collectCompositeElementsInfoRecursive(
239-
const Module &M, Type *Ty, const unsigned *&IDIter, unsigned &Offset,
247+
const Module &M, Type *Ty, const ID *&IDIter, unsigned &Offset,
240248
std::vector<SpecConstantDescriptor> &Result) {
249+
if (IDIter->Undef) {
250+
// We can just skip undefined values because every such value is just a
251+
// padding and will be handled in a different manner.
252+
return;
253+
}
241254
if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
242255
for (size_t I = 0; I < ArrTy->getNumElements(); ++I) {
243256
// TODO: this is a spot for potential optimization: for arrays we could
@@ -246,7 +259,9 @@ void collectCompositeElementsInfoRecursive(
246259
collectCompositeElementsInfoRecursive(M, ArrTy->getElementType(), IDIter,
247260
Offset, Result);
248261
}
249-
} else if (auto *StructTy = dyn_cast<StructType>(Ty)) {
262+
return;
263+
}
264+
if (auto *StructTy = dyn_cast<StructType>(Ty)) {
250265
const StructLayout *SL = M.getDataLayout().getStructLayout(StructTy);
251266
const unsigned BaseOffset = Offset;
252267
unsigned LocalOffset = Offset;
@@ -267,7 +282,12 @@ void collectCompositeElementsInfoRecursive(
267282
BaseOffset + SL->getSizeInBytes() - LocalOffset;
268283
if (PostStructPadding > 0) {
269284
SpecConstantDescriptor Desc;
270-
// ID of padding descriptors is the max value possible.
285+
// ID of padding descriptors is the max value possible. This value is a
286+
// magic value for the runtime and will just be skipped. Even if there
287+
// are many specialization constants and every constant has padding of
288+
// a different length, everything will work regardless rewriting
289+
// the descriptions with Desc.ID equals to the max value: they will just
290+
// be ignored at all.
271291
Desc.ID = std::numeric_limits<unsigned>::max();
272292
Desc.Offset = LocalOffset;
273293
Desc.Size = PostStructPadding;
@@ -277,25 +297,29 @@ void collectCompositeElementsInfoRecursive(
277297
// Update "global" offset according to the total size of a handled struct
278298
// type.
279299
Offset += SL->getSizeInBytes();
280-
} else if (auto *VecTy = dyn_cast<FixedVectorType>(Ty)) {
300+
return;
301+
}
302+
if (auto *VecTy = dyn_cast<FixedVectorType>(Ty)) {
281303
for (size_t I = 0; I < VecTy->getNumElements(); ++I) {
282304
// TODO: this is a spot for potential optimization: for vectors we could
283305
// just make a single recursive call here and use it to populate Result
284306
// in a loop.
285307
collectCompositeElementsInfoRecursive(M, VecTy->getElementType(), IDIter,
286308
Offset, Result);
287309
}
288-
} else { // Assume that we encountered some scalar element
289-
SpecConstantDescriptor Desc;
290-
Desc.ID = *IDIter;
291-
Desc.Offset = Offset;
292-
Desc.Size = M.getDataLayout().getTypeStoreSize(Ty);
293-
Result.push_back(Desc);
294-
295-
// Move current ID and offset
296-
++IDIter;
297-
Offset += Desc.Size;
310+
return;
298311
}
312+
313+
// Assume that we encountered some scalar element
314+
SpecConstantDescriptor Desc;
315+
Desc.ID = IDIter->ID;
316+
Desc.Offset = Offset;
317+
Desc.Size = M.getDataLayout().getTypeStoreSize(Ty);
318+
Result.push_back(Desc);
319+
320+
// Move current ID and offset
321+
++IDIter;
322+
Offset += Desc.Size;
299323
}
300324

301325
/// Recursively iterates over a composite type in order to collect information
@@ -306,8 +330,8 @@ void collectCompositeElementsInfoRecursive(
306330
void collectCompositeElementsDefaultValuesRecursive(
307331
const Module &M, Constant *C, unsigned &Offset,
308332
std::vector<char> &DefaultValues) {
309-
if (isa<ConstantAggregateZero>(C)) {
310-
// This code is generic for zeroinitializer for both arrays and structs
333+
if (isa<ConstantAggregateZero>(C) || isa<UndefValue>(C)) {
334+
// This code is generic for both arrays and structs
311335
size_t NumBytes = M.getDataLayout().getTypeStoreSize(C->getType());
312336
std::fill_n(std::back_inserter(DefaultValues), NumBytes, 0);
313337
Offset += NumBytes;
@@ -400,7 +424,7 @@ void collectCompositeElementsDefaultValuesRecursive(
400424
}
401425

402426
MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID,
403-
Type *SCTy, ArrayRef<unsigned> IDs,
427+
Type *SCTy, ArrayRef<ID> IDs,
404428
bool IsNativeSpecConstant) {
405429
SmallVector<Metadata *, 16> MDOps;
406430
LLVMContext &Ctx = M.getContext();
@@ -413,7 +437,7 @@ MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID,
413437
std::vector<SpecConstantDescriptor> Result;
414438
Result.reserve(IDs.size());
415439
unsigned Offset = 0;
416-
const unsigned *IDPtr = IDs.data();
440+
const ID *IDPtr = IDs.data();
417441
collectCompositeElementsInfoRecursive(M, SCTy, IDPtr, Offset, Result);
418442

419443
// We may have padding elements so size should be at least the same size as
@@ -432,7 +456,7 @@ MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID,
432456
assert(IDs.size() == 1 &&
433457
"There must be a single ID for emulated spec constant");
434458
MDOps.push_back(ConstantAsMetadata::get(
435-
Constant::getIntegerValue(Int32Ty, APInt(32, IDs[0]))));
459+
Constant::getIntegerValue(Int32Ty, APInt(32, IDs[0].ID))));
436460
// Second element is always zero here
437461
MDOps.push_back(ConstantAsMetadata::get(
438462
Constant::getIntegerValue(Int32Ty, APInt(32, 0))));
@@ -519,14 +543,9 @@ Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,
519543
return emitCall(Ty, SPIRV_GET_SPEC_CONST_VAL, Args, InsertBefore);
520544
}
521545

522-
Instruction *emitSpecConstantComposite(Type *Ty,
523-
ArrayRef<Instruction *> Elements,
546+
Instruction *emitSpecConstantComposite(Type *Ty, ArrayRef<Value *> Elements,
524547
Instruction *InsertBefore) {
525-
SmallVector<Value *, 8> Args(Elements.size());
526-
for (unsigned I = 0; I < Elements.size(); ++I) {
527-
Args[I] = cast<Value>(Elements[I]);
528-
}
529-
return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Args, InsertBefore);
548+
return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Elements, InsertBefore);
530549
}
531550

532551
/// For specified specialization constant type emits LLVM IR which is required
@@ -553,28 +572,46 @@ Instruction *emitSpecConstantComposite(Type *Ty,
553572
/// composite (plus for the top-level composite). Also enumerates all
554573
/// encountered scalars and assigns them IDs (or re-uses existing ones).
555574
Instruction *emitSpecConstantRecursiveImpl(Type *Ty, Instruction *InsertBefore,
556-
SmallVectorImpl<unsigned> &IDs,
575+
SmallVectorImpl<ID> &IDs,
557576
unsigned &Index,
558577
Constant *DefaultValue) {
559578
if (!Ty->isArrayTy() && !Ty->isStructTy() && !Ty->isVectorTy()) { // Scalar
560579
if (Index >= IDs.size()) {
561580
// If it is a new specialization constant, we need to generate IDs for
562581
// scalar elements, starting with the second one.
563-
IDs.push_back(IDs.back() + 1);
582+
assert(!isa_and_nonnull<UndefValue>(DefaultValue) &&
583+
"All scalar values should be defined");
584+
IDs.push_back({IDs.back().ID + 1, false});
564585
}
565-
return emitSpecConstant(IDs[Index++], Ty, InsertBefore, DefaultValue);
586+
return emitSpecConstant(IDs[Index++].ID, Ty, InsertBefore, DefaultValue);
566587
}
567588

568-
SmallVector<Instruction *, 8> Elements;
589+
SmallVector<Value *, 8> Elements;
590+
auto HandleUndef = [&](Constant *Def) {
591+
if (Index >= IDs.size()) {
592+
// If it is a new specialization constant, we need to generate IDs for
593+
// the whole undef value.
594+
IDs.push_back({IDs.back().ID + 1, true});
595+
}
596+
Elements.push_back(Def);
597+
};
569598
auto LoopIteration = [&](Type *Ty, unsigned LocalIndex) {
570599
// Select corresponding element of the default value if it was provided
571600
Constant *Def =
572601
DefaultValue ? DefaultValue->getAggregateElement(LocalIndex) : nullptr;
573-
Elements.push_back(
574-
emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index, Def));
602+
if (isa_and_nonnull<UndefValue>(Def))
603+
HandleUndef(Def);
604+
else
605+
Elements.push_back(
606+
emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index, Def));
575607
};
576608

577-
if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
609+
if (isa_and_nonnull<UndefValue>(DefaultValue)) {
610+
// If the default value is a composite and has the value 'undef', we should
611+
// not generate a bunch of __spirv_SpecConstant for its elements but
612+
// pass it into __spirv_SpecConstantComposite as is.
613+
HandleUndef(DefaultValue);
614+
} else if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
578615
for (size_t I = 0; I < ArrTy->getNumElements(); ++I) {
579616
LoopIteration(ArrTy->getElementType(), I);
580617
}
@@ -596,7 +633,7 @@ Instruction *emitSpecConstantRecursiveImpl(Type *Ty, Instruction *InsertBefore,
596633

597634
/// Wrapper intended to hide IsFirstElement argument from the caller
598635
Instruction *emitSpecConstantRecursive(Type *Ty, Instruction *InsertBefore,
599-
SmallVectorImpl<unsigned> &IDs,
636+
SmallVectorImpl<ID> &IDs,
600637
Constant *DefaultValue) {
601638
unsigned Index = 0;
602639
return emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index,
@@ -607,9 +644,9 @@ Instruction *emitSpecConstantRecursive(Type *Ty, Instruction *InsertBefore,
607644

608645
PreservedAnalyses SpecConstantsPass::run(Module &M,
609646
ModuleAnalysisManager &MAM) {
610-
unsigned NextID = 0;
647+
ID NextID = {0, false};
611648
unsigned NextOffset = 0;
612-
StringMap<SmallVector<unsigned, 1>> IDMap;
649+
StringMap<SmallVector<ID, 1>> IDMap;
613650
StringMap<unsigned> OffsetMap;
614651
MapVector<StringRef, MDNode *> SCMetadata;
615652
SmallVector<MDNode *, 4> DefaultsMetadata;
@@ -690,9 +727,8 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
690727
if (SetValAtRT) {
691728
// 2. Spec constant value will be set at run time - then add the literal
692729
// to a "spec const string literal ID" -> "vector of integer IDs" map,
693-
// uniquing the integer IDs if this is a new literal
694-
auto Ins =
695-
IDMap.insert(std::make_pair(SymID, SmallVector<unsigned, 1>{}));
730+
// making the integer IDs unique if this is a new literal
731+
auto Ins = IDMap.insert(std::make_pair(SymID, SmallVector<ID, 1>{}));
696732
IsNewSpecConstant = Ins.second;
697733
auto &IDs = Ins.first->second;
698734
if (IsNewSpecConstant) {
@@ -708,7 +744,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
708744
// emitSpecConstantRecursive might emit more than one spec constant
709745
// (because of composite types) and therefore, we need to adjust
710746
// NextID according to the actual amount of emitted spec constants.
711-
NextID += IDs.size();
747+
NextID.ID += IDs.size();
712748

713749
// Generate necessary metadata which later will be pulled by
714750
// sycl-post-link and transformed into device image properties
@@ -740,7 +776,7 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
740776
SCMetadata[SymID] = generateSpecConstantMetadata(
741777
M, SymID, SCTy, NextID, /* is native spec constant */ false);
742778

743-
++NextID;
779+
++NextID.ID;
744780
NextOffset += Size;
745781
}
746782

0 commit comments

Comments
 (0)