[NFCI][SYCL] Assert KernelHasName in StoreLambda

I don't understand what not having a name was supposed to mean and it
doesn't make any sense to me. The only failures I've seen after
implementing this change is when we compile SYCL code with non-SYCL
compiler without any SYCL knowledge (as in `clang++ -fsycl
-fsycl-host-compiler=<...>` does have SYCL knowledge provided via
integration headers).

If anyone knows why this change would be incorrect, please let me know.

Author: aelovikov-intel

sycl/include/sycl/handler.hpp

@@ -825,16 +825,24 @@ class __SYCL_EXPORT handler {
 #endif
     constexpr auto Info = detail::CompileTimeKernelInfo<KernelName>;
 
-    constexpr bool KernelHasName = (Info.Name != std::string_view{});
+    // SYCL unittests are built without sycl compiler, so "host" information
+    // about kernels isn't provided (e.g., via integration headers or compiler
+    // builtins).
+    //
+    // However, some copy/fill USM operation are implemented via SYCL kernels
+    // and are instantiated resulting in all the `static_assert` checks being
+    // exercised. Without kernel information that would fail, so we explicitly
+    // disable such checks when this macro is defined. Note that the unittests
+    // don't actually execute those operation, that's why disabling
+    // unconditional `static_asserts`s is enough for now.
+#ifndef __SYCL_UNITTESTS
+    static_assert(Info.Name != std::string_view{}, "Kernel must have a name!");
 
     // Some host compilers may have different captures from Clang. Currently
     // there is no stable way of handling this when extracting the captures,
     // so a static assert is made to fail for incompatible kernel lambdas.
-
-    // TODO remove the ifdef once the kernel size builtin is supported.
-#ifdef __INTEL_SYCL_USE_INTEGRATION_HEADERS
     static_assert(
-        !KernelHasName || sizeof(KernelType) == Info.KernelSize,
+        sizeof(KernelType) == Info.KernelSize,
         "Unexpected kernel lambda size. This can be caused by an "
         "external host compiler producing a lambda with an "
         "unexpected layout. This is a limitation of the compiler."
@@ -846,25 +854,13 @@ class __SYCL_EXPORT handler {
         "-fsycl-host-compiler-options='/std:c++latest' "
         "might also help.");
 #endif
-    // Empty name indicates that the compilation happens without integration
-    // header, so don't perform things that require it.
-    if constexpr (KernelHasName) {
-      // TODO support ESIMD in no-integration-header case too.
-
-      // Force hasSpecialCaptures to be evaluated at compile-time.
-      setKernelInfo((void *)MHostKernel->getPtr(), Info.NumParams,
-                    Info.ParamDescGetter, Info.IsESIMD,
-                    Info.HasSpecialCaptures);
-
-      MKernelName = Info.Name;
-      setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<KernelName>());
-    } else {
-      // In case w/o the integration header it is necessary to process
-      // accessors from the list(which are associated with this handler) as
-      // arguments. We must copy the associated accessors as they are checked
-      // later during finalize.
-      setArgsToAssociatedAccessors();
-    }
+
+    // Force hasSpecialCaptures to be evaluated at compile-time.
+    setKernelInfo((void *)MHostKernel->getPtr(), Info.NumParams,
+                  Info.ParamDescGetter, Info.IsESIMD, Info.HasSpecialCaptures);
+
+    MKernelName = Info.Name;
+    setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<KernelName>());
 
     // If the kernel lambda is callable with a kernel_handler argument, manifest
     // the associated kernel handler.

sycl/test/basic_tests/fp-accuracy.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx -%fsycl-host-only -c -ffp-accuracy=high -faltmathlib=SVMLAltMathLibrary -fno-math-errno %s
+// RUN: %clangxx -fsycl -c -ffp-accuracy=high -faltmathlib=SVMLAltMathLibrary -fno-math-errno %s
 
 #include <sycl/sycl.hpp>
 

sycl/test/basic_tests/single_task_error_message.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+// RUN: %clangxx -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
 #include <iostream>
 #include <sycl/sycl.hpp>
 int main() {
@@ -11,7 +11,7 @@ int main() {
       myQueue
           .single_task([&](sycl::handler &cgh) {
             // expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}}
-            // expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
+            // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{no matching function for call to object of type 'const (lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
           })
           .wait();
     }
@@ -27,7 +27,7 @@ int main() {
           .single_task(e,
                        [&](sycl::handler &cgh) {
                          // expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}}
-                         // expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
+                         // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{no matching function for call to object of type 'const (lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
                        })
           .wait();
     }
@@ -43,7 +43,7 @@ int main() {
           .single_task(vector_event,
                        [&](sycl::handler &cgh) {
                          // expected-error-re@sycl/queue.hpp:* {{static assertion failed due to requirement '{{.*}}': sycl::queue.single_task() requires a kernel instead of command group.{{.*}} Use queue.submit() instead}}
-                         // expected-error-re@sycl/detail/cg_types.hpp:* {{no matching function for call to object of type '(lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
+                         // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{no matching function for call to object of type 'const (lambda at {{.*}}single_task_error_message.cpp:{{.*}})'}}
                        })
           .wait();
     }

sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
+// RUN: %clangxx -D__SYCL_UNITTESTS %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
 
 #include <sycl/sycl.hpp>
 

sycl/test/virtual-functions/properties-negative.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
+// RUN: %clangxx -D__SYCL_UNITTESTS %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
 
 #include <sycl/sycl.hpp>
 

sycl/test/warnings/deprecated_get_backend_info.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+// RUN: %clangxx -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/kernel_bundle.hpp>

sycl/test/warnings/sycl_2020_deprecations.cpp

@@ -1,4 +1,4 @@
-// RUN: %clangxx %fsycl-host-only -fsyntax-only -ferror-limit=0 -sycl-std=2020 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+// RUN: %clangxx -fsycl-device-only -fsyntax-only -ferror-limit=0 -sycl-std=2020 -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
 
 // expected-warning@CL/sycl.hpp:* {{CL/sycl.hpp is deprecated, use sycl/sycl.hpp}}
 #include <CL/sycl.hpp>
@@ -283,23 +283,23 @@ int main() {
 
           // expected-warning@+8{{'get_pointer' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
           // expected-warning@+7{{'get_pointer<sycl::access::target::device, void>' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
-          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::global_space, sycl::access::decorated::legacy, void>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
+          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::global_space, sycl::access::decorated::legacy>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
           sycl::multi_ptr<int, sycl::access::address_space::global_space,
                           sycl::access::decorated::legacy>
               LegacyGlobalMptr =
                   sycl::make_ptr<int, sycl::access::address_space::global_space,
                                  sycl::access::decorated::legacy>(
                       GlobalAcc.get_pointer());
           // expected-warning@+7{{'get_pointer' is deprecated: local_accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
-          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy, void>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
+          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
           sycl::multi_ptr<int, sycl::access::address_space::local_space,
                           sycl::access::decorated::legacy>
               LegacyLocalMptr =
                   sycl::make_ptr<int, sycl::access::address_space::local_space,
                                  sycl::access::decorated::legacy>(
                       LocalAcc.get_pointer());
 
-          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy, void>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
+          // expected-warning@+4{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
           sycl::multi_ptr<int, sycl::access::address_space::private_space,
                           sycl::access::decorated::legacy>
               LegacyPrivateMptr =
@@ -329,27 +329,27 @@ int main() {
                           sycl::access::decorated::yes>
               UndecoratedPrivateMptr = DecoratedPrivateMptr;
 
-          // expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+          // expected-warning@+2{{'operator __global int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
           auto DecoratedGlobalPtr =
               static_cast<typename decltype(DecoratedGlobalMptr)::pointer>(
                   DecoratedGlobalMptr);
-          // expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+          // expected-warning@+2{{'operator __local int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
           auto DecoratedLocalPtr =
               static_cast<typename decltype(DecoratedLocalMptr)::pointer>(
                   DecoratedLocalMptr);
-          // expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+          // expected-warning@+2{{'operator __private int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
           auto DecoratedPrivatePtr =
               static_cast<typename decltype(DecoratedPrivateMptr)::pointer>(
                   DecoratedPrivateMptr);
-          // expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+          // expected-warning@+2{{'operator __global int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
           auto UndecoratedGlobalPtr =
               static_cast<typename decltype(UndecoratedGlobalMptr)::pointer>(
                   UndecoratedGlobalMptr);
-          // expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+          // expected-warning@+2{{'operator __local int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
           auto UndecoratedLocalPtr =
               static_cast<typename decltype(UndecoratedLocalMptr)::pointer>(
                   UndecoratedLocalMptr);
-          // expected-warning@+2{{'operator int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
+          // expected-warning@+2{{'operator __private int *' is deprecated: Conversion to pointer type is deprecated since SYCL 2020. Please use get() instead.}}
           auto UndecoratedPrivatePtr =
               static_cast<typename decltype(UndecoratedPrivateMptr)::pointer>(
                   UndecoratedPrivateMptr);

sycl/unittests/CMakeLists.txt

@@ -9,6 +9,8 @@ endforeach()
 
 add_compile_definitions(SYCL2020_DISABLE_DEPRECATION_WARNINGS SYCL_DISABLE_FSYCL_SYCLHPP_WARNING)
 
+add_compile_definitions(__SYCL_UNITTESTS)
+
 # suppress warnings which came from Google Test sources
 if (CXX_SUPPORTS_SUGGEST_OVERRIDE_FLAG)
   add_compile_options("-Wno-suggest-override")