Skip to content

[SYCL][ESIMD] Fix the crash in sycl-post-link while processing global spirv functions. #7590

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 5 commits into from
Dec 5, 2022
Merged

[SYCL][ESIMD] Fix the crash in sycl-post-link while processing global spirv functions. #7590

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
62b8ac6
Fix a crash of sycl-post-link
fineg74 Nov 29, 2022
b06b5b4
Fix the test
fineg74 Nov 30, 2022
da12a1f
Address PR comments
fineg74 Nov 30, 2022
37c1616
Address PR comments
fineg74 Dec 1, 2022
fdb6ac5
Address PR comments
fineg74 Dec 2, 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
25 changes: 15 additions & 10 deletions llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1142,9 +1142,11 @@ static uint64_t getIndexFromExtract(ExtractElementInst *EEI) {
/// right before the given extract element instruction \p EEI using the result
/// of vector load. The parameter \p IsVectorCall tells what version of GenX
/// intrinsic (scalar or vector) to use to lower the load from SPIRV global.
static Instruction *generateGenXCall(ExtractElementInst *EEI,
StringRef IntrinName, bool IsVectorCall) {
uint64_t IndexValue = getIndexFromExtract(EEI);
static Instruction *generateGenXCall(Instruction *EEI, StringRef IntrinName,
bool IsVectorCall) {
uint64_t IndexValue = isa<ExtractElementInst>(EEI)
? getIndexFromExtract(cast<ExtractElementInst>(EEI))
: 0;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think using '0' for the getelementptr case is wrong. Instead, index from the GEP instruction should be extracted (marked below)

%0 = load i64, i64 addrspace(1)* getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0), align 32

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is actually for trunc instruction. In the pathological case there is no getElementptr i.e. for

%0 = load i64, i64 addrspace(1)* getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0), align 32
  %conv = trunc i64 %0 to i32

case and that is why index doesn't matter much here.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't understand why index does not matter.

std::string Suffix =
IsVectorCall
? ".v3i32"
Expand Down Expand Up @@ -1244,12 +1246,15 @@ translateSpirvGlobalUses(LoadInst *LI, StringRef SpirvGlobalName,
}

// Only loads from _vector_ SPIRV globals reach here now. Their users are
// expected to be ExtractElementInst only, and they are replaced in this loop.
// When loads from _scalar_ SPIRV globals are handled here as well, the users
// will not be replaced by new instructions, but the GenX call replacing the
// original load 'LI' should be inserted before each user.
// expected to be ExtractElementInst or TruncInst only, and they are replaced
// in this loop. When loads from _scalar_ SPIRV globals are handled here as
// well, the users will not be replaced by new instructions, but the GenX call
// replacing the original load 'LI' should be inserted before each user.
for (User *LU : LI->users()) {
ExtractElementInst *EEI = cast<ExtractElementInst>(LU);
assert(
(isa<ExtractElementInst>(LU) || isa<TruncInst>(LU)) &&
"SPIRV global users should be either ExtractElementInst or TruncInst");
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry for being late for the review, I don't see such restrictions in SPIR-V spec, could you please help me finding it?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These restrictions are to protect the code to make sure it processes only the code it is expected to process and not to enforce SPIR-V spec.

Instruction *EEI = cast<Instruction>(LU);
NewInst = nullptr;

if (SpirvGlobalName == "WorkgroupSize") {
Expand Down Expand Up @@ -1786,10 +1791,10 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
if (LI) {
Value *LoadPtrOp = LI->getPointerOperand();
Value *SpirvGlobal = nullptr;
// Look through casts to find SPIRV builtin globals
// Look through constant expressions to find SPIRV builtin globals
// It may come with or without cast.
auto *CE = dyn_cast<ConstantExpr>(LoadPtrOp);
if (CE) {
assert(CE->isCast() && "ConstExpr should be a cast");
SpirvGlobal = CE->getOperand(0);
} else {
SpirvGlobal = LoadPtrOp;
Expand Down
153 changes: 153 additions & 0 deletions sycl/test/esimd/sycl-post-link-test.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,153 @@
; RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %s -o %t.table
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

please move the test to llvm/test/tools/sycl-post-link

; RUN: FileCheck %s -input-file=%t_esimd_0.ll
; This test checks that IR code below can be successfully processed by
; sycl-post-link. In this IR no extractelement instruction and no casting are used

; ModuleID = 'sycl-post-link-test.cpp'
source_filename = "sycl-post-link-test.cpp"
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.sycl::_V1::ext::intel::esimd::simd" = type { %"class.sycl::_V1::ext::intel::esimd::detail::simd_obj_impl" }
%"class.sycl::_V1::ext::intel::esimd::detail::simd_obj_impl" = type { <8 x float> }
%"class.sycl::_V1::ext::intel::esimd::simd.0" = type { %"class.sycl::_V1::ext::intel::esimd::detail::simd_obj_impl.1" }
%"class.sycl::_V1::ext::intel::esimd::detail::simd_obj_impl.1" = type { <8 x i32> }

@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32

; Function Attrs: convergent norecurse
define dso_local spir_func void @vmult2(%"class.sycl::_V1::ext::intel::esimd::simd"* noundef %a) local_unnamed_addr #0 !srcloc !47 !sycl_explicit_simd !45 !intel_reqd_sub_group_size !48 !sycl_fixed_targets !45 {
entry:
%Res.i.i.i = alloca %"class.sycl::_V1::ext::intel::esimd::simd", align 32
%Res.i.i = alloca %"class.sycl::_V1::ext::intel::esimd::simd", align 32
%ref.tmp.i = alloca %"class.sycl::_V1::ext::intel::esimd::simd.0", align 32
%0 = load i64, i64 addrspace(1)* getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0), align 32
%conv = trunc i64 %0 to i32
%1 = bitcast %"class.sycl::_V1::ext::intel::esimd::simd.0"* %ref.tmp.i to i8*
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

  1. Please simplify the test. Only instructions participating in pattern matching should be left.
  2. Please inline //CHECK: directives right after source IR instruction they are supposed to verify
    transformation of

call void @llvm.lifetime.start.p0i8(i64 32, i8* nonnull %1) #3
%splat.splatinsert.i.i.i = insertelement <8 x i32> poison, i32 %conv, i64 0
%splat.splat.i.i.i = shufflevector <8 x i32> %splat.splatinsert.i.i.i, <8 x i32> poison, <8 x i32> zeroinitializer
%M_data.i.i.i = getelementptr inbounds %"class.sycl::_V1::ext::intel::esimd::simd.0", %"class.sycl::_V1::ext::intel::esimd::simd.0"* %ref.tmp.i, i64 0, i32 0, i32 0
%2 = addrspacecast <8 x i32>* %M_data.i.i.i to <8 x i32> addrspace(4)*
store <8 x i32> %splat.splat.i.i.i, <8 x i32>* %M_data.i.i.i, align 32, !tbaa !49
%3 = bitcast %"class.sycl::_V1::ext::intel::esimd::simd"* %Res.i.i to i8*
call void @llvm.lifetime.start.p0i8(i64 32, i8* nonnull %3) #3
%4 = bitcast %"class.sycl::_V1::ext::intel::esimd::simd"* %Res.i.i.i to i8*
call void @llvm.lifetime.start.p0i8(i64 32, i8* nonnull %4) #3, !noalias !52
%M_data.i.i.i.i = getelementptr inbounds %"class.sycl::_V1::ext::intel::esimd::simd", %"class.sycl::_V1::ext::intel::esimd::simd"* %a, i64 0, i32 0, i32 0
%5 = addrspacecast <8 x float>* %M_data.i.i.i.i to <8 x float> addrspace(4)*
%6 = addrspacecast <8 x float>* %M_data.i.i.i.i to <8 x float> addrspace(4)*
%call.i.i.i.i = tail call spir_func noundef <8 x float> @_Z13__esimd_vloadIfLi8EEN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeEPKS9_(<8 x float> addrspace(4)* noundef %6) #4, !noalias !52
%call.i6.i.i.i = call spir_func noundef <8 x i32> @_Z13__esimd_vloadIiLi8EEN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeEPKS9_(<8 x i32> addrspace(4)* noundef %2) #4, !noalias !52
%conv.i.i.i.i = sitofp <8 x i32> %call.i6.i.i.i to <8 x float>
%mul.i.i.i.i.i.i = fmul <8 x float> %call.i.i.i.i, %conv.i.i.i.i
%M_data.i.i.i.i.i.i = getelementptr inbounds %"class.sycl::_V1::ext::intel::esimd::simd", %"class.sycl::_V1::ext::intel::esimd::simd"* %Res.i.i.i, i64 0, i32 0, i32 0
%7 = addrspacecast <8 x float>* %M_data.i.i.i.i.i.i to <8 x float> addrspace(4)*
%8 = addrspacecast <8 x float>* %M_data.i.i.i.i.i.i to <8 x float> addrspace(4)*
call spir_func void @_Z14__esimd_vstoreIfLi8EEvPN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeES9_(<8 x float> addrspace(4)* noundef %7, <8 x float> noundef %mul.i.i.i.i.i.i) #4, !noalias !52
%call.i.i.i.i.i.i = call spir_func noundef <8 x float> @_Z13__esimd_vloadIfLi8EEN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeEPKS9_(<8 x float> addrspace(4)* noundef %8) #4, !noalias !52
%M_data.i2.i.i.i.i.i = getelementptr inbounds %"class.sycl::_V1::ext::intel::esimd::simd", %"class.sycl::_V1::ext::intel::esimd::simd"* %Res.i.i, i64 0, i32 0, i32 0
%9 = addrspacecast <8 x float>* %M_data.i2.i.i.i.i.i to <8 x float> addrspace(4)*
%10 = addrspacecast <8 x float>* %M_data.i2.i.i.i.i.i to <8 x float> addrspace(4)*
call spir_func void @_Z14__esimd_vstoreIfLi8EEvPN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeES9_(<8 x float> addrspace(4)* noundef %9, <8 x float> noundef %call.i.i.i.i.i.i) #4
call void @llvm.lifetime.end.p0i8(i64 32, i8* nonnull %4) #3, !noalias !52
%call.i.i.i = call spir_func noundef <8 x float> @_Z13__esimd_vloadIfLi8EEN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeEPKS9_(<8 x float> addrspace(4)* noundef %10) #4
call spir_func void @_Z14__esimd_vstoreIfLi8EEvPN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeES9_(<8 x float> addrspace(4)* noundef %5, <8 x float> noundef %call.i.i.i) #4
call void @llvm.lifetime.end.p0i8(i64 32, i8* nonnull %3) #3
call void @llvm.lifetime.end.p0i8(i64 32, i8* nonnull %1) #3
ret void
}

; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.start.p0i8(i64 immarg, i8* nocapture) #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.end.p0i8(i64 immarg, i8* nocapture) #1

; Function Attrs: convergent
declare dso_local spir_func noundef <8 x i32> @_Z13__esimd_vloadIiLi8EEN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeEPKS9_(<8 x i32> addrspace(4)* noundef) local_unnamed_addr #2

; Function Attrs: convergent
declare dso_local spir_func void @_Z14__esimd_vstoreIfLi8EEvPN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeES9_(<8 x float> addrspace(4)* noundef, <8 x float> noundef) local_unnamed_addr #2

; Function Attrs: convergent
declare dso_local spir_func noundef <8 x float> @_Z13__esimd_vloadIfLi8EEN4sycl3_V13ext5intel5esimd6detail15raw_vector_typeIT_XT0_EE4typeEPKS9_(<8 x float> addrspace(4)* noundef) local_unnamed_addr #2

declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...)

attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="256" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="sycl-post-link-test.cpp" }
attributes #1 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #3 = { nounwind }
attributes #4 = { convergent }

!llvm.module.flags = !{!0, !1}
!opencl.spir.version = !{!2}
!spirv.Source = !{!3}
!sycl_aspects = !{!4, !5, !6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39, !40, !41, !42, !43, !44}
!opencl.compiler.options = !{!45}
!llvm.ident = !{!46}

!0 = !{i32 1, !"wchar_size", i32 4}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

likewise, all unnecessary metadata should be removed.

!1 = !{i32 7, !"frame-pointer", i32 2}
!2 = !{i32 1, i32 2}
!3 = !{i32 0, i32 100000}
!4 = !{!"host", i32 0}
!5 = !{!"cpu", i32 1}
!6 = !{!"gpu", i32 2}
!7 = !{!"accelerator", i32 3}
!8 = !{!"custom", i32 4}
!9 = !{!"fp16", i32 5}
!10 = !{!"fp64", i32 6}
!11 = !{!"image", i32 9}
!12 = !{!"online_compiler", i32 10}
!13 = !{!"online_linker", i32 11}
!14 = !{!"queue_profiling", i32 12}
!15 = !{!"usm_device_allocations", i32 13}
!16 = !{!"usm_host_allocations", i32 14}
!17 = !{!"usm_shared_allocations", i32 15}
!18 = !{!"usm_restricted_shared_allocations", i32 16}
!19 = !{!"usm_system_allocations", i32 17}
!20 = !{!"ext_intel_pci_address", i32 18}
!21 = !{!"ext_intel_gpu_eu_count", i32 19}
!22 = !{!"ext_intel_gpu_eu_simd_width", i32 20}
!23 = !{!"ext_intel_gpu_slices", i32 21}
!24 = !{!"ext_intel_gpu_subslices_per_slice", i32 22}
!25 = !{!"ext_intel_gpu_eu_count_per_subslice", i32 23}
!26 = !{!"ext_intel_max_mem_bandwidth", i32 24}
!27 = !{!"ext_intel_mem_channel", i32 25}
!28 = !{!"usm_atomic_host_allocations", i32 26}
!29 = !{!"usm_atomic_shared_allocations", i32 27}
!30 = !{!"atomic64", i32 28}
!31 = !{!"ext_intel_device_info_uuid", i32 29}
!32 = !{!"ext_oneapi_srgb", i32 30}
!33 = !{!"ext_oneapi_native_assert", i32 31}
!34 = !{!"host_debuggable", i32 32}
!35 = !{!"ext_intel_gpu_hw_threads_per_eu", i32 33}
!36 = !{!"ext_oneapi_cuda_async_barrier", i32 34}
!37 = !{!"ext_oneapi_bfloat16", i32 35}
!38 = !{!"ext_intel_free_memory", i32 36}
!39 = !{!"ext_intel_device_id", i32 37}
!40 = !{!"ext_intel_memory_clock_rate", i32 38}
!41 = !{!"ext_intel_memory_bus_width", i32 39}
!42 = !{!"int64_base_atomics", i32 7}
!43 = !{!"int64_extended_atomics", i32 8}
!44 = !{!"usm_system_allocator", i32 17}
!45 = !{}
!46 = !{!"Intel(R) oneAPI DPC++/C++ Compiler 2023.1.0 (2023.x.0.YYYYMMDD)"}
!47 = !{i32 606}
!48 = !{i32 1}
!49 = !{!50, !50, i64 0}
!50 = !{!"omnipotent char", !51, i64 0}
!51 = !{!"Simple C++ TBAA"}
!52 = !{!53}
!53 = distinct !{!53, !54, !"_ZN4sycl3_V13ext5intel5esimd6detailmlIfiLi8ENS3_4simdENS6_IfLi8EEEvEEDaRKNS4_13simd_obj_implINS4_19element_type_traitsIT_vE4RawTEXT1_ET2_ISA_XT1_EEvEERKNS8_INS9_IT0_vE4RawTEXT1_ESD_ISI_XT1_EEvEE: %agg.result"}
!54 = distinct !{!54, !"_ZN4sycl3_V13ext5intel5esimd6detailmlIfiLi8ENS3_4simdENS6_IfLi8EEEvEEDaRKNS4_13simd_obj_implINS4_19element_type_traitsIT_vE4RawTEXT1_ET2_ISA_XT1_EEvEERKNS8_INS9_IT0_vE4RawTEXT1_ESD_ISI_XT1_EEvEE"}


; CHECK: define dso_local spir_func void @vmult2
; CHECK: call <3 x i32> @llvm.genx.local.id.v3i32()
; CHECK: call <3 x i32> @llvm.genx.local.size.v3i32()
; CHECK: call i32 @llvm.genx.group.id.x()
; CHECK: ret void
; CHECK: }
2 changes: 1 addition & 1 deletion sycl/test/lit.cfg.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@
config.test_format = lit.formats.ShTest()

# suffixes: A list of file extensions to treat as test files.
config.suffixes = ['.c', '.cpp', '.dump', '.test'] #add .spv. Currently not clear what to do with those
config.suffixes = ['.c', '.cpp', '.dump', '.test', '.ll'] #add .spv. Currently not clear what to do with those

# feature tests are considered not so lightweight, so, they are excluded by default
config.excludes = ['Inputs', 'feature-tests']
Expand Down