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

krasznaa · 2020-08-27T07:00:11Z

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

The text was updated successfully, but these errors were encountered:

Fznamznon · 2020-08-27T07:35:03Z

@AlexeySachkov , @AlexeySotkin FYI

AlexeySachkov · 2020-08-27T07:43:03Z

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 · 2021-03-12T20:30:41Z

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 · 2021-05-31T10:28:32Z

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 · 2021-06-01T12:27:43Z

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 · 2021-06-11T06:52:32Z

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.

Enable adapter tests to run on all discovered adapters.

AlexeySachkov assigned AlexeySotkin Aug 27, 2020

AlexeySachkov added the bug Something isn't working label Aug 27, 2020

AlexeySachkov added the SPIR-V Issues related to SPIRV-LLVM-Translator label Dec 25, 2020

illuhad mentioned this issue Mar 12, 2021

WIP: Add initial Intel GPU support / Level Zero runtime backend / SPIR-V compilation AdaptiveCpp/AdaptiveCpp#480

Merged

Fznamznon unassigned AlexeySotkin Jun 11, 2021

AlexeySachkov assigned premanandrao Jan 31, 2022

AlexeySachkov added confirmed and removed SPIR-V Issues related to SPIRV-LLVM-Translator labels Jan 31, 2022

Fznamznon assigned Fznamznon and unassigned premanandrao Feb 3, 2022

Fznamznon mentioned this issue Mar 29, 2022

[SYCL] Fix device code outlining for static local variables #5915

Merged

bader closed this as completed in #5915 Apr 1, 2022

Chenyang-L pushed a commit that referenced this issue Feb 18, 2025

Merge pull request #2376 from aarongreig/aaron/parameterizeAdapterTests

44da86d

Enable adapter tests to run on all discovered adapters.

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

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

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

krasznaa commented Aug 27, 2020

Fznamznon commented Aug 27, 2020

AlexeySachkov commented Aug 27, 2020

illuhad commented Mar 12, 2021 •

edited

Loading

Fznamznon commented May 31, 2021

AaronBallman commented Jun 1, 2021

Fznamznon commented Jun 11, 2021 •

edited

Loading

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

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

Comments

krasznaa commented Aug 27, 2020

Fznamznon commented Aug 27, 2020

AlexeySachkov commented Aug 27, 2020

illuhad commented Mar 12, 2021 • edited Loading

Fznamznon commented May 31, 2021

AaronBallman commented Jun 1, 2021

Fznamznon commented Jun 11, 2021 • edited Loading

illuhad commented Mar 12, 2021 •

edited

Loading

Fznamznon commented Jun 11, 2021 •

edited

Loading