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

[Maetveis]

The SYCLVirtualFunctionsAnalysis pass was crashing (calling llvm_unreachable) when it encountered a non-constant non-global use of an indirectly callable function.
Instead of that just ignore any user we don't know how to handle and log a debug message. The LLVM User hierarchy is mostly closed, so in practice the ignored users are mostly going to be instructions.

[bader]

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.

[Maetveis]

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.

[bader]

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.

[Maetveis]

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.

[AlexeySachkov]

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).

[Maetveis]

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)

[Maetveis]

ping @intel/dpcpp-tools-reviewers

[Maetveis]

@intel/dpcpp-tools-reviewers Can you please look at this?

BTW I believe the test failures are unrelated: #17075

[Maetveis]

Ping.

[asudarsa]

Hi,

I had a couple of questions (Sorry, new to virtual functions support in SYCL).

1. Do we know if the SYCL program (that was provided as a comment in this PR) has all the required markings that are required to support SYCL virtual functions?
2. Do we know if the frontend device compiler (till the call to analyze virtual functions) lowers everything (w.r.t virtual functions) correctly and if the LLVM IR fed into the virtual function is as per expectations?

Thanks