SYCL program build fails with --fsycl-host-compiler

[vpirogov]

Describe the bug

This issue was initially reported in uxlfoundation/oneDNN#2035

Investigation shows that attempt to build a SYCL program with --fsycl-host-compiler option fails with "error: ‘KernelInfoData’ is not a class template". This behavior can be reproduced with

To reproduce

Reproducer uses vector-add-buffers.cpp sample

Command:

clang++ -fsycl -fsycl-unnamed-lambda -fsycl-host-compiler=/usr/bin/g++ -fsycl-host-compiler-options="-std=c++17" vector-add-buffers.cpp

Output:

In file included from <command-line>:
/tmp/vector-add-buffers-header-805685.h:42:20: error: ‘KernelInfoData’ is not a class template
   42 | template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'N', '4', 's', 'y', 'c', 'l', '3', '_', 'V', '1', '6', 'd', 'e', 't', 'a', 'i', 'l', '1', '8', 'R', 'o', 'u', 'n', 'd', 'e', 'd', 'R', 'a', 'n', 'g', 'e', 'K', 'e', 'r', 'n', 'e', 'l', 'I', 'N', 'S', '0', '_', '4', 'i', 't', 'e', 'm', 'I', 'L', 'i', '1', 'E', 'L', 'b', '1', 'E', 'E', 'E', 'L', 'i', '1', 'E', 'Z', 'Z', '9', 'V', 'e', 'c', 't', 'o', 'r', 'A', 'd', 'd', 'R', 'N', 'S', '0', '_', '5', 'q', 'u', 'e', 'u', 'e', 'E', 'R', 'K', 'S', 't', '6', 'v', 'e', 'c', 't', 'o', 'r', 'I', 'i', 'S', 'a', 'I', 'i', 'E', 'E', 'S', 'B', '_', 'R', 'S', '9', '_', 'E', 'N', 'K', 'U', 'l', 'R', 'N', 'S', '0', '_', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', 'E', '_', 'E', 'U', 'l', 'T', '_', 'E', '_', 'E', 'E'> {
      |                    ^~~~~~~~~~~~~~
In file included from <command-line>:
/tmp/vector-add-buffers-header-805685.h:42:795: error: explicit specialization of non-template ‘sycl::_V1::detail::KernelInfoData’
   42 | template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'N', '4', 's', 'y', 'c', 'l', '3', '_', 'V', '1', '6', 'd', 'e', 't', 'a', 'i', 'l', '1', '8', 'R', 'o', 'u', 'n', 'd', 'e', 'd', 'R', 'a', 'n', 'g', 'e', 'K', 'e', 'r', 'n', 'e', 'l', 'I', 'N', 'S', '0', '_', '4', 'i', 't', 'e', 'm', 'I', 'L', 'i', '1', 'E', 'L', 'b', '1', 'E', 'E', 'E', 'L', 'i', '1', 'E', 'Z', 'Z', '9', 'V', 'e', 'c', 't', 'o', 'r', 'A', 'd', 'd', 'R', 'N', 'S', '0', '_', '5', 'q', 'u', 'e', 'u', 'e', 'E', 'R', 'K', 'S', 't', '6', 'v', 'e', 'c', 't', 'o', 'r', 'I', 'i', 'S', 'a', 'I', 'i', 'E', 'E', 'S', 'B', '_', 'R', 'S', '9', '_', 'E', 'N', 'K', 'U', 'l', 'R', 'N', 'S', '0', '_', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', 'E', '_', 'E', 'U', 'l', 'T', '_', 'E', '_', 'E', 'E'> {
      |                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                           ^
In file included from <command-line>:
/tmp/vector-add-buffers-header-805685.h:89:20: error: ‘KernelInfoData’ is not a class template
   89 | template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '9', 'V', 'e', 'c', 't', 'o', 'r', 'A', 'd', 'd', 'R', 'N', '4', 's', 'y', 'c', 'l', '3', '_', 'V', '1', '5', 'q', 'u', 'e', 'u', 'e', 'E', 'R', 'K', 'S', 't', '6', 'v', 'e', 'c', 't', 'o', 'r', 'I', 'i', 'S', 'a', 'I', 'i', 'E', 'E', 'S', '7', '_', 'R', 'S', '5', '_', 'E', 'N', 'K', 'U', 'l', 'R', 'N', 'S', '0', '_', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', 'A', '_', 'E', 'U', 'l', 'T', '_', 'E', '_'> {
      |                    ^~~~~~~~~~~~~~
In file included from <command-line>:
/tmp/vector-add-buffers-header-805685.h:89:503: error: ‘sycl::_V1::detail::KernelInfoData’ is not a template
   89 | template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '9', 'V', 'e', 'c', 't', 'o', 'r', 'A', 'd', 'd', 'R', 'N', '4', 's', 'y', 'c', 'l', '3', '_', 'V', '1', '5', 'q', 'u', 'e', 'u', 'e', 'E', 'R', 'K', 'S', 't', '6', 'v', 'e', 'c', 't', 'o', 'r', 'I', 'i', 'S', 'a', 'I', 'i', 'E', 'E', 'S', '7', '_', 'R', 'S', '5', '_', 'E', 'N', 'K', 'U', 'l', 'R', 'N', 'S', '0', '_', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', 'A', '_', 'E', 'U', 'l', 'T', '_', 'E', '_'> {
      |                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                                       ^
In file included from <command-line>:
/tmp/vector-add-buffers-header-805685.h:42:20: note: previous declaration here
   42 | template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'N', '4', 's', 'y', 'c', 'l', '3', '_', 'V', '1', '6', 'd', 'e', 't', 'a', 'i', 'l', '1', '8', 'R', 'o', 'u', 'n', 'd', 'e', 'd', 'R', 'a', 'n', 'g', 'e', 'K', 'e', 'r', 'n', 'e', 'l', 'I', 'N', 'S', '0', '_', '4', 'i', 't', 'e', 'm', 'I', 'L', 'i', '1', 'E', 'L', 'b', '1', 'E', 'E', 'E', 'L', 'i', '1', 'E', 'Z', 'Z', '9', 'V', 'e', 'c', 't', 'o', 'r', 'A', 'd', 'd', 'R', 'N', 'S', '0', '_', '5', 'q', 'u', 'e', 'u', 'e', 'E', 'R', 'K', 'S', 't', '6', 'v', 'e', 'c', 't', 'o', 'r', 'I', 'i', 'S', 'a', 'I', 'i', 'E', 'E', 'S', 'B', '_', 'R', 'S', '9', '_', 'E', 'N', 'K', 'U', 'l', 'R', 'N', 'S', '0', '_', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', 'E', '_', 'E', 'U', 'l', 'T', '_', 'E', '_', 'E', 'E'> {
      |                    ^~~~~~~~~~~~~~
In file included from /nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/queue.hpp:35,
                 from /nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/detail/core.hpp:23,
                 from /nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/sycl.hpp:11,
                 from vector-add-buffers.cpp:22:
/nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/handler.hpp:1586:76: warning: ‘clang::sycl_kernel’ scoped attribute directive ignored [-Wattributes]
 1586 |   __SYCL_KERNEL_ATTR__ void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) {
      |                                                                            ^
/nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/handler.hpp:1605:65: warning: ‘clang::sycl_kernel’ scoped attribute directive ignored [-Wattributes]
 1605 |                                                kernel_handler KH) {
      |                                                                 ^
/nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/handler.hpp:1623:77: warning: ‘clang::sycl_kernel’ scoped attribute directive ignored [-Wattributes]
 1623 |   __SYCL_KERNEL_ATTR__ void kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc)) {
      |                                                                             ^
/nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/handler.hpp:1641:66: warning: ‘clang::sycl_kernel’ scoped attribute directive ignored [-Wattributes]
 1641 |                                                 kernel_handler KH) {
      |                                                                  ^
/nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/handler.hpp:1660:62: warning: ‘clang::sycl_kernel’ scoped attribute directive ignored [-Wattributes]
 1660 |   kernel_parallel_for_work_group(_KERNELFUNCPARAM(KernelFunc)) {
      |                                                              ^
/nfs/pdx/disks/hal9000/vpirogov/bugs/onednn-2035/bin/../include/sycl/handler.hpp:1679:51: warning: ‘clang::sycl_kernel’ scoped attribute directive ignored [-Wattributes]
 1679 |                                  kernel_handler KH) {
      |                                                   ^

Environment

• OS: Linux
• Target device and vendor: Intel GPU
• DPC++ version: DPC++ daily 2024-08-16

clang version 19.0.0git (https://github.com/intel/llvm ef13a991cc6355f427d7eba089079d66da77f211)
Target: x86_64-unknown-linux-gnu
Thread model: posix
Build config: +assertions

• Dependencies version:

$ g++ --version
g++ (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0
Copyright (C) 2021 Free Software Foundation, Inc.
This is free software; see the source for copying conditions.  There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

Additional context

No response

[vpirogov]

Intel(R) oneAPI DPC++/C++ Compiler 2024.2.1 icpx driver throws an error for this case indicating that -fsycl-unnamed-lambda is not supported with -fsycl-host-compiler. It would be helpful if clang++ had similar diagnostics.

Updating the sample to use named kernel allows it to build correctly.

Original issue with oneDNN build does not seem to suffer from this scenario though...

[bader]

@mdtoguchi, FYI.

[vpirogov]

Another observation is that icx driver adds -fno-sycl-unnamed-lambda option implicitly when -fsycl-host-compiler is specified, but clang++ let's it through and fails the build.