ICE when array used in Functor

[agozillon]

#include <CL/sycl.hpp>

constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;

struct Functor {
  int X;
  float test[2];

  cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;

  Functor(
      int X_,
      cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_)
      : X(X_), Acc(Acc_) {}

  void operator()() { Acc[0] += X; }
};

int main() {
  int A[] = { 10 };

  {
    cl::sycl::queue Q;
    cl::sycl::buffer<int, 1> Buf(A, 1);

    Q.submit([&](cl::sycl::handler &cgh) {
      auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
      Functor F(1, Acc);

      cgh.single_task(F);
    });
  }

  return 0;
}

Stealing an example from the kernel_functor.cpp test in this case to show the issue.

If you use a raw array in a functor like above it will kill the compiler in the device outliner. It's likely the use case is illegal or unspecified in the Specification, but I imagine it shouldn't ICE the compiler.

Here's the compiler spew:

error: array initializer must be an initializer list
functor_test.cpp:31:11: note: in instantiation of function template specialization 'cl::sycl::handler::single_task<cl::sycl::detail::auto_name, Functor>' requested here
      cgh.single_task(F);
          ^
Stack dump:
0.	Program arguments: /storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10 -cc1 -triple spir64-unknown-linux-sycldevice -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -std=c++11 -disable-llvm-passes -sycl-std=1.2.1 -emit-llvm-bc -disable-free -disable-llvm-verifier -discard-value-names -main-file-name functor_test.cpp -mrelocation-model static -mthread-model posix -mframe-pointer=all -fmath-errno -mconstructor-aliases -fuse-init-array -dwarf-column-info -debugger-tuning=gdb -resource-dir /storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8/backward -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8/backward -internal-isystem /usr/local/include -internal-isystem /storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /usr/local/include -internal-isystem /storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -fdeprecated-macro -fdebug-compilation-dir /storage/ogozillo/intel-sycl/no-mod/sycl/sycl/test/functor -ferror-limit 19 -fmessage-length 0 -fobjc-runtime=gcc -fcxx-exceptions -fexceptions -fdiagnostics-show-option -faddrsig -o /tmp/functor_test-486772.o -x c++ functor_test.cpp 
1.	<eof> parser at end of file
2.	/storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0/include/CL/sycl/handler.hpp:668:8: instantiating function definition 'cl::sycl::handler::single_task<cl::sycl::detail::auto_name, Functor>'
 #0 0x000055dc47191eea llvm::sys::PrintStackTrace(llvm::raw_ostream&) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1753eea)
 #1 0x000055dc4718fc04 llvm::sys::RunSignalHandlers() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1751c04)
 #2 0x000055dc4718fd85 SignalHandler(int) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1751d85)
 #3 0x00007f8ab6128f40 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x13f40)
 #4 0x000055dc4940ea58 clang::InitListExpr::InitListExpr(clang::ASTContext const&, clang::SourceLocation, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x39d0a58)
 #5 0x000055dc48ffb780 CreateOpenCLKernelBody(clang::Sema&, clang::FunctionDecl*, clang::DeclContext*) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x35bd780)
 #6 0x000055dc48ffd651 clang::Sema::ConstructOpenCLKernel(clang::FunctionDecl*, clang::MangleContext&) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x35bf651)
 #7 0x000055dc49126bf7 clang::Sema::PerformPendingInstantiations(bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x36e8bf7)
 #8 0x000055dc4912786c clang::Sema::InstantiateFunctionDefinition(clang::SourceLocation, clang::FunctionDecl*, bool, bool, bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x36e986c)
 #9 0x000055dc49126867 clang::Sema::PerformPendingInstantiations(bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x36e8867)
#10 0x000055dc48b01605 clang::Sema::ActOnEndOfTranslationUnitFragment(clang::Sema::TUFragmentKind) (.part.1433) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x30c3605)
#11 0x000055dc48b017b9 clang::Sema::ActOnEndOfTranslationUnit() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x30c37b9)
#12 0x000055dc489efe51 clang::Parser::ParseTopLevelDecl(clang::OpaquePtr<clang::DeclGroupRef>&, bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x2fb1e51)
#13 0x000055dc489e3838 clang::ParseAST(clang::Sema&, bool, bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x2fa5838)
#14 0x000055dc47dff6f2 clang::CodeGenAction::ExecuteAction() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x23c16f2)
#15 0x000055dc47890329 clang::FrontendAction::Execute() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1e52329)
#16 0x000055dc47852ffb clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1e14ffb)
#17 0x000055dc47956eab clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1f18eab)
#18 0x000055dc4632e9c3 cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x8f09c3)
#19 0x000055dc462d326f main (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x89526f)
#20 0x00007f8ab5bb4b6b __libc_start_main /build/glibc-KRRWSm/glibc-2.29/csu/../csu/libc-start.c:342:3
#21 0x000055dc4632c88a _start (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x8ee88a)
clang-10: error: unable to execute command: Segmentation fault
error: array initializer must be an initializer list
functor_test.cpp:31:11: note: in instantiation of function template specialization 'cl::sycl::handler::single_task<cl::sycl::detail::auto_name, Functor>' requested here
      cgh.single_task(F);
          ^
Stack dump:
0.	Program arguments: /storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10 -cc1 -triple spir64-unknown-linux-sycldevice -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -std=c++11 -disable-llvm-passes -sycl-std=1.2.1 -fsyntax-only -disable-free -disable-llvm-verifier -discard-value-names -main-file-name functor_test.cpp -mrelocation-model static -mthread-model posix -mframe-pointer=all -fmath-errno -mconstructor-aliases -fuse-init-array -dwarf-column-info -debugger-tuning=gdb -resource-dir /storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8/backward -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/x86_64-linux-gnu/c++/8 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/8/../../../../include/c++/8/backward -internal-isystem /usr/local/include -internal-isystem /storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -internal-isystem /usr/local/include -internal-isystem /storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -fdeprecated-macro -fdebug-compilation-dir /storage/ogozillo/intel-sycl/no-mod/sycl/sycl/test/functor -ferror-limit 19 -fmessage-length 0 -fobjc-runtime=gcc -fcxx-exceptions -fexceptions -fdiagnostics-show-option -fsycl-int-header=/tmp/functor_test-7a9d11.h -faddrsig -o /tmp/functor_test-7a9d11.h -x c++ functor_test.cpp 
1.	<eof> parser at end of file
2.	/storage/ogozillo/intel-sycl/no-mod/sycl/build/lib/clang/10.0.0/include/CL/sycl/handler.hpp:668:8: instantiating function definition 'cl::sycl::handler::single_task<cl::sycl::detail::auto_name, Functor>'
 #0 0x00005620d0661eea llvm::sys::PrintStackTrace(llvm::raw_ostream&) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1753eea)
 #1 0x00005620d065fc04 llvm::sys::RunSignalHandlers() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1751c04)
 #2 0x00005620d065fd85 SignalHandler(int) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1751d85)
 #3 0x00007fc27eb97f40 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x13f40)
 #4 0x00005620d28dea58 clang::InitListExpr::InitListExpr(clang::ASTContext const&, clang::SourceLocation, llvm::ArrayRef<clang::Expr*>, clang::SourceLocation) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x39d0a58)
 #5 0x00005620d24cb780 CreateOpenCLKernelBody(clang::Sema&, clang::FunctionDecl*, clang::DeclContext*) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x35bd780)
 #6 0x00005620d24cd651 clang::Sema::ConstructOpenCLKernel(clang::FunctionDecl*, clang::MangleContext&) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x35bf651)
 #7 0x00005620d25f6bf7 clang::Sema::PerformPendingInstantiations(bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x36e8bf7)
 #8 0x00005620d25f786c clang::Sema::InstantiateFunctionDefinition(clang::SourceLocation, clang::FunctionDecl*, bool, bool, bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x36e986c)
 #9 0x00005620d25f6867 clang::Sema::PerformPendingInstantiations(bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x36e8867)
#10 0x00005620d1fd1605 clang::Sema::ActOnEndOfTranslationUnitFragment(clang::Sema::TUFragmentKind) (.part.1433) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x30c3605)
#11 0x00005620d1fd17b9 clang::Sema::ActOnEndOfTranslationUnit() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x30c37b9)
#12 0x00005620d1ebfe51 clang::Parser::ParseTopLevelDecl(clang::OpaquePtr<clang::DeclGroupRef>&, bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x2fb1e51)
#13 0x00005620d1eb3838 clang::ParseAST(clang::Sema&, bool, bool) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x2fa5838)
#14 0x00005620d0d60329 clang::FrontendAction::Execute() (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1e52329)
#15 0x00005620d0d22ffb clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1e14ffb)
#16 0x00005620d0e26eab clang::ExecuteCompilerInvocation(clang::CompilerInstance*) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x1f18eab)
#17 0x00005620cf7fe9c3 cc1_main(llvm::ArrayRef<char const*>, char const*, void*) (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x8f09c3)
#18 0x00005620cf7a326f main (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x89526f)
#19 0x00007fc27e623b6b __libc_start_main /build/glibc-KRRWSm/glibc-2.29/csu/../csu/libc-start.c:342:3
#20 0x00005620cf7fc88a _start (/storage/ogozillo/intel-sycl/no-mod/sycl/build/bin/clang-10+0x8ee88a)
clang-10: error: unable to execute command: Segmentation fault
clang-10: error: clang frontend command failed due to signal (use -v to see invocation)
clang-10: error: clang frontend command failed due to signal (use -v to see invocation)
clang version 10.0.0 (https://github.com/intel/llvm.git 1edd21722794ddd72d5390e36a04db786fc68019)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /storage/ogozillo/intel-sycl/no-mod/sycl/build/bin
clang-10: note: diagnostic msg: PLEASE submit a bug report to https://bugs.llvm.org/ and include the crash backtrace, preprocessed source, and associated run script.
clang-10: note: diagnostic msg: Error generating preprocessed source(s).

[Fznamznon]

From looking to the log I assume that the problem in our new initialization of kernel object, #488 also was caused by it. I'll take a look what we can do with it...

[Fznamznon]

I have feeling that #1423 will fix this issue. @rdeodhar , is it possible to try code sample from the description with your patch?

[rdeodhar]

Yes, the program compiles successfully with PR 1423 applied, and by adding prints I confirmed that it runs correctly too.