[SYCL] Split throwOnKernelParameterMisuse to improve compile time (#17670)

For the following code (note no submit, just bare minimum
`cgh.single_task`):

```

int main() {
  int *p;
  sycl::detail::loop<2>([&](auto outer_idx) {
    sycl::detail::loop<200>([&](auto idx) {
      auto krn = [=]() { *p = 42; };
      auto s = [&](sycl::handler &cgh) { cgh.single_task(krn); };
      (void)sycl::detail::type_erased_cgfo_ty{s};
    });
  });
}
```

compiled as

` $ time clang++ -isystem ~/sycl/build/include a.cpp -c -o /dev/null
-Wno-deprecated`

(to simulate host-only compilation), improves 3.4s -> 3.0s.

Author: aelovikov-intel

sycl/include/sycl/handler.hpp

@@ -3779,13 +3779,13 @@ class __SYCL_EXPORT handler {
   /// According to section 4.7.6.11. of the SYCL specification, a local accessor
   /// must not be used in a SYCL kernel function that is invoked via single_task
   /// or via the simple form of parallel_for that takes a range parameter.
-  template <typename KernelName, typename KernelType>
-  void throwOnKernelParameterMisuse() const {
-    using NameT =
-        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-    for (unsigned I = 0; I < detail::getKernelNumParams<NameT>(); ++I) {
-      const detail::kernel_param_desc_t ParamDesc =
-          detail::getKernelParamDesc<NameT>(I);
+  //
+  // Exception handling generates lots of code, outline it out of template
+  // method to improve compilation times.
+  void throwOnKernelParameterMisuseHelper(
+      int N, detail::kernel_param_desc_t (*f)(int)) const {
+    for (int I = 0; I < N; ++I) {
+      detail::kernel_param_desc_t ParamDesc = (*f)(I);
       const detail::kernel_param_kind_t &Kind = ParamDesc.kind;
       const access::target AccTarget =
           static_cast<access::target>(ParamDesc.info & AccessTargetMask);
@@ -3805,6 +3805,13 @@ class __SYCL_EXPORT handler {
             "of parallel_for that takes a range parameter.");
     }
   }
+  template <typename KernelName, typename KernelType>
+  void throwOnKernelParameterMisuse() const {
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    throwOnKernelParameterMisuseHelper(detail::getKernelNumParams<NameT>(),
+                                       &detail::getKernelParamDesc<NameT>);
+  }
 
   template <typename T, int Dims, access::mode AccessMode,
             access::target AccessTarget,

sycl/test/abi/sycl_symbols_windows.dump

@@ -4408,6 +4408,7 @@
 ?supportsUSMMemset2D@handler@_V1@sycl@@AEAA_NXZ
 ?sycl_category@_V1@sycl@@YAAEBVerror_category@std@@XZ
 ?throwIfActionIsCreated@handler@_V1@sycl@@AEAAXXZ
+?throwOnKernelParameterMisuseHelper@handler@_V1@sycl@@AEBAXHP6A?AUkernel_param_desc_t@detail@23@H@Z@Z
 ?throw_asynchronous@queue@_V1@sycl@@QEAAXXZ
 ?unmap@experimental@oneapi@ext@_V1@sycl@@YAXPEBX_KAEBVcontext@45@@Z
 ?unsampledImageConstructorNotification@detail@_V1@sycl@@YAXPEAX0AEBV?$optional@W4image_target@_V1@sycl@@@std@@W4mode@access@23@PEBXIAEBUcode_location@123@@Z