Skip to content

[sycl-post-link] Fix a crash during spec-constant properties generation #5538

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
bader merged 7 commits into from Feb 18, 2022
Merged

[sycl-post-link] Fix a crash during spec-constant properties generation #5538

Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
f0e14b5
[sycl-post-link] Fix a crash during spec-constant properties generation
Feb 10, 2022
5b6fd9a
[sycl-post-link] Optimize __spirv_SpecConstantComposite generation
Feb 10, 2022
564ffd4
[sycl-post-link] Add a lit test
Feb 10, 2022
c779c3e
[sycl-post-link][NFC] Fix formatting issues
Feb 10, 2022
b1d7f1d
[sycl-post-link] Add a lit test to check a primitive padding
Feb 11, 2022
7ac0966
[sycl-post-link] Add undefined value handling
Feb 15, 2022
b7ffefe
[sycl-post-link] Check spec. constant metadata generation in the lit
Feb 15, 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
82 changes: 82 additions & 0 deletions llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,82 @@
; 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 structs
; have an 'undef' value for padding 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 { %struct.coeff_str_aligned_t }
%"class.cl::sycl::specialization_id.1" = type { %struct.coeff2_str_aligned_t }
%struct.coeff_str_aligned_t = type { %"class.std::array", i64, [8 x i8] }
%struct.coeff2_str_aligned_t = type { %"class.std::array", i64, [7 x i8], i8 }
%"class.std::array" = type { [3 x float] }

$_ZTSZ4mainEUlN2cl4sycl14kernel_handlerEE_ = comdat any

@__usid_str = private unnamed_addr constant [32 x i8] c"ef880fa09cf7a9d7____ZL8coeff_id\00", align 1
@_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
@_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

; Function Attrs: convergent norecurse
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 {
%1 = alloca %struct.coeff_str_aligned_t, align 32
%2 = addrspacecast %struct.coeff_str_aligned_t* %1 to %struct.coeff_str_aligned_t addrspace(4)*
%3 = bitcast %struct.coeff_str_aligned_t* %1 to i8*
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
; 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" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS3]])
; CHECK-IR: %[[#NS5:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID3:]], i64 0)
; 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)

%4 = alloca %struct.coeff2_str_aligned_t, align 32
%5 = addrspacecast %struct.coeff2_str_aligned_t* %4 to %struct.coeff2_str_aligned_t addrspace(4)*
%6 = bitcast %struct.coeff2_str_aligned_t* %4 to i8*
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 ([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.1" addrspace(1)* @_ZL8coeff_id2 to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4
; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
; CHECK-IR: %[[#NS8:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
; CHECK-IR: %[[#NS9:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
; CHECK-IR: %[[#NS10:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS7]], float %[[#NS8]], float %[[#NS9]])
; CHECK-IR: %[[#NS11:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS10]])
; CHECK-IR: %[[#NS12:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID7:]], i64 0)
; 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)

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

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

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" }
attributes #1 = { argmemonly mustprogress nofree nosync nounwind willreturn }
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #3 = { nounwind }
attributes #4 = { convergent }

!llvm.dependent-libraries = !{!0}
!llvm.module.flags = !{!1, !2}
!opencl.spir.version = !{!3}
!spirv.Source = !{!4}
!llvm.ident = !{!5}

!0 = !{!"libcpmt"}
!1 = !{i32 1, !"wchar_size", i32 2}
!2 = !{i32 7, !"frame-pointer", i32 2}
!3 = !{i32 1, i32 2}
!4 = !{i32 4, i32 100000}
!5 = !{!"clang version 14.0.0"}
!6 = !{i32 -1}
!7 = !{i1 true}

; CHECK-PROP: [SYCL/specialization constants]
; CHECK-PROP-NEXT: ef880fa09cf7a9d7____ZL8coeff_id=2|

; CHECK-PROP: [SYCL/specialization constants default values]
; CHECK-PROP-NEXT: all=2|
29 changes: 16 additions & 13 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 @@ -306,8 +306,8 @@ void collectCompositeElementsInfoRecursive(
void collectCompositeElementsDefaultValuesRecursive(
const Module &M, Constant *C, unsigned &Offset,
std::vector<char> &DefaultValues) {
if (isa<ConstantAggregateZero>(C)) {
// This code is generic for zeroinitializer for both arrays and structs
if (isa<ConstantAggregateZero>(C) || isa<UndefValue>(C)) {
// This code is generic for both arrays and structs
size_t NumBytes = M.getDataLayout().getTypeStoreSize(C->getType());
std::fill_n(std::back_inserter(DefaultValues), NumBytes, 0);
Offset += NumBytes;
Expand Down Expand Up @@ -519,14 +519,9 @@ Instruction *emitSpecConstant(unsigned NumericID, Type *Ty,
return emitCall(Ty, SPIRV_GET_SPEC_CONST_VAL, Args, InsertBefore);
}

Instruction *emitSpecConstantComposite(Type *Ty,
ArrayRef<Instruction *> Elements,
Instruction *emitSpecConstantComposite(Type *Ty, ArrayRef<Value *> Elements,
Instruction *InsertBefore) {
SmallVector<Value *, 8> Args(Elements.size());
for (unsigned I = 0; I < Elements.size(); ++I) {
Args[I] = cast<Value>(Elements[I]);
}
return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Args, InsertBefore);
return emitCall(Ty, SPIRV_GET_SPEC_CONST_COMPOSITE, Elements, InsertBefore);
}

/// For specified specialization constant type emits LLVM IR which is required
Expand Down Expand Up @@ -565,16 +560,24 @@ Instruction *emitSpecConstantRecursiveImpl(Type *Ty, Instruction *InsertBefore,
return emitSpecConstant(IDs[Index++], Ty, InsertBefore, DefaultValue);
}

SmallVector<Instruction *, 8> Elements;
SmallVector<Value *, 8> Elements;
auto LoopIteration = [&](Type *Ty, unsigned LocalIndex) {
// Select corresponding element of the default value if it was provided
Constant *Def =
DefaultValue ? DefaultValue->getAggregateElement(LocalIndex) : nullptr;
Elements.push_back(
emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index, Def));
if (isa_and_nonnull<UndefValue>(Def))
Elements.push_back(Def);
else
Elements.push_back(
emitSpecConstantRecursiveImpl(Ty, InsertBefore, IDs, Index, Def));
};

if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
if (isa_and_nonnull<UndefValue>(DefaultValue)) {
// If the default value is a composite and has the value 'undef', we should
// not generate a bunch of __spirv_SpecConstant for its elements but
// pass it into __spirv_SpecConstantComposite as is.
Elements.push_back(DefaultValue);
} else if (auto *ArrTy = dyn_cast<ArrayType>(Ty)) {
for (size_t I = 0; I < ArrTy->getNumElements(); ++I) {
LoopIteration(ArrTy->getElementType(), I);
}
Expand Down