Merge remote-tracking branches 'my/setDeviceKernelInfo' and 'my/test-asserts' into HEAD

Author: aelovikov-intel

sycl/include/sycl/handler.hpp

@@ -491,6 +491,25 @@ class __SYCL_EXPORT handler {
                             "a single kernel or explicit memory operation.");
   }
 
+  template <class Kernel> void setDeviceKernelInfo(void *KernelFuncPtr) {
+    constexpr auto Info = detail::CompileTimeKernelInfo<Kernel>;
+    MKernelName = Info.Name;
+    // TODO support ESIMD in no-integration-header case too.
+    setKernelInfo(KernelFuncPtr, Info.NumParams, Info.ParamDescGetter,
+                  Info.IsESIMD, Info.HasSpecialCaptures);
+    setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<Kernel>());
+    setType(detail::CGType::Kernel);
+  }
+
+  void setDeviceKernelInfo(kernel &&Kernel) {
+    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    MKernelName = getKernelName();
+    setType(detail::CGType::Kernel);
+
+    // If any extra actions are added here make sure that logic around
+    // `lambdaAndKernelHaveEqualName` calls can handle that.
+  }
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   // TODO: Those functions are not used anymore, remove it in the next
   // ABI-breaking window.
@@ -823,18 +842,27 @@ class __SYCL_EXPORT handler {
         detail::GetInstantiateKernelOnHostPtr<KernelType, LambdaArgType,
                                               Dims>());
 #endif
+
+    // 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
     constexpr auto Info = detail::CompileTimeKernelInfo<KernelName>;
 
-    constexpr bool KernelHasName = (Info.Name != std::string_view{});
+    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 +874,8 @@ 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();
-    }
+
+    setDeviceKernelInfo<KernelName>((void *)MHostKernel->getPtr());
 
     // If the kernel lambda is callable with a kernel_handler argument, manifest
     // the associated kernel handler.
@@ -1306,7 +1317,6 @@ class __SYCL_EXPORT handler {
       setNDRangeDescriptor(RoundedRange);
       StoreLambda<KName, decltype(Wrapper), Dims, TransformedArgType>(
           std::move(Wrapper));
-      setType(detail::CGType::Kernel);
 #endif
     } else
 #endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ &&
@@ -1328,7 +1338,6 @@ class __SYCL_EXPORT handler {
       setNDRangeDescriptor(std::move(UserRange));
       StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
           std::move(KernelFunc));
-      setType(detail::CGType::Kernel);
 #endif
 #else
       (void)KernelFunc;
@@ -1350,13 +1359,11 @@ class __SYCL_EXPORT handler {
                          [[maybe_unused]] kernel Kernel) {
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    setDeviceKernelInfo(std::move(Kernel));
     detail::checkValueRange<Dims>(NumWorkItems);
     setNDRangeDescriptor(std::move(NumWorkItems));
     processLaunchProperties<PropertiesT>(Props);
-    setType(detail::CGType::Kernel);
     extractArgsAndReqs();
-    MKernelName = getKernelName();
 #endif
   }
 
@@ -1375,13 +1382,11 @@ class __SYCL_EXPORT handler {
                          [[maybe_unused]] kernel Kernel) {
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    setDeviceKernelInfo(std::move(Kernel));
     detail::checkValueRange<Dims>(NDRange);
     setNDRangeDescriptor(std::move(NDRange));
     processLaunchProperties(Props);
-    setType(detail::CGType::Kernel);
     extractArgsAndReqs();
-    MKernelName = getKernelName();
 #endif
   }
 
@@ -1408,7 +1413,6 @@ class __SYCL_EXPORT handler {
     }
     throwIfActionIsCreated();
     verifyUsedKernelBundleInternal(Info.Name);
-    setType(detail::CGType::Kernel);
 
     detail::checkValueRange<Dims>(params...);
     if constexpr (SetNumWorkGroups) {
@@ -1454,7 +1458,6 @@ class __SYCL_EXPORT handler {
     // kernel.
     setHandlerKernelBundle(Kernel);
     verifyUsedKernelBundleInternal(Info.Name);
-    setType(detail::CGType::Kernel);
 
     detail::checkValueRange<Dims>(params...);
     if constexpr (SetNumWorkGroups) {
@@ -1464,12 +1467,11 @@ class __SYCL_EXPORT handler {
       setNDRangeDescriptor(std::move(params)...);
     }
 
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    if (!lambdaAndKernelHaveEqualName<NameT>()) {
-      extractArgsAndReqs();
-      MKernelName = getKernelName();
-    } else {
+    setDeviceKernelInfo(std::move(Kernel));
+    if (lambdaAndKernelHaveEqualName<NameT>()) {
       StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
+    } else {
+      extractArgsAndReqs();
     }
     processProperties<Info.IsESIMD, PropertiesT>(Props);
 #endif
@@ -1849,10 +1851,8 @@ class __SYCL_EXPORT handler {
     // No need to check if range is out of INT_MAX limits as it's compile-time
     // known constant
     setNDRangeDescriptor(range<1>{1});
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    setType(detail::CGType::Kernel);
+    setDeviceKernelInfo(std::move(Kernel));
     extractArgsAndReqs();
-    MKernelName = getKernelName();
   }
 
   void parallel_for(range<1> NumWorkItems, kernel Kernel) {
@@ -1885,12 +1885,10 @@ class __SYCL_EXPORT handler {
                     [[maybe_unused]] kernel Kernel) {
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    setDeviceKernelInfo(std::move(Kernel));
     detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
     setNDRangeDescriptor(std::move(NumWorkItems), std::move(WorkItemOffset));
-    setType(detail::CGType::Kernel);
     extractArgsAndReqs();
-    MKernelName = getKernelName();
 #endif
   }
 
@@ -1928,13 +1926,12 @@ class __SYCL_EXPORT handler {
     // No need to check if range is out of INT_MAX limits as it's compile-time
     // known constant
     setNDRangeDescriptor(range<1>{1});
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    setType(detail::CGType::Kernel);
-    if (!lambdaAndKernelHaveEqualName<NameT>()) {
-      extractArgsAndReqs();
-      MKernelName = getKernelName();
-    } else
+    setDeviceKernelInfo(std::move(Kernel));
+    if (lambdaAndKernelHaveEqualName<NameT>()) {
       StoreLambda<NameT, KernelType, /*Dims*/ 1, void>(std::move(KernelFunc));
+    } else {
+      extractArgsAndReqs();
+    }
 #else
     detail::CheckDeviceCopyable<KernelType>();
 #endif

sycl/source/detail/device_kernel_info.cpp

@@ -32,9 +32,13 @@ void DeviceKernelInfo::init(KernelNameStrRefT KernelName) {
 }
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-void DeviceKernelInfo::initIfNeeded(KernelNameStrRefT KernelName) {
-  if (!MInitialized.load())
-    init(KernelName);
+void DeviceKernelInfo::initIfEmpty(const CompileTimeKernelInfoTy &Info) {
+  if (MInitialized.load())
+    return;
+
+  CompileTimeKernelInfoTy::operator=(Info);
+  Name = Info.Name.data();
+  init(Name.data());
 }
 #endif
 

sycl/source/detail/device_kernel_info.hpp

@@ -102,7 +102,8 @@ class DeviceKernelInfo : public CompileTimeKernelInfoTy {
 
   void init(KernelNameStrRefT KernelName);
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-  void initIfNeeded(KernelNameStrRefT KernelName);
+  // Initialize default-created entry that has no data recorded:
+  void initIfEmpty(const CompileTimeKernelInfoTy &Info);
 #endif
   void setCompileTimeInfoIfNeeded(const CompileTimeKernelInfoTy &Info);
 

sycl/source/handler.cpp

@@ -544,7 +544,14 @@ event handler::finalize() {
   if (type == detail::CGType::Kernel) {
     if (impl->MDeviceKernelInfoPtr) {
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
-      impl->MDeviceKernelInfoPtr->initIfNeeded(toKernelNameStrT(MKernelName));
+      detail::CompileTimeKernelInfoTy HandlerInfo;
+      HandlerInfo.Name = MKernelName;
+      HandlerInfo.NumParams = impl->MKernelNumArgs;
+      HandlerInfo.ParamDescGetter = impl->MKernelParamDescGetter;
+      HandlerInfo.IsESIMD = impl->MKernelIsESIMD;
+      HandlerInfo.HasSpecialCaptures = impl->MKernelHasSpecialCaptures;
+
+      impl->MDeviceKernelInfoPtr->initIfEmpty(HandlerInfo);
 #endif
     } else {
       // Fetch the device kernel info pointer if it hasn't been set (e.g.
@@ -553,6 +560,20 @@ event handler::finalize() {
           &detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
               toKernelNameStrT(MKernelName));
     }
+
+    detail::DeviceKernelInfo &Info = *impl->MDeviceKernelInfoPtr;
+    (void)Info;
+
+    // Make sure that information set by old binaries gets properly copied to
+    // the `DeviceKernelInfo` so that the rest of the code base could use it as
+    // the single source of all the data:
+    assert(MKernelName == static_cast<std::string_view>(Info.Name));
+    assert(static_cast<unsigned>(impl->MKernelNumArgs) == Info.NumParams);
+    assert(impl->MKernelParamDescGetter == Info.ParamDescGetter ||
+           impl->MKernelParamDescGetter == nullptr);
+    assert(impl->MKernelIsESIMD == Info.IsESIMD);
+    assert(impl->MKernelHasSpecialCaptures == Info.HasSpecialCaptures);
+
     // If there were uses of set_specialization_constant build the kernel_bundle
     detail::kernel_bundle_impl *KernelBundleImpPtr =
         getOrInsertHandlerKernelBundlePtr(/*Insert=*/false);

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>