Skip to content

Commit eec22ed

Browse files
author
Pavel Samolysov
authored
[sycl-post-link] Add handling for zero-initialized array constants (#5321)
The problem is the following, when the `getOperand` method is invoked for a `zeroinitializer` array: `C->getOperand(i)`, the assert from `User.h` was triggered: ``` assert(i < NumUserOperands && "getOperand() out of range!"); ``` This is because `zeroinitializer` has no operands. The problem occurs only when the specialization constant was not initialized directly in the point of definition. When a specialization constant is defined in the IR as `zeroinitializer`, its default value is just an array of zeros of the arbitrary size. No recursive calls of the `collectCompositeElementsDefaultValuesRecursive` function is required in the case. Also, to make processing constants with different types more robust, dynamic casting of the constant itself (`dyn_cast<ConstantType>(C)`) should be used instead of casting the constant's type since the type can be inconsistent with the constant behavior, for example not every constant with the type equals to `ArrayType` has any components (the `zeroinitializer` has not).
1 parent 2359e81 commit eec22ed

File tree

2 files changed

+168
-41
lines changed

2 files changed

+168
-41
lines changed

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

Lines changed: 116 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,116 @@
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 arrays
7+
; within arrays have zeroinitializer 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 { %"class.std::array" }
13+
%"class.std::array" = type { [3 x %"class.std::array.1"] }
14+
%"class.std::array.1" = type { [3 x float] }
15+
%"class.cl::sycl::kernel_handler" = type { i8 addrspace(4)* }
16+
%"class.cl::sycl::specialization_id.1" = type { %struct.coeff_str_t }
17+
%struct.coeff_str_t = type { %"class.std::array.1", i64 }
18+
19+
@__usid_str.1 = private unnamed_addr constant [32 x i8] c"9f47062a80eecfa7____ZL8coeff_id\00", align 1
20+
@_ZL8coeff_id = internal addrspace(1) constant %"class.cl::sycl::specialization_id" zeroinitializer, align 4
21+
22+
@__usid_str.2 = private unnamed_addr constant [33 x i8] c"405761736d5a1797____ZL9coeff_id2\00", align 1
23+
@_ZL9coeff_id2 = internal addrspace(1) constant %"class.cl::sycl::specialization_id" { %"class.std::array" { [3 x %"class.std::array.1"] [%"class.std::array.1" zeroinitializer, %"class.std::array.1" { [3 x float] [float 0.000000e+00, float 1.000000e+00, float 2.000000e+00] }, %"class.std::array.1" { [3 x float] [float 0x4010666660000000, float 0x4014666660000000, float 0x4018CCCCC0000000] }] } }, align 4
24+
25+
@__usid_str.3 = private unnamed_addr constant [33 x i8] c"6da74a122db9f35d____ZL9coeff_id3\00", align 1
26+
@_ZL9coeff_id3 = internal addrspace(1) constant %"class.cl::sycl::specialization_id.1" zeroinitializer, align 8
27+
28+
; Function Attrs: convergent mustprogress norecurse
29+
define internal spir_func void @_ZN2cl4sycl14kernel_handler33getSpecializationConstantOnDeviceIL_ZL8coeff_idESt5arrayIS3_IfLy3EELy3EELPv0EEET0_v(%"class.std::array" addrspace(4)* noalias sret(%"class.std::array") align 4 %0, %"class.cl::sycl::kernel_handler" addrspace(4)* align 8 dereferenceable_or_null(8) %1) #0 align 2 {
30+
%3 = alloca %"class.cl::sycl::kernel_handler" addrspace(4)*, align 8
31+
%4 = alloca i8 addrspace(4)*, align 8
32+
%5 = addrspacecast %"class.cl::sycl::kernel_handler" addrspace(4)** %3 to %"class.cl::sycl::kernel_handler" addrspace(4)* addrspace(4)*
33+
%6 = addrspacecast i8 addrspace(4)** %4 to i8 addrspace(4)* addrspace(4)*
34+
store %"class.cl::sycl::kernel_handler" addrspace(4)* %1, %"class.cl::sycl::kernel_handler" addrspace(4)* addrspace(4)* %5, align 8, !tbaa !4
35+
%7 = load %"class.cl::sycl::kernel_handler" addrspace(4)*, %"class.cl::sycl::kernel_handler" addrspace(4)* addrspace(4)* %5, align 8
36+
store i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([32 x i8], [32 x i8]* @__usid_str.1, i32 0, i32 0) to i8 addrspace(4)*), i8 addrspace(4)* addrspace(4)* %6, align 8, !tbaa !4
37+
%8 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %6, align 8, !tbaa !4
38+
%9 = getelementptr inbounds %"class.cl::sycl::kernel_handler", %"class.cl::sycl::kernel_handler" addrspace(4)* %7, i32 0, i32 0
39+
%10 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %9, align 8, !tbaa !8
40+
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
41+
; CHECK-IR: %[[#NS0:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00)
42+
; CHECK-IR: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00)
43+
; CHECK-IR: %[[#NS2:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
44+
; CHECK-IR: %[[#NS3:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]])
45+
; CHECK-IR: %[[#NS4:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS3]])
46+
; CHECK-IR: %[[#NS5:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID3:]], float 0.000000e+00)
47+
; CHECK-IR: %[[#NS6:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
48+
; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
49+
; CHECK-IR: %[[#NS8:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS5]], float %[[#NS6]], float %[[#NS7]])
50+
; CHECK-IR: %[[#NS9:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS8]])
51+
; CHECK-IR: %[[#NS10:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
52+
; CHECK-IR: %[[#NS11:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID7:]], float 0.000000e+00)
53+
; CHECK-IR: %[[#NS12:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID8:]], float 0.000000e+00)
54+
; CHECK-IR: %[[#NS13:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS10]], float %[[#NS11]], float %[[#NS12]])
55+
; CHECK-IR: %[[#NS14:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS13]])
56+
; 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]])
57+
; CHECK-IR: %[[#NS16:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS15]])
58+
59+
%11 = alloca %"class.std::array", align 4
60+
%12 = addrspacecast %"class.std::array"* %11 to %"class.std::array" addrspace(4)*
61+
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
62+
; CHECK-IR: %[[#NS17:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID9:]], float 0.000000e+00)
63+
; CHECK-IR: %[[#NS18:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID10:]], float 0.000000e+00)
64+
; CHECK-IR: %[[#NS19:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID11:]], float 0.000000e+00)
65+
; CHECK-IR: %[[#NS20:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS17]], float %[[#NS18]], float %[[#NS19]])
66+
; CHECK-IR: %[[#NS21:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS20]])
67+
; CHECK-IR: %[[#NS22:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID12:]], float 0.000000e+00)
68+
; CHECK-IR: %[[#NS23:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID13:]], float 1.000000e+00)
69+
; CHECK-IR: %[[#NS24:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID14:]], float 2.000000e+00)
70+
; CHECK-IR: %[[#NS25:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS22]], float %[[#NS23]], float %[[#NS24]])
71+
; CHECK-IR: %[[#NS26:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS25]])
72+
; CHECK-IR: %[[#NS27:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID15:]], float 0x4010666660000000)
73+
; CHECK-IR: %[[#NS28:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID16:]], float 0x4014666660000000)
74+
; CHECK-IR: %[[#NS29:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID17:]], float 0x4018CCCCC0000000)
75+
; CHECK-IR: %[[#NS30:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS27]], float %[[#NS28]], float %[[#NS29]])
76+
; CHECK-IR: %[[#NS31:]] = call %"class.std::array.1" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array.1"([3 x float] %[[#NS30]])
77+
; 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]])
78+
; CHECK-IR: %[[#NS33:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_class.std::array.1_Rclass.std::array"([3 x %"class.std::array.1"] %[[#NS32]])
79+
80+
%13 = alloca %struct.coeff_str_t, align 8
81+
%14 = addrspacecast %struct.coeff_str_t* %13 to %struct.coeff_str_t addrspace(4)*
82+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI11coeff_str_tET_PKcPKvS5_(%struct.coeff_str_t addrspace(4)* sret(%struct.coeff_str_t) align 8 %14, i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([33 x i8], [33 x i8]* @__usid_str.3, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id.1" addrspace(1)* @_ZL9coeff_id3 to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* null) #13
83+
ret void
84+
}
85+
86+
; Function Attrs: convergent
87+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueISt5arrayIS0_IfLy3EELy3EEET_PKcPKvS7_(%"class.std::array" addrspace(4)* sret(%"class.std::array") align 4, i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) #1
88+
89+
; Function Attrs: convergent
90+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI11coeff_str_tET_PKcPKvS5_(%struct.coeff_str_t addrspace(4)* sret(%struct.coeff_str_t) align 8, i8 addrspace(4)*, i8 addrspace(4)*, i8 addrspace(4)*) local_unnamed_addr #2
91+
92+
attributes #0 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
93+
attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
94+
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
95+
96+
!opencl.spir.version = !{!1}
97+
!spirv.Source = !{!2}
98+
!llvm.ident = !{!3}
99+
100+
!1 = !{i32 1, i32 2}
101+
!2 = !{i32 4, i32 100000}
102+
!3 = !{!"clang version 14.0.0"}
103+
!4 = !{!5, !5, i64 0}
104+
!5 = !{!"any pointer", !6, i64 0}
105+
!6 = !{!"omnipotent char", !7, i64 0}
106+
!7 = !{!"Simple C++ TBAA"}
107+
!8 = !{!9, !5, i64 0}
108+
!9 = !{!"_ZTSN2cl4sycl14kernel_handlerE", !5, i64 0}
109+
110+
; CHECK-PROP: [SYCL/specialization constants]
111+
; CHECK-PROP-NEXT: 9f47062a80eecfa7____ZL8coeff_id=2|gNAAAAAAAAAAAAAAAAAAAQAAAAQAAAAAEAAAAQAAAAgAAAAAIAAAAQAAAAwAAAAAMAAAAQAAAAABAAAAQAAAAQAAAAQBAAAAUAAAAQAAAAgBAAAAYAAAAQAAAAwBAAAAcAAAAQAAAAACAAAAgAAAAQAAAAA
112+
; CHECK-PROP-NEXT: 405761736d5a1797____ZL9coeff_id2=2|gNAAAAAAAAQCAAAAAAAAAQAAAAgCAAAAEAAAAQAAAAwCAAAAIAAAAQAAAAADAAAAMAAAAQAAAAQDAAAAQAAAAQAAAAgDAAAAUAAAAQAAAAwDAAAAYAAAAQAAAAAEAAAAcAAAAQAAAAQEAAAAgAAAAQAAAAA
113+
; CHECK-PROP-NEXT: 6da74a122db9f35d____ZL9coeff_id3=2|AGAAAAAAAAgEAAAAAAAAAQAAAAwEAAAAEAAAAQAAAAAFAAAAIAAAAQAAAAQFAAAAQAAAAgAAAAA
114+
115+
; CHECK-PROP: [SYCL/specialization constants default values]
116+
; CHECK-PROP-NEXT: all=2|AMAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAg/AAAAA0MzMIQzMzoAZmZGDEAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA

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

Lines changed: 52 additions & 41 deletions
Original file line numberDiff line numberDiff line change
@@ -293,30 +293,40 @@ void collectCompositeElementsInfoRecursive(
293293
void collectCompositeElementsDefaultValuesRecursive(
294294
const Module &M, Constant *C, unsigned &Offset,
295295
std::vector<char> &DefaultValues) {
296-
Type *Ty = C->getType();
296+
if (isa<ConstantAggregateZero>(C)) {
297+
// This code is generic for zeroinitializer for both arrays and structs
298+
size_t NumBytes = M.getDataLayout().getTypeStoreSize(C->getType());
299+
std::fill_n(std::back_inserter(DefaultValues), NumBytes, 0);
300+
Offset += NumBytes;
301+
return;
302+
}
303+
297304
if (auto *DataSeqC = dyn_cast<ConstantDataSequential>(C)) {
298305
// This code is generic for both vectors and arrays of scalars
299306
for (size_t I = 0; I < DataSeqC->getNumElements(); ++I) {
300307
Constant *El = cast<Constant>(DataSeqC->getElementAsConstant(I));
301308
collectCompositeElementsDefaultValuesRecursive(M, El, Offset,
302309
DefaultValues);
303310
}
304-
} else if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
311+
return;
312+
}
313+
314+
if (auto *ArrayC = dyn_cast<ConstantArray>(C)) {
305315
// This branch handles arrays of composite types (structs, arrays, etc.)
306-
for (size_t I = 0; I < ArrTy->getNumElements(); ++I) {
307-
Constant *El = cast<Constant>(C->getOperand(I));
308-
collectCompositeElementsDefaultValuesRecursive(M, El, Offset,
309-
DefaultValues);
316+
assert(!C->isZeroValue() && "C must not be a zeroinitializer");
317+
for (size_t I = 0; I < ArrayC->getType()->getNumElements(); ++I) {
318+
collectCompositeElementsDefaultValuesRecursive(M, ArrayC->getOperand(I),
319+
Offset, DefaultValues);
310320
}
311-
} else if (auto *StructTy = dyn_cast<StructType>(Ty)) {
321+
return;
322+
}
323+
324+
if (auto *StructC = dyn_cast<ConstantStruct>(C)) {
325+
assert(!C->isZeroValue() && "C must not be a zeroinitializer");
326+
auto *StructTy = StructC->getType();
312327
const StructLayout *SL = M.getDataLayout().getStructLayout(StructTy);
313328
const size_t BaseDefaultValueOffset = DefaultValues.size();
314329
for (size_t I = 0, E = StructTy->getNumElements(); I < E; ++I) {
315-
Constant *El = nullptr;
316-
if (C->isZeroValue())
317-
El = Constant::getNullValue(StructTy->getElementType(I));
318-
else
319-
El = cast<Constant>(C->getOperand(I));
320330
// When handling elements of a structure, we do not use manually
321331
// calculated offsets (which are sum of sizes of all previously
322332
// encountered elements), but instead rely on data provided for us by
@@ -328,8 +338,8 @@ void collectCompositeElementsDefaultValuesRecursive(
328338
while (LocalOffset != DefaultValues.size())
329339
DefaultValues.push_back(0);
330340

331-
collectCompositeElementsDefaultValuesRecursive(M, El, LocalOffset,
332-
DefaultValues);
341+
collectCompositeElementsDefaultValuesRecursive(
342+
M, StructC->getOperand(I), LocalOffset, DefaultValues);
333343
}
334344
const size_t SLSize = SL->getSizeInBytes();
335345

@@ -341,38 +351,39 @@ void collectCompositeElementsDefaultValuesRecursive(
341351
// Update "global" offset according to the total size of a handled struct
342352
// type.
343353
Offset += SLSize;
344-
} else { // Assume that we encountered some scalar element
345-
int NumBytes = M.getDataLayout().getTypeStoreSize(Ty);
354+
return;
355+
}
346356

347-
if (auto IntConst = dyn_cast<ConstantInt>(C)) {
348-
auto Val = IntConst->getValue().getZExtValue();
349-
std::copy_n(reinterpret_cast<char *>(&Val), NumBytes,
357+
// Assume that we encountered some scalar element
358+
size_t NumBytes = M.getDataLayout().getTypeStoreSize(C->getType());
359+
if (auto IntConst = dyn_cast<ConstantInt>(C)) {
360+
auto Val = IntConst->getValue().getZExtValue();
361+
std::copy_n(reinterpret_cast<char *>(&Val), NumBytes,
362+
std::back_inserter(DefaultValues));
363+
} else if (auto FPConst = dyn_cast<ConstantFP>(C)) {
364+
auto Val = FPConst->getValue();
365+
366+
if (NumBytes == 2) {
367+
auto IVal = Val.bitcastToAPInt();
368+
assert(IVal.getBitWidth() == 16);
369+
auto Storage = static_cast<uint16_t>(IVal.getZExtValue());
370+
std::copy_n(reinterpret_cast<char *>(&Storage), NumBytes,
371+
std::back_inserter(DefaultValues));
372+
} else if (NumBytes == 4) {
373+
float v = Val.convertToFloat();
374+
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
375+
std::back_inserter(DefaultValues));
376+
} else if (NumBytes == 8) {
377+
double v = Val.convertToDouble();
378+
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
350379
std::back_inserter(DefaultValues));
351-
} else if (auto FPConst = dyn_cast<ConstantFP>(C)) {
352-
auto Val = FPConst->getValue();
353-
354-
if (NumBytes == 2) {
355-
auto IVal = Val.bitcastToAPInt();
356-
assert(IVal.getBitWidth() == 16);
357-
auto Storage = static_cast<uint16_t>(IVal.getZExtValue());
358-
std::copy_n(reinterpret_cast<char *>(&Storage), NumBytes,
359-
std::back_inserter(DefaultValues));
360-
} else if (NumBytes == 4) {
361-
float v = Val.convertToFloat();
362-
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
363-
std::back_inserter(DefaultValues));
364-
} else if (NumBytes == 8) {
365-
double v = Val.convertToDouble();
366-
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
367-
std::back_inserter(DefaultValues));
368-
} else {
369-
llvm_unreachable("Unexpected constant floating point type");
370-
}
371380
} else {
372-
llvm_unreachable("Unexpected constant scalar type");
381+
llvm_unreachable("Unexpected constant floating point type");
373382
}
374-
Offset += NumBytes;
383+
} else {
384+
llvm_unreachable("Unexpected constant scalar type");
375385
}
386+
Offset += NumBytes;
376387
}
377388

378389
MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID,

0 commit comments

Comments
 (0)