[SYCL] Using a lambda in a kernel causes a linker crash

[krasznaa]

Dear All,

Similar to #2353 (but definitely not the same thing), I now ran into a linking issue while using lambdas. 😦

Take the following code:

// SYCL include(s).
#include <CL/sycl.hpp>

// System include(s).
#include <vector>

int main() {

   // Allocate a small buffer.
   static constexpr std::size_t BUFFER_SIZE = 1000;
   std::vector< float > testVector( BUFFER_SIZE, 1.23f );
   cl::sycl::buffer< float, 1 > testBuffer( testVector.data(), BUFFER_SIZE );

   // Set up the queue.
   cl::sycl::default_selector selector;
   cl::sycl::queue queue( selector );

   // Let the user know what device they're running on.
   cl::sycl::device device = queue.get_device();
   std::cout << "Running on device:" << std::endl;
   std::cout << "  Name  : "
             << device.get_info< cl::sycl::info::device::name >() << std::endl;
   std::cout << "  Vendor: "
             << device.get_info< cl::sycl::info::device::vendor >()
             << std::endl;

   // Set up a test lambda.
   auto testLambda = []( float value ) -> float {
      return value * 2.0f;
   };

   // A helper declaration.
   using am = cl::sycl::access::mode;

   // Submit a job on the buffer.
   const cl::sycl::range< 1 > workRange( BUFFER_SIZE );
   queue.submit( [&]( cl::sycl::handler& h ) {
         auto acc = testBuffer.get_access< am::read_write >( h );
         h.parallel_for< class func >( workRange,
                                       [=]( cl::sycl::id< 1 > id ) {
                                          if( id < BUFFER_SIZE ) {
                                             acc[ id ] =
                                                testLambda( acc[ id ] );
                                          }
                                       } );
      } );

   // Check the result.
   {
      auto acc = testBuffer.get_access< am::read >();
      std::cout << "acc[ 0 ] = " << acc[ 0 ] << std::endl;
   }

   return 0;
}

With a few-day-old version of the compiler, I get the following while trying to build it:

[bash][Elrond]:lambda > source /data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/setup.sh 
Configured CUDA from: /data/software/cuda/10.1.243/x86_64-ubuntu1804
Configured Clang from: /data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt
[bash][Elrond]:lambda > make
clang++ -fsycl -c -o test.o test.cxx
clang++ -fsycl -o test test.o
LLVM ERROR: llvm.memmove of non-constant length not supported
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace.
Stack dump:
0.      Program arguments: /data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv -o /tmp/test-50fa17-72de44.spv -spirv-max-version=1.1 -spirv-ext=+all,-SPV_INTEL_usm_storage_classes /tmp/test-0eb506.bc 
1.      Running pass 'Lower llvm.memmove into llvm.memcpy' on module '/tmp/test-0eb506.bc'.
 #0 0x00005613b2bfcd4a llvm::sys::PrintStackTrace(llvm::raw_ostream&) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x58cd4a)
 #1 0x00005613b2bfac14 llvm::sys::RunSignalHandlers() (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x58ac14)
 #2 0x00005613b2bfad58 SignalHandler(int) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x58ad58)
 #3 0x00007f6905d918a0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x128a0)
 #4 0x00007f6904c46f47 raise /build/glibc-2ORdQG/glibc-2.27/signal/../sysdeps/unix/sysv/linux/raise.c:51:0
 #5 0x00007f6904c488b1 abort /build/glibc-2ORdQG/glibc-2.27/stdlib/abort.c:81:0
 #6 0x00005613b2bc2426 llvm::report_fatal_error(llvm::Twine const&, bool) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x552426)
 #7 0x00005613b2bc2558 (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x552558)
 #8 0x00005613b2901ff5 SPIRV::SPIRVLowerMemmove::visitMemMoveInst(llvm::MemMoveInst&) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x291ff5)
 #9 0x00005613b2900998 SPIRV::SPIRVLowerMemmove::runOnModule(llvm::Module&) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x290998)
#10 0x00005613b2b48981 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x4d8981)
#11 0x00005613b28366c2 llvm::writeSpirv(llvm::Module*, SPIRV::TranslatorOpts const&, std::ostream&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x1c66c2)
#12 0x00005613b277f8a4 main (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x10f8a4)
#13 0x00007f6904c29b97 __libc_start_main /build/glibc-2ORdQG/glibc-2.27/csu/../csu/libc-start.c:344:0
#14 0x00005613b278edda _start (/data/software/intel/clang/12.0.0-2020-08-24/x86_64-ubuntu1804-gcc8-opt/bin/llvm-spirv+0x11edda)
llvm-foreach: Aborted (core dumped)
clang-12: error: llvm-spirv command failed with exit code 1 (use -v to see invocation)
Makefile:6: recipe for target 'test' failed
make: *** [test] Error 1
[bash][Elrond]:lambda >

I may be wrong, but I'm 99% sure that code like this did work in the past... 😕 Do you have any ideas on what's going wrong?

Cheers,
Attila

[Fznamznon]

@AlexeySachkov , @AlexeySotkin FYI

[AlexeySachkov]

Do you have any ideas on what's going wrong?

I bet that FE optimizations changed the IR in a way that we cannot translate it anymore. Lack of non-constant memove support is a known issue of the translator, @AlexeySotkin, could you please take a look as you already have a PR for non-constant memset in flight?

[illuhad]

Any update on this? The hipSYCL unit tests also trigger this ICE when compiling with hipSYCL's new SPIR-V backend (which basically uses the clang SYCL frontend with hipSYCL header files and runtime)

EDIT: In my case the error is slightly different, complaining about the size of the memcpy:

LLVM ERROR: Size of the memcpy should match the allocated memory
PLEASE submit a bug report to https://software.intel.com/en-us/support/priority-support and include the crash backtrace.
Stack dump:
0.      Program arguments: /opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv -o /tmp/hipSYCL-syclcc-_q35icp7/hipsycl-kernels.spv -spirv-max-version=1.1 -spirv-ext=+all,-SPV_INTEL_usm_storage_classes /tmp/math-c16338.bc 
1.      Running pass 'Lower llvm.memmove into llvm.memcpy' on module '/tmp/math-c16338.bc'.
 #0 0x000055e86bd19e37 llvm::sys::PrintStackTrace(llvm::raw_ostream&) (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x237e37)
 #1 0x000055e86bd18680 llvm::sys::RunSignalHandlers() (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x236680)
 #2 0x000055e86bd1a0f1 SignalHandler(int) (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x2380f1)
 #3 0x00007f22080603c0 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x153c0)
 #4 0x00007f2207e9f18b raise /build/glibc-eX1tMB/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1
 #5 0x00007f2207e7e859 abort /build/glibc-eX1tMB/glibc-2.31/stdlib/abort.c:81:7
 #6 0x000055e86bcffdea (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x21ddea)
 #7 0x000055e86bcffc47 (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x21dc47)
 #8 0x000055e86bfe3d1c (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x501d1c)
 #9 0x000055e86bfe34a9 SPIRV::SPIRVLowerMemmove::runOnModule(llvm::Module&) (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x5014a9)
#10 0x000055e86c09a848 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x5b8848)
#11 0x000055e86bf3310e llvm::writeSpirv(llvm::Module*, SPIRV::TranslatorOpts const&, std::__1::basic_ostream<char, std::__1::char_traits<char> >&, std::__1::basic_string<char, std::__1::char_traits<char>, std::__1::allocator<char> >&) (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x45110e)
#12 0x000055e86bec5898 main (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x3e3898)
#13 0x00007f2207e800b3 __libc_start_main /build/glibc-eX1tMB/glibc-2.31/csu/../csu/libc-start.c:342:3
#14 0x000055e86bec0393 _start (/opt/intel/oneapi/compiler/latest/linux/bin/llvm-spirv+0x3de393)
clang++: error: unable to execute command: Aborted (core dumped)
clang++: error: llvm-spirv command failed due to signal (use -v to see invocation)
Intel(R) oneAPI DPC++ Compiler 2021.1.2 (2020.10.0.1214)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/intel/oneapi/compiler/latest/linux/bin
clang++: note: diagnostic msg: 
********************

[Fznamznon]

I've tried test case from @krasznaa with the latest compiler sources and it fails with an assertion from clang FE. @elizabethandrews , @premanandrao , @AaronBallman , FYI.

clang++: llvm/llvm/lib/IR/Constants.cpp:2224: static llvm::Constant* llvm::ConstantExpr::getBitCast(llvm::Constant*, llvm::Type*, bool): Assertion `CastInst::castIsValid(Instruction::BitCast, C, DstTy) && "Invalid constantexpr bitcast!"' failed.
PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.
Stack dump:
0.	Program arguments: clang++ -fsycl -fsycl-device-only test1.cpp
1.	<eof> parser at end of file
2.	Per-file LLVM IR generation
3.	/usr/lib/gcc/x86_64-linux-gnu/7.5.0/../../../../include/x86_64-linux-gnu/c++/7.5.0/bits/gthr-default.h:247:1: Generating code for declaration '__gthread_active_p'
Stack dump without symbol names (ensure you have llvm-symbolizer in your PATH or set the environment var `LLVM_SYMBOLIZER_PATH` to point to it):
clang++(_ZN4llvm3sys15PrintStackTraceERNS_11raw_ostreamEi+0x2c)[0x55b5adda2ecc]
clang++(_ZN4llvm3sys17RunSignalHandlersEv+0x34)[0x55b5adda0d84]
clang++(_ZN4llvm3sys15CleanupOnSignalEm+0xb5)[0x55b5adda1005]
clang++(+0x2048d08)[0x55b5add02d08]
/lib/x86_64-linux-gnu/libpthread.so.0(+0x12890)[0x7f4ad3db5890]
/lib/x86_64-linux-gnu/libc.so.6(gsignal+0xc7)[0x7f4ad2a66e97]
/lib/x86_64-linux-gnu/libc.so.6(abort+0x141)[0x7f4ad2a68801]
/lib/x86_64-linux-gnu/libc.so.6(+0x3039a)[0x7f4ad2a5839a]
/lib/x86_64-linux-gnu/libc.so.6(+0x30412)[0x7f4ad2a58412]
clang++(+0x1973107)[0x55b5ad62d107]
clang++(_ZN5clang7CodeGen13CodeGenModule19GetWeakRefReferenceEPKNS_9ValueDeclE+0x1ab)[0x55b5ae19afbb]
clang++(_ZN5clang7CodeGen15ConstantEmitter14tryEmitPrivateERKNS_7APValueENS_8QualTypeE+0xe28)[0x55b5ae0c8e78]
clang++(_ZN5clang7CodeGen15ConstantEmitter12emitAbstractENS_14SourceLocationERKNS_7APValueENS_8QualTypeE+0x2d)[0x55b5ae0c94ed]
clang++(_ZN5clang7CodeGen15CodeGenFunction17tryEmitAsConstantEPNS_11DeclRefExprE+0x320)[0x55b5ae3a1890]
clang++(+0x27284b6)[0x55b5ae3e24b6]
clang++(+0xd62595)[0x55b5aca1c595]
clang++(+0x2727dea)[0x55b5ae3e1dea]
clang++(+0x2729776)[0x55b5ae3e3776]
clang++(+0x2729c68)[0x55b5ae3e3c68]
clang++(+0x2726cae)[0x55b5ae3e0cae]
clang++(+0xd62b25)[0x55b5aca1cb25]
clang++(+0x2727dea)[0x55b5ae3e1dea]
clang++(_ZN5clang7CodeGen15CodeGenFunction14EmitScalarExprEPKNS_4ExprEb+0x66)[0x55b5ae3e2bf6]
clang++(_ZN5clang7CodeGen15CodeGenFunction14EmitReturnStmtERKNS_10ReturnStmtE+0x7b5)[0x55b5ae0e0c35]
clang++(_ZN5clang7CodeGen15CodeGenFunction8EmitStmtEPKNS_4StmtEN4llvm8ArrayRefIPKNS_4AttrEEE+0x6ed)[0x55b5ae0e7bbd]
clang++(_ZN5clang7CodeGen15CodeGenFunction28EmitCompoundStmtWithoutScopeERKNS_12CompoundStmtEbNS0_12AggValueSlotE+0xcc)[0x55b5ae0e7dbc]
clang++(_ZN5clang7CodeGen15CodeGenFunction16EmitFunctionBodyEPKNS_4StmtE+0x8f)[0x55b5ae138b5f]
clang++(_ZN5clang7CodeGen15CodeGenFunction12GenerateCodeENS_10GlobalDeclEPN4llvm8FunctionERKNS0_14CGFunctionInfoE+0x8c3)[0x55b5ae14cbb3]
clang++(_ZN5clang7CodeGen13CodeGenModule28EmitGlobalFunctionDefinitionENS_10GlobalDeclEPN4llvm11GlobalValueE+0x1a3)[0x55b5ae19bcf3]
clang++(_ZN5clang7CodeGen13CodeGenModule20EmitGlobalDefinitionENS_10GlobalDeclEPN4llvm11GlobalValueE+0x265)[0x55b5ae1994b5]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x202)[0x55b5ae1a0d72]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule12EmitDeferredEv+0x288)[0x55b5ae1a0df8]
clang++(_ZN5clang7CodeGen13CodeGenModule7ReleaseEv+0x2c)[0x55b5ae1a1cac]
clang++(+0x308cd07)[0x55b5aed46d07]
clang++(+0x308b655)[0x55b5aed45655]
clang++(_ZN5clang8ParseASTERNS_4SemaEbb+0x499)[0x55b5afa2de09]
clang++(_ZN5clang13CodeGenAction13ExecuteActionEv+0x2b4)[0x55b5aed44504]
clang++(_ZN5clang14FrontendAction7ExecuteEv+0xe1)[0x55b5ae6e7291]
clang++(_ZN5clang16CompilerInstance13ExecuteActionERNS_14FrontendActionE+0x162)[0x55b5ae681092]
clang++(_ZN5clang25ExecuteCompilerInvocationEPNS_16CompilerInstanceE+0xb0a)[0x55b5ae7b558a]
clang++(_Z8cc1_mainN4llvm8ArrayRefIPKcEES2_Pv+0x120c)[0x55b5acb5a04c]
clang++(+0xe9b4a9)[0x55b5acb554a9]
clang++(+0x2866c35)[0x55b5ae520c35]
clang++(_ZN4llvm20CrashRecoveryContext9RunSafelyENS_12function_refIFvvEEE+0x23)[0x55b5add02de3]
clang++(+0x28682f4)[0x55b5ae5222f4]
clang++(_ZNK5clang6driver11Compilation14ExecuteCommandERKNS0_7CommandERPS3_+0x9a)[0x55b5ae4ec3ba]
clang++(_ZNK5clang6driver11Compilation11ExecuteJobsERKNS0_7JobListERN4llvm15SmallVectorImplISt4pairIiPKNS0_7CommandEEEE+0xae)[0x55b5ae4ece6e]
clang++(_ZN5clang6driver6Driver18ExecuteCompilationERNS0_11CompilationERN4llvm15SmallVectorImplISt4pairIiPKNS0_7CommandEEEE+0x9a)[0x55b5ae4f6a0a]
clang++(main+0x1c38)[0x55b5acad3f28]
/lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0xe7)[0x7f4ad2a49b97]
clang++(_start+0x2a)[0x55b5acb5501a]
clang-13: error: clang frontend command failed with exit code 134 (use -v to see invocation)
clang version 13.0.0 (https://github.com/intel/llvm.git 5b7d56afffbb90877ed8fec3007cd36828be7b6d)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: llvm/build/bin
clang-13: note: diagnostic msg: Error generating preprocessed source(s).

[AaronBallman]

When trying to reduce the example, I noticed that something strange is happening. I had gotten it down to:

// SYCL include(s).
#include <CL/sycl.hpp>

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
  kernelFunc(1);
}

int main() {
   // Allocate a small buffer.
   static constexpr std::size_t BUFFER_SIZE = 1000;

   // Set up the queue.
   cl::sycl::queue queue;

   // Submit a job on the buffer.
   kernel< class func >([=]( cl::sycl::id<1> id ) {
     (void)(id < BUFFER_SIZE);
   });

   return 0;
}

I could still reproduce the failed assertion with that. Notice how the queue object is unused? When I remove that... the assertion goes away. But if I leave queue alone and remove the kernel call instead, the crash goes away. If I change BUFFER_SIZE to a literal, the crash also goes away. If I change cl::sycl::id<1> to be an int, the crash also goes away.

[Fznamznon]

FYI: the translator's problem with llvm.memmove of non-constant length should be fixed now KhronosGroup/SPIRV-LLVM-Translator#1060 , the fix will come here with one of the next pull downs, it doesn't fix the front-end assertion fail though.