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

fineg74 · 2022-11-30T03:51:03Z

Currently for int i = __spirv_GlobalInvocationId_x(); c++ code followinf IR code is generated:

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

This IR is what sycl-post-link is expecting and is able to process successfully.
However, following IR code was generated during the testing:

%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

and it caused sycl-post-link to crash.
The fix resolves the issue

kbobrovs · 2022-11-30T05:42:43Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

+                                     bool IsVectorCall) {
+  uint64_t IndexValue = isa<ExtractElementInst>(EEI)
+                            ? getIndexFromExtract(cast<ExtractElementInst>(EEI))
+                            : 0;


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

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.

I don't understand why index does not matter.

kbobrovs · 2022-11-30T05:46:22Z

sycl/test/esimd/sycl-post-link-test.ll

+  %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*


Please simplify the test. Only instructions participating in pattern matching should be left.

Please inline //CHECK: directives right after source IR instruction they are supposed to verify
transformation of

kbobrovs · 2022-11-30T05:46:52Z

sycl/test/esimd/sycl-post-link-test.ll

+!opencl.compiler.options = !{!45}
+!llvm.ident = !{!46}
+
+!0 = !{i32 1, !"wchar_size", i32 4}


likewise, all unnecessary metadata should be removed.

kbobrovs · 2022-11-30T05:49:12Z

sycl/test/esimd/sycl-post-link-test.ll

@@ -0,0 +1,153 @@
+; RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %s -o %t.table


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

MrSidims · 2022-12-01T14:25:09Z

BTW can we instead of handling this GEP(0,0) in sycl-post-link - remove this GEP at all here? It seems to be incorrect to access a vector via GEP and GEP(0,0) is quite useless anyway. Though, unfortunately, we can't just call replaceAllUsesWith...

kbobrovs · 2022-12-01T18:35:42Z

It seems to be incorrect to access a vector via GEP and GEP(0,0) is quite useless anyway.

Good catch. extractelement should be used for accessing vector elements normally, even though for Intel GPU target the GEP will generate correct address. I wonder who generates the questionable IR now:
%0 = load i64, i64 addrspace(1)* getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0), align 32

kbobrovs · 2022-12-01T18:36:15Z

kbobrovs closed this 17 minutes ago

sorry, hit wrong button

Fznamznon · 2022-12-01T18:38:58Z

It seems to be incorrect to access a vector via GEP and GEP(0,0) is quite useless anyway.

Good catch. extractelement should be used for accessing vector elements normally, even though for Intel GPU target the GEP will generate correct address. I wonder who generates the questionable IR now: %0 = load i64, i64 addrspace(1)* getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0), align 32

This IR appeared during the pulldown after llvm/llvm-project@163bb6d.

kbobrovs · 2022-12-01T18:45:28Z

It seems to be incorrect to access a vector via GEP and GEP(0,0) is quite useless anyway.

Good catch. extractelement should be used for accessing vector elements normally, even though for Intel GPU target the GEP will generate correct address. I wonder who generates the questionable IR now: %0 = load i64, i64 addrspace(1)* getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0), align 32

I take this back.
Actually, I believe getelementptr (<3 x i64>, <3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId, i64 0, i64 0) is OK.
extractelement, of course, is required to be used to access elements, but we are calculating vector element's offset in memory. Vector's layout in memory is defined by the target, so for SPIRV target this seems perfectly OK code.

Fznamznon · 2022-12-01T20:22:59Z

Can we proceed with this patch? The pulldown is blocked by the crash being fixed here.

kbobrovs · 2022-12-02T04:42:38Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

+      auto *GEPCE =
+          dyn_cast<GetElementPtrConstantExpr>(LI->getPointerOperand());
+      if (GEPCE) {
+        IndexValue = cast<Constant>(GEPCE->getOperand(2))
+                         ->getUniqueInteger()
+                         .getZExtValue();


The above code silently produces invalid IR if it gets unexpected IR. Please fix.

Suggested change

auto *GEPCE =

dyn_cast<GetElementPtrConstantExpr>(LI->getPointerOperand());

if (GEPCE) {

IndexValue = cast<Constant>(GEPCE->getOperand(2))

->getUniqueInteger()

.getZExtValue();

auto *GEPCE =

cast<GetElementPtrConstantExpr>(LI->getPointerOperand());

IndexValue = cast<Constant>(GEPCE->getOperand(2))

->getUniqueInteger()

.getZExtValue();

kbobrovs · 2022-12-02T04:45:04Z

Can we proceed with this patch? The pulldown is blocked by the crash being fixed here.

not yet, there is still a bug to be fixed. (see above comment). I believe the patch can be merged once it is addressed.

kbobrovs

This comment hasn't been addressed:
Please inline //CHECK: directives right after source IR instruction they are supposed to verify transformation of

Approving for the sake of urgency. Please address with follow-up patch.

fineg74 · 2022-12-05T19:25:33Z

@intel/dpcpp-tools-reviewers Could you please take a look ?

sarnex

Approving for urgency, leaving correctness review to Konst and other reviewers

kbobrovs · 2022-12-05T21:03:49Z

HIP Failures are unrelated:

Unresolved Tests (5):
SYCL :: Assert/assert_in_kernels.cpp
SYCL :: Assert/assert_in_multiple_tus.cpp
SYCL :: Assert/assert_in_multiple_tus_one_ndebug.cpp
SYCL :: Assert/assert_in_one_kernel.cpp
SYCL :: Assert/assert_in_simultaneously_multiple_tus_one_ndebug.cpp

Timed Out Tests (1):
SYCL :: AtomicRef/max_local.cpp

Fznamznon · 2022-12-05T21:07:25Z

The guilty llvm change that caused crash just enables an additional optimization. To unblock pulldown the change that disables the optimization for sycl device code was merged. But this is not a proper solution. I've been testing workaround in SPIR-V translator together with reverting change that disables guilty optimization and also cherry-picked changes from this PR for testing. I've started testing here https://github.com/intel/llvm/actions/runs/3623046607/jobs/6108488479 . The thing is, with changes from this PR, I still see one of the tests failing, but with another error:

sycl-post-link: /__w/llvm/llvm/src/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp:1252: void {anonymous}::translateSpirvGlobalUses(llvm::LoadInst*, llvm::StringRef, llvm::SmallVectorImpl<llvm::Instruction*>&): Assertion `(isa<ExtractElementInst>(LU) || isa<TruncInst>(LU)) && "SPIRV global users should be either ExtractElementInst or TruncInst"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: sycl-post-link -split-esimd -lower-esimd -O0 -S /__w/llvm/llvm/build/tools/sycl/test/esimd/Output/spirv_intrins_trans.cpp.tmp -o /__w/llvm/llvm/build/tools/sycl/test/esimd/Output/spirv_intrins_trans.cpp.tmp.table
 #0 0x000055774af13ff4 PrintStackTraceSignalHandler(void*) Signals.cpp:0:0
 #1 0x000055774af11fe4 SignalHandler(int) Signals.cpp:0:0
 #2 0x00007fc293202420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)
 #3 0x00007fc292ca500b raise (/lib/x86_64-linux-gnu/libc.so.6+0x4300b)
 #4 0x00007fc292c84859 abort (/lib/x86_64-linux-gnu/libc.so.6+0x22859)
 #5 0x00007fc292c84729 (/lib/x86_64-linux-gnu/libc.so.6+0x22729)
 #6 0x00007fc292c95fd6 (/lib/x86_64-linux-gnu/libc.so.6+0x33fd6)
 #7 0x000055774af8fdd1 llvm::SYCLLowerESIMDPass::runOnFunction(llvm::Function&, llvm::SmallPtrSet<llvm::Type*, 4u>&) (/__w/llvm/llvm/build/bin/sycl-post-link+0x3a5dd1)
 #8 0x000055774af92180 llvm::SYCLLowerESIMDPass::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/__w/llvm/llvm/build/bin/sycl-post-link+0x3a8180)
 #9 0x000055774acb05b6 llvm::detail::PassModel<llvm::Module, llvm::SYCLLowerESIMDPass, llvm::PreservedAnalyses, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/__w/llvm/llvm/build/bin/sycl-post-link+0xc65b6)
#10 0x000055774ae30c69 llvm::PassManager<llvm::Module, llvm::AnalysisManager<llvm::Module>>::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (/__w/llvm/llvm/build/bin/sycl-post-link+0x246c69)
#11 0x000055774acbc574 (anonymous namespace)::lowerEsimdConstructs(llvm::module_split::ModuleDesc&) sycl-post-link.cpp:0:0
#12 0x000055774acc2b7c (anonymous namespace)::processInputModule(std::unique_ptr<llvm::Module, std::default_delete<llvm::Module>>) sycl-post-link.cpp:0:0
#13 0x000055774ac5e0b6 main (/__w/llvm/llvm/build/bin/sycl-post-link+0x740b6)
#14 0x00007fc292c86083 __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x24083)
#15 0x000055774acaf85e _start (/__w/llvm/llvm/build/bin/sycl-post-link+0xc585e)

error: command failed with exit status: -6

--

********************
Testing:  0.. 10.. 20.. 30.. 40.. 50.. 60.. 70.. 80.. 90.. 
********************
Failed Tests (1):
  SYCL :: esimd/spirv_intrins_trans.cpp

Could you please double check that everything is correct?

kbobrovs · 2022-12-05T21:30:26Z

@fineg74, please take a look at the issue reported above.

MrSidims · 2022-12-06T15:14:03Z

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

-    ExtractElementInst *EEI = cast<ExtractElementInst>(LU);
+    assert(
+        (isa<ExtractElementInst>(LU) || isa<TruncInst>(LU)) &&
+        "SPIRV global users should be either ExtractElementInst or TruncInst");


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

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.

fineg74 · 2022-12-07T01:57:42Z

@fineg74, please take a look at the issue reported above.

Folow up PR: #7673

fineg74 and others added 2 commits November 29, 2022 14:56

Fix a crash of sycl-post-link

62b8ac6

Fix the test

b06b5b4

fineg74 requested review from a team as code owners November 30, 2022 03:51

fineg74 requested a review from againull November 30, 2022 03:51

kbobrovs reviewed Nov 30, 2022

View reviewed changes

Address PR comments

da12a1f

fineg74 requested a review from a team as a code owner November 30, 2022 07:51

Address PR comments

37c1616

kbobrovs closed this Dec 1, 2022

kbobrovs reopened this Dec 1, 2022

fineg74 mentioned this pull request Dec 2, 2022

[SYCL][ESIMD] Introduce a test to validate correct processing of spirv global builtin calls intel/llvm-test-suite#1431

Closed

kbobrovs reviewed Dec 2, 2022

View reviewed changes

Address PR comments

fdb6ac5

againull removed request for a team and againull December 2, 2022 19:55

kbobrovs approved these changes Dec 5, 2022

View reviewed changes

sarnex approved these changes Dec 5, 2022

View reviewed changes

kbobrovs merged commit 63c749c into intel:sycl Dec 5, 2022

MrSidims reviewed Dec 6, 2022

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

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

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

fineg74 commented Nov 30, 2022

kbobrovs Nov 30, 2022

fineg74 Nov 30, 2022

kbobrovs Nov 30, 2022

kbobrovs Nov 30, 2022

kbobrovs Nov 30, 2022

kbobrovs Nov 30, 2022

MrSidims commented Dec 1, 2022 •

edited

Loading

kbobrovs commented Dec 1, 2022

kbobrovs commented Dec 1, 2022

Fznamznon commented Dec 1, 2022

kbobrovs commented Dec 1, 2022

Fznamznon commented Dec 1, 2022

kbobrovs Dec 2, 2022

kbobrovs commented Dec 2, 2022 •

edited

Loading

kbobrovs left a comment

fineg74 commented Dec 5, 2022

sarnex left a comment

kbobrovs commented Dec 5, 2022

Fznamznon commented Dec 5, 2022

kbobrovs commented Dec 5, 2022

MrSidims Dec 6, 2022

fineg74 Dec 7, 2022

fineg74 commented Dec 7, 2022

		@@ -0,0 +1,153 @@
		; RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %s -o %t.table

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

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

Conversation

fineg74 commented Nov 30, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MrSidims commented Dec 1, 2022 • edited Loading

kbobrovs commented Dec 1, 2022

kbobrovs commented Dec 1, 2022

Fznamznon commented Dec 1, 2022

kbobrovs commented Dec 1, 2022

Fznamznon commented Dec 1, 2022

Choose a reason for hiding this comment

kbobrovs commented Dec 2, 2022 • edited Loading

kbobrovs left a comment

Choose a reason for hiding this comment

fineg74 commented Dec 5, 2022

sarnex left a comment

Choose a reason for hiding this comment

kbobrovs commented Dec 5, 2022

Fznamznon commented Dec 5, 2022

kbobrovs commented Dec 5, 2022

Choose a reason for hiding this comment

Choose a reason for hiding this comment

fineg74 commented Dec 7, 2022

MrSidims commented Dec 1, 2022 •

edited

Loading

kbobrovs commented Dec 2, 2022 •

edited

Loading