Skip to content

[SYCL] Ignore unknown users of virtual functions instead of crashing #17397

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

Open
wants to merge 1 commit into
base: sycl
Choose a base branch
from
Open

[SYCL] Ignore unknown users of virtual functions instead of crashing #17397

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
5 changes: 4 additions & 1 deletion llvm/lib/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,8 @@
#include "llvm/IR/Operator.h"
#include "llvm/Pass.h"

#define DEBUG_TYPE "sycl-virtual-functions-analysis"

using namespace llvm;

namespace {
Expand Down Expand Up @@ -110,7 +112,8 @@ void collectVTablesThatUseFunction(
// ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4)), ...]
collectVTablesThatUseFunction(U, VTables);
} else {
llvm_unreachable("Unhandled type of user");
Copy link
Contributor

Choose a reason for hiding this comment

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

According to my understanding, llvm_unreachable is put here to catch cases which were missed during development. Ignoring "unhandled" cases might lead to bugs which are hard to analyze. The crash here is easy to analyze.

@Maetveis, do you have test case leading to running into llvm_unreachable? If so, please, add a regression test and one more if to handle it.

If would prefer to have llvm_unreachable("Unhandled type of user"); for unhandled users instead of silently ignoring them.

Copy link
Contributor Author

@Maetveis Maetveis Mar 11, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

See the attached test, it was reduced from basically:

struct C {
  /* sycl_indirectly_callable */
  void foo(){}
};

SYCL_EXTERNAL void direct_caller() {
  C c;
  c.foo();
}

Should this be added as test too?

I can add an isa<Instruction> branch if thats preferred.

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 it's preferred and will help us detecting incompatible LLVM community changes (e.g. adding more types of Users we need to handle).

Alternatively, we can consider replacing llvm_unreachable with an assert, but I'd like code owners/authors to comment on that first.

Copy link
Contributor Author

@Maetveis Maetveis Mar 12, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

The code is trying to look through constant expressions to find global variables that use the function pointers in their initializers. Personally, I don't really expect LLVM to start allowing non constants in global variable initializers. If other Users appear they are either going to be subclasses of Constant or probably won't appear in global initializers.
BTW looking at it again even the current code is redundant, ConstantExpr is a subclass of Constant.

In summary IMO having an allowlist here is unlikely to catch real bugs, but can easily be a source of churn or hidden crashes.

Copy link
Contributor

Choose a reason for hiding this comment

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

struct C {
  /* sycl_indirectly_callable */
  void foo(){}
};

SYCL_EXTERNAL void direct_caller() {
  C c;
  c.foo();
}

Should this be added as test too?

What is the form of sycl_indirectly_callable? Right know there is only a single (more or less properly) documented way of using indirect calls in SYCL - that is through virtual functions and the corresponding extension (#10540).

Copy link
Contributor Author

@Maetveis Maetveis Mar 19, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

What is the form of sycl_indirectly_callable? Right know there is only a single (more or less properly) documented way of using indirect calls in SYCL - that is through virtual functions and the corresponding extension (#10540).

I wasn't at my PC when I wrote that message, and didn't know the full syntax OTOH.
This is the full reproducer

// a.cpp
#include <sycl/sycl.hpp>
namespace syclext = sycl::ext::oneapi::experimental;
struct C {
  SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(syclext::indirectly_callable)
  virtual void foo(){}
};
SYCL_EXTERNAL void direct_caller() {
  C c;
  c.foo();
}
❯ build/bin/clang++ -fsycl -fsycl-targets=spir64 -O3 asd.cpp
Unhandled type of user
UNREACHABLE executed at llvm/llvm/lib/SYCLLowerIR/SYCLVirtualFunctionsAnalysis.cpp:145!
PLEASE append the compiler options "-save-temps -v", rebuild the application to get the full command which is failing and submit a bug report to https://software.intel.com/en-us/support/priority-support which includes the failing command, input files for the command and the crash backtrace (if any).
Stack dump:
(...)
1.      <eof> parser at end of file
2.      Optimizer
3.      Running pass "sycl-virtual-functions-analysis" on module "asd.cpp"
 #0 0x00007f4b2879b258 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (llvm/bin/../lib/libLLVMSupport.so.21.0git+0x267258)
 #1 0x00007f4b2879bbee SignalHandler(int) Signals.cpp:0:0
 #2 0x00007f4b28013520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #3 0x00007f4b280679fc __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
 #4 0x00007f4b280679fc __pthread_kill_internal ./nptl/pthread_kill.c:78:10
 #5 0x00007f4b280679fc pthread_kill ./nptl/pthread_kill.c:89:10
 #6 0x00007f4b28013476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #7 0x00007f4b27ff97f3 abort ./stdlib/abort.c:81:7
 #8 0x00007f4b286b807f (llvm/bin/../lib/libLLVMSupport.so.21.0git+0x18407f)
 #9 0x00007f4b2e9ad1e6 (anonymous namespace)::collectVTablesThatUseFunction(llvm::Value const*, llvm::SmallVectorImpl<llvm::GlobalVariable const*>&) SYCLVirtualFunctionsAnalysis.cpp:0:0
#10 0x00007f4b2e9ac0c7 llvm::SYCLVirtualFunctionsAnalysisPass::run(llvm::Module&, llvm::AnalysisManager<llvm::Module>&) (llvm/bin/../lib/libLLVMSYCLLowerIR.so.21.0git+0x12b0c7)

LLVM_DEBUG(dbgs() << "Ignoring unknown indirectly callable func user: '"
<< *U << "'\n");
}
}
}
Expand Down
36 changes: 36 additions & 0 deletions ...test/SYCLLowerIR/SYCLVirtualFunctionsAnalysis/calls-indirectly-propagation-other-users.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,36 @@
; RUN: opt -S -passes=sycl-virtual-functions-analysis %s | FileCheck %s
;
; Check if non-vtable uses of a function marked with "indirectly-callable" attribute
; are ignored. This used to crash.

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

@vtable = linkonce_odr dso_local unnamed_addr addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) null, ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4))] }, align 8

define linkonce_odr dso_local spir_func void @foo() unnamed_addr #0 {
ret void
}

define weak_odr dso_local spir_kernel void @kernel(ptr addrspace(1) noundef align 8 %_arg_StorageAcc) #1 {
entry:
store ptr addrspace(1) getelementptr inbounds inrange(-16, 8) (i8, ptr addrspace(1) @vtable, i64 16), ptr addrspace(1) %_arg_StorageAcc, align 8
ret void
}

define weak_odr dso_local spir_kernel void @kernel2() #1 {
entry:
; using the symbol in an instruction is ignored
%ptr = addrspacecast ptr @foo to ptr addrspace(4)
; as a special case of above using the symbol in a call should be ignored
call spir_func void @foo()
ret void
}

; CHECK: @kernel({{.*}} #[[#KERNEL_ATTRS:]]
; CHECK: @kernel2({{.*}} #[[#KERNEL2_ATTRS:]]
; CHECK: attributes #[[#KERNEL_ATTRS]] = {{.*}}"calls-indirectly"="set-foo"
; CHECK-NOT: attributes #[[#KERNEL2_ATTRS]] = {{.*}}"calls-indirectly"

attributes #0 = { "indirectly-callable"="set-foo" "sycl-module-id"="v.cpp" }
attributes #1 = { "sycl-module-id"="v.cpp" }
Loading