[SYCL] Optimize handling of compile-time properties

We perform correctness check of compile-time properties in
`handler::processProperties` helper function. That function was
specialized by kernel name, meaning that if we have two kernels with
absolutely the same properties applied to them there will be two
instantiations of `processProperties`. That consumes compilation time
for both front-end which has to emit extra isntantiations and for host
compilation pass which gets more functions (with the same body!) to
handle.

The only use of kernel name within `processProperties` is to check if
that kernel is a ESIMD kernel. That can be done at caller side and
propagated to `processProperties` as a simple boolean, thus reducing
amount of instantiations of that function.

This patch does exactly that.

Please note that this is technically a *functional* change: even though
we still process the properties in the same way, we will now emit
diagnostics in a slightly different manner: if there are two kernels
with the same set of illegal properties there will be only one
diagnostic for the first kernel.
Once that first kernel is fixed, the diagnostic will be displayed for
the second kernel, i.e. we won't display _all_ violations in a single
compilation run.

It seems like we only have one test which exposes that behavior and
considering that with C++ it is almost always you should fix the first
error first to see what happens to the rest, I don't think that such
change in diagnostics is bad enough to outweight potential (even if they
are the slightest) compilation time improvements.

Author: AlexeySachkov

sycl/include/sycl/handler.hpp

@@ -1047,7 +1047,7 @@ class __SYCL_EXPORT handler {
   ///
   /// Stores information about kernel properties into the handler.
   template <
-      typename KernelName,
+      bool IsESIMDKernel,
       typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
   void processProperties(PropertiesT Props) {
     static_assert(
@@ -1058,7 +1058,7 @@ class __SYCL_EXPORT handler {
             sycl::ext::intel::experimental::fp_control_key>() ||
             (PropertiesT::template has_property<
                  sycl::ext::intel::experimental::fp_control_key>() &&
-             detail::isKernelESIMD<KernelName>()),
+             IsESIMDKernel),
         "Floating point control property is supported for ESIMD kernels only.");
     static_assert(
         !PropertiesT::template has_property<
@@ -1398,7 +1398,7 @@ class __SYCL_EXPORT handler {
       kernel_parallel_for_wrapper<NameT, TransformedArgType, KernelType,
                                   PropertiesT>(KernelFunc);
 #ifndef __SYCL_DEVICE_ONLY__
-      processProperties<NameT, PropertiesT>(Props);
+      processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
       detail::checkValueRange<Dims>(UserRange);
       setNDRangeDescriptor(std::move(UserRange));
       StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
@@ -1450,7 +1450,7 @@ class __SYCL_EXPORT handler {
 #ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(ExecutionRange);
     setNDRangeDescriptor(std::move(ExecutionRange));
-    processProperties<NameT, PropertiesT>(Props);
+    processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
     StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
         std::move(KernelFunc));
     setType(detail::CGType::Kernel);
@@ -1532,7 +1532,7 @@ class __SYCL_EXPORT handler {
     kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType, KernelType,
                                            PropertiesT>(KernelFunc);
 #ifndef __SYCL_DEVICE_ONLY__
-    processProperties<NameT, PropertiesT>(Props);
+    processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
     detail::checkValueRange<Dims>(NumWorkGroups);
     setNDRangeDescriptor(NumWorkGroups, /*SetNumWorkGroups=*/true);
     StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
@@ -1574,7 +1574,7 @@ class __SYCL_EXPORT handler {
     kernel_parallel_for_work_group_wrapper<NameT, LambdaArgType, KernelType,
                                            PropertiesT>(KernelFunc);
 #ifndef __SYCL_DEVICE_ONLY__
-    processProperties<NameT, PropertiesT>(Props);
+    processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
     nd_range<Dims> ExecRange =
         nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
     detail::checkValueRange<Dims>(ExecRange);
@@ -1774,7 +1774,7 @@ class __SYCL_EXPORT handler {
     if constexpr (ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
                           _KERNELFUNCPARAMTYPE>::value) {
-      processProperties<KernelName>(
+      processProperties<detail::isKernelESIMD<KernelName>()>(
           KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
     }
 #endif
@@ -1853,7 +1853,7 @@ 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});
-    processProperties<NameT, PropertiesT>(Props);
+    processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
     StoreLambda<NameT, KernelType, /*Dims*/ 1, void>(KernelFunc);
     setType(detail::CGType::Kernel);
 #endif

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

@@ -19,7 +19,9 @@ int main() {
 
   // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
   q.single_task(props_empty, [=]() {});
-  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  // When both "props_empty" and "props_void" are in use, we won't see the
+  // static assert firing for the second one, because there will be only one
+  // instantiation of handler::processProperties.
   q.single_task(props_void, [=]() {});
   // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
   q.single_task(props_int, [=]() {});