Skip to content

[sycl-post-link] Add handling for zero-initialized array constants #5321

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 11 commits into from Jan 20, 2022
Merged

[sycl-post-link] Add handling for zero-initialized array constants #5321

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
11 commits
Select commit Hold shift + click to select a range
7bfc53c
[sycl-post-link] Add handling for zero-initialized array constants
Jan 17, 2022
47e3837
[sycl-post-link] Add a lit test to check array-of-arrays handling
Jan 17, 2022
3875a0f
[sycl-post-link] Simplify the lit test
Jan 18, 2022
9687d41
[sycl-post-link] Simplify the lit test again
Jan 18, 2022
6d80223
[sycl-post-link] Add a lit test to check a partially initialized array
Jan 18, 2022
a84ae51
[sycl-post-link] Add a lit test to check a structure with an array
Jan 18, 2022
320b83a
[sycl-post-link] Use dynamic casting to Constant subtypes
Jan 18, 2022
1708dae
[sycl-post-link][NFC] Fix formatting issue
Jan 18, 2022
662aeba
[sycl-post-link] Using early exit to make code more readable
Jan 19, 2022
a5a3f85
[sycl-post-link][NFC] Fix formatting issues
Jan 19, 2022
657aa24
[sycl-post-link] Clean up a bit the RUN lines for the lit test
Jan 19, 2022
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
116 changes: 116 additions & 0 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
@@ -0,0 +1,116 @@
; RUN: sycl-post-link --spec-const=rt -S %s -o %t.files.table
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP
;
; This test is intended to check that SpecConstantsPass is able to handle the
; situation where specialization constants with complex types such as arrays
; within arrays have zeroinitializer in LLVM IR

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"
target triple = "spir64-unknown-unknown"

%"class.cl::sycl::specialization_id" = type { %"class.std::array" }
%"class.std::array" = type { [3 x %"class.std::array.1"] }
%"class.std::array.1" = type { [3 x float] }
%"class.cl::sycl::kernel_handler" = type { i8 addrspace(4)* }
%"class.cl::sycl::specialization_id.1" = type { %struct.coeff_str_t }
%struct.coeff_str_t = type { %"class.std::array.1", i64 }

@__usid_str.1 = private unnamed_addr constant [32 x i8] c"9f47062a80eecfa7____ZL8coeff_id\00", align 1
@_ZL8coeff_id = internal addrspace(1) constant %"class.cl::sycl::specialization_id" zeroinitializer, align 4

@__usid_str.2 = private unnamed_addr constant [33 x i8] c"405761736d5a1797____ZL9coeff_id2\00", align 1
@_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

@__usid_str.3 = private unnamed_addr constant [33 x i8] c"6da74a122db9f35d____ZL9coeff_id3\00", align 1
@_ZL9coeff_id3 = internal addrspace(1) constant %"class.cl::sycl::specialization_id.1" zeroinitializer, align 8

; Function Attrs: convergent mustprogress norecurse
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 {
%3 = alloca %"class.cl::sycl::kernel_handler" addrspace(4)*, align 8
%4 = alloca i8 addrspace(4)*, align 8
%5 = addrspacecast %"class.cl::sycl::kernel_handler" addrspace(4)** %3 to %"class.cl::sycl::kernel_handler" addrspace(4)* addrspace(4)*
%6 = addrspacecast i8 addrspace(4)** %4 to i8 addrspace(4)* addrspace(4)*
store %"class.cl::sycl::kernel_handler" addrspace(4)* %1, %"class.cl::sycl::kernel_handler" addrspace(4)* addrspace(4)* %5, align 8, !tbaa !4
%7 = load %"class.cl::sycl::kernel_handler" addrspace(4)*, %"class.cl::sycl::kernel_handler" addrspace(4)* addrspace(4)* %5, align 8
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
%8 = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %6, align 8, !tbaa !4
%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]])

%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]])

%13 = alloca %struct.coeff_str_t, align 8
%14 = addrspacecast %struct.coeff_str_t* %13 to %struct.coeff_str_t addrspace(4)*
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
ret void
}

; Function Attrs: convergent
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

; Function Attrs: convergent
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

attributes #0 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }

!opencl.spir.version = !{!1}
!spirv.Source = !{!2}
!llvm.ident = !{!3}

!1 = !{i32 1, i32 2}
!2 = !{i32 4, i32 100000}
!3 = !{!"clang version 14.0.0"}
!4 = !{!5, !5, i64 0}
!5 = !{!"any pointer", !6, i64 0}
!6 = !{!"omnipotent char", !7, i64 0}
!7 = !{!"Simple C++ TBAA"}
!8 = !{!9, !5, i64 0}
!9 = !{!"_ZTSN2cl4sycl14kernel_handlerE", !5, i64 0}

; CHECK-PROP: [SYCL/specialization constants]
; CHECK-PROP-NEXT: 9f47062a80eecfa7____ZL8coeff_id=2|gNAAAAAAAAAAAAAAAAAAAQAAAAQAAAAAEAAAAQAAAAgAAAAAIAAAAQAAAAwAAAAAMAAAAQAAAAABAAAAQAAAAQAAAAQBAAAAUAAAAQAAAAgBAAAAYAAAAQAAAAwBAAAAcAAAAQAAAAACAAAAgAAAAQAAAAA
; CHECK-PROP-NEXT: 405761736d5a1797____ZL9coeff_id2=2|gNAAAAAAAAQCAAAAAAAAAQAAAAgCAAAAEAAAAQAAAAwCAAAAIAAAAQAAAAADAAAAMAAAAQAAAAQDAAAAQAAAAQAAAAgDAAAAUAAAAQAAAAwDAAAAYAAAAQAAAAAEAAAAcAAAAQAAAAQEAAAAgAAAAQAAAAA
; CHECK-PROP-NEXT: 6da74a122db9f35d____ZL9coeff_id3=2|AGAAAAAAAAgEAAAAAAAAAQAAAAwEAAAAEAAAAQAAAAAFAAAAIAAAAQAAAAQFAAAAQAAAAgAAAAA

; CHECK-PROP: [SYCL/specialization constants default values]
; CHECK-PROP-NEXT: all=2|AMAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAg/AAAAA0MzMIQzMzoAZmZGDEAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAAA
93 changes: 52 additions & 41 deletions llvm/tools/sycl-post-link/SpecConstants.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -293,30 +293,40 @@ void collectCompositeElementsInfoRecursive(
void collectCompositeElementsDefaultValuesRecursive(
const Module &M, Constant *C, unsigned &Offset,
std::vector<char> &DefaultValues) {
Type *Ty = C->getType();
if (isa<ConstantAggregateZero>(C)) {
// This code is generic for zeroinitializer for both arrays and structs
size_t NumBytes = M.getDataLayout().getTypeStoreSize(C->getType());
std::fill_n(std::back_inserter(DefaultValues), NumBytes, 0);
Offset += NumBytes;
return;
}

if (auto *DataSeqC = dyn_cast<ConstantDataSequential>(C)) {
// This code is generic for both vectors and arrays of scalars
for (size_t I = 0; I < DataSeqC->getNumElements(); ++I) {
Constant *El = cast<Constant>(DataSeqC->getElementAsConstant(I));
collectCompositeElementsDefaultValuesRecursive(M, El, Offset,
DefaultValues);
}
} else if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
return;
}

if (auto *ArrayC = dyn_cast<ConstantArray>(C)) {
// This branch handles arrays of composite types (structs, arrays, etc.)
for (size_t I = 0; I < ArrTy->getNumElements(); ++I) {
Constant *El = cast<Constant>(C->getOperand(I));
collectCompositeElementsDefaultValuesRecursive(M, El, Offset,
DefaultValues);
assert(!C->isZeroValue() && "C must not be a zeroinitializer");
for (size_t I = 0; I < ArrayC->getType()->getNumElements(); ++I) {
collectCompositeElementsDefaultValuesRecursive(M, ArrayC->getOperand(I),
Offset, DefaultValues);
}
} else if (auto *StructTy = dyn_cast<StructType>(Ty)) {
return;
}

if (auto *StructC = dyn_cast<ConstantStruct>(C)) {
assert(!C->isZeroValue() && "C must not be a zeroinitializer");
auto *StructTy = StructC->getType();
const StructLayout *SL = M.getDataLayout().getStructLayout(StructTy);
const size_t BaseDefaultValueOffset = DefaultValues.size();
for (size_t I = 0, E = StructTy->getNumElements(); I < E; ++I) {
Constant *El = nullptr;
if (C->isZeroValue())
El = Constant::getNullValue(StructTy->getElementType(I));
else
El = cast<Constant>(C->getOperand(I));
// When handling elements of a structure, we do not use manually
// calculated offsets (which are sum of sizes of all previously
// encountered elements), but instead rely on data provided for us by
Expand All @@ -328,8 +338,8 @@ void collectCompositeElementsDefaultValuesRecursive(
while (LocalOffset != DefaultValues.size())
DefaultValues.push_back(0);

collectCompositeElementsDefaultValuesRecursive(M, El, LocalOffset,
DefaultValues);
collectCompositeElementsDefaultValuesRecursive(
M, StructC->getOperand(I), LocalOffset, DefaultValues);
}
const size_t SLSize = SL->getSizeInBytes();

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

if (auto IntConst = dyn_cast<ConstantInt>(C)) {
auto Val = IntConst->getValue().getZExtValue();
std::copy_n(reinterpret_cast<char *>(&Val), NumBytes,
// Assume that we encountered some scalar element
size_t NumBytes = M.getDataLayout().getTypeStoreSize(C->getType());
if (auto IntConst = dyn_cast<ConstantInt>(C)) {
auto Val = IntConst->getValue().getZExtValue();
std::copy_n(reinterpret_cast<char *>(&Val), NumBytes,
std::back_inserter(DefaultValues));
} else if (auto FPConst = dyn_cast<ConstantFP>(C)) {
auto Val = FPConst->getValue();

if (NumBytes == 2) {
auto IVal = Val.bitcastToAPInt();
assert(IVal.getBitWidth() == 16);
auto Storage = static_cast<uint16_t>(IVal.getZExtValue());
std::copy_n(reinterpret_cast<char *>(&Storage), NumBytes,
std::back_inserter(DefaultValues));
} else if (NumBytes == 4) {
float v = Val.convertToFloat();
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
std::back_inserter(DefaultValues));
} else if (NumBytes == 8) {
double v = Val.convertToDouble();
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
std::back_inserter(DefaultValues));
} else if (auto FPConst = dyn_cast<ConstantFP>(C)) {
auto Val = FPConst->getValue();

if (NumBytes == 2) {
auto IVal = Val.bitcastToAPInt();
assert(IVal.getBitWidth() == 16);
auto Storage = static_cast<uint16_t>(IVal.getZExtValue());
std::copy_n(reinterpret_cast<char *>(&Storage), NumBytes,
std::back_inserter(DefaultValues));
} else if (NumBytes == 4) {
float v = Val.convertToFloat();
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
std::back_inserter(DefaultValues));
} else if (NumBytes == 8) {
double v = Val.convertToDouble();
std::copy_n(reinterpret_cast<char *>(&v), NumBytes,
std::back_inserter(DefaultValues));
} else {
llvm_unreachable("Unexpected constant floating point type");
}
} else {
llvm_unreachable("Unexpected constant scalar type");
llvm_unreachable("Unexpected constant floating point type");
}
Offset += NumBytes;
} else {
llvm_unreachable("Unexpected constant scalar type");
}
Offset += NumBytes;
}

MDNode *generateSpecConstantMetadata(const Module &M, StringRef SymbolicID,
Expand Down