[NFC][SYCL] Use CompileTimeKernelInfo instead of individual properties

Factored out from #19929.

Author: aelovikov-intel

sycl/include/sycl/detail/compile_time_kernel_info.hpp

@@ -13,38 +13,85 @@
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
-inline namespace compile_time_kernel_info_v1 {
 
+template <typename KernelNameType>
+constexpr kernel_param_desc_t getKernelParamDesc(int Idx) {
+#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
+  kernel_param_desc_t ParamDesc;
+  ParamDesc.kind =
+      __builtin_sycl_kernel_param_kind(KernelIdentity<KernelNameType>(), Idx);
+  ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor
+                       ? __builtin_sycl_kernel_param_access_target(
+                             KernelIdentity<KernelNameType>(), Idx)
+                       : __builtin_sycl_kernel_param_size(
+                             KernelIdentity<KernelNameType>(), Idx);
+  ParamDesc.offset =
+      __builtin_sycl_kernel_param_offset(KernelIdentity<KernelNameType>(), Idx);
+  return ParamDesc;
+#else
+  return KernelInfo<KernelNameType>::getParamDesc(Idx);
+#endif
+}
+
+inline namespace compile_time_kernel_info_v1 {
 // This is being passed across ABI boundary, so we don't use std::string_view,
 // at least for as long as we support user apps built with GNU libstdc++'s
 // pre-C++11 ABI.
 struct CompileTimeKernelInfoTy {
-  detail::string_view Name;
+  detail::string_view Name{};
   unsigned NumParams = 0;
   bool IsESIMD = false;
+  // TODO: Can we just have code_location here?
   detail::string_view FileName{};
   detail::string_view FunctionName{};
   unsigned LineNumber = 0;
   unsigned ColumnNumber = 0;
   int64_t KernelSize = 0;
   using ParamDescGetterT = kernel_param_desc_t (*)(int);
   ParamDescGetterT ParamDescGetter = nullptr;
-  bool HasSpecialCaptures = true;
+
+  bool HasSpecialCaptures = [this]() constexpr {
+    // No-compile time info for the kernel (i.e., kernel_bundle/interop/etc.),
+    // be conservative:
+    if (NumParams == 0)
+      return true;
+
+    for (unsigned I = 0; I < NumParams; ++I) {
+      auto ParamDesc = ParamDescGetter(I);
+      if (ParamDesc.kind != kernel_param_kind_t::kind_std_layout &&
+          ParamDesc.kind != kernel_param_kind_t::kind_pointer)
+        return true;
+    }
+
+    return false;
+  }();
 };
 
 template <class Kernel>
 inline constexpr CompileTimeKernelInfoTy CompileTimeKernelInfo{
-    std::string_view(getKernelName<Kernel>()),
-    getKernelNumParams<Kernel>(),
-    isKernelESIMD<Kernel>(),
-    std::string_view(getKernelFileName<Kernel>()),
-    std::string_view(getKernelFunctionName<Kernel>()),
-    getKernelLineNumber<Kernel>(),
-    getKernelColumnNumber<Kernel>(),
-    getKernelSize<Kernel>(),
-    &getKernelParamDesc<Kernel>,
-    hasSpecialCaptures<Kernel>()};
-
+#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
+    __builtin_sycl_kernel_name(KernelIdentity<Kernel>()),
+    __builtin_sycl_kernel_param_count(KernelIdentity<Kernel>()),
+    false /*IsESIMD*/, // TODO needs a builtin counterpart
+    __builtin_sycl_kernel_file_name(KernelIdentity<Kernel>()),
+    __builtin_sycl_kernel_function_name(KernelIdentity<Kernel>()),
+    __builtin_sycl_kernel_line_number(KernelIdentity<Kernel>()),
+    __builtin_sycl_kernel_column_number(KernelIdentity<Kernel>()),
+    // TODO needs a builtin counterpart, but is currently only used for checking
+    // cases with external host compiler, which use integration headers.
+    0 /* KernelSize */, &getKernelParamDesc<Kernel>
+#else
+    detail::string_view{KernelInfo<Kernel>::getName()},
+    KernelInfo<Kernel>::getNumParams(), KernelInfo<Kernel>::isESIMD(),
+    detail::string_view{KernelInfo<Kernel>::getFileName()},
+    detail::string_view{KernelInfo<Kernel>::getFunctionName()},
+    KernelInfo<Kernel>::getLineNumber(), KernelInfo<Kernel>::getColumnNumber(),
+    KernelInfo<Kernel>::getKernelSize(),
+    // Can't use KernelInfo::getParamDesc due to different return type (const
+    // ref vs. by val):
+    &getKernelParamDesc<Kernel>
+#endif
+};
 } // namespace compile_time_kernel_info_v1
 } // namespace detail
 } // namespace _V1

sycl/include/sycl/detail/kernel_desc.hpp

@@ -187,96 +187,6 @@ template <typename KNT> struct KernelIdentity {
   using type = KNT;
 };
 
-template <typename KernelNameType> constexpr unsigned getKernelNumParams() {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  return __builtin_sycl_kernel_param_count(KernelIdentity<KernelNameType>());
-#else
-  return KernelInfo<KernelNameType>::getNumParams();
-#endif
-}
-
-template <typename KernelNameType>
-constexpr kernel_param_desc_t getKernelParamDesc(int Idx) {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  kernel_param_desc_t ParamDesc;
-  ParamDesc.kind =
-      __builtin_sycl_kernel_param_kind(KernelIdentity<KernelNameType>(), Idx);
-  ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor
-                       ? __builtin_sycl_kernel_param_access_target(
-                             KernelIdentity<KernelNameType>(), Idx)
-                       : __builtin_sycl_kernel_param_size(
-                             KernelIdentity<KernelNameType>(), Idx);
-  ParamDesc.offset =
-      __builtin_sycl_kernel_param_offset(KernelIdentity<KernelNameType>(), Idx);
-  return ParamDesc;
-#else
-  return KernelInfo<KernelNameType>::getParamDesc(Idx);
-#endif
-}
-
-template <typename KernelNameType> constexpr const char *getKernelName() {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  return __builtin_sycl_kernel_name(KernelIdentity<KernelNameType>());
-#else
-  return KernelInfo<KernelNameType>::getName();
-#endif
-}
-
-template <typename KernelNameType> constexpr bool isKernelESIMD() {
-  // TODO Needs a builtin counterpart
-  return KernelInfo<KernelNameType>::isESIMD();
-}
-
-template <typename KernelNameType> constexpr const char *getKernelFileName() {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  return __builtin_sycl_kernel_file_name(KernelIdentity<KernelNameType>());
-#else
-  return KernelInfo<KernelNameType>::getFileName();
-#endif
-}
-
-template <typename KernelNameType>
-constexpr const char *getKernelFunctionName() {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  return __builtin_sycl_kernel_function_name(KernelIdentity<KernelNameType>());
-#else
-  return KernelInfo<KernelNameType>::getFunctionName();
-#endif
-}
-
-template <typename KernelNameType> constexpr unsigned getKernelLineNumber() {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  return __builtin_sycl_kernel_line_number(KernelIdentity<KernelNameType>());
-#else
-  return KernelInfo<KernelNameType>::getLineNumber();
-#endif
-}
-
-template <typename KernelNameType> constexpr unsigned getKernelColumnNumber() {
-#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
-  return __builtin_sycl_kernel_column_number(KernelIdentity<KernelNameType>());
-#else
-  return KernelInfo<KernelNameType>::getColumnNumber();
-#endif
-}
-
-template <typename KernelNameType> constexpr int64_t getKernelSize() {
-  // TODO needs a builtin counterpart, but is currently only used for checking
-  // cases with external host compiler, which use integration headers.
-  return KernelInfo<KernelNameType>::getKernelSize();
-}
-
-template <typename KernelNameType> constexpr bool hasSpecialCaptures() {
-  bool FoundSpecialCapture = false;
-  for (unsigned I = 0; I < getKernelNumParams<KernelNameType>(); ++I) {
-    auto ParamDesc = getKernelParamDesc<KernelNameType>(I);
-    bool IsSpecialCapture =
-        (ParamDesc.kind != kernel_param_kind_t::kind_std_layout &&
-         ParamDesc.kind != kernel_param_kind_t::kind_pointer);
-    FoundSpecialCapture |= IsSpecialCapture;
-  }
-  return FoundSpecialCapture;
-}
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/detail/kernel_launch_helper.hpp

@@ -9,6 +9,7 @@
 #pragma once
 
 #include <sycl/detail/cg_types.hpp>
+#include <sycl/detail/compile_time_kernel_info.hpp>
 #include <sycl/detail/helpers.hpp>
 #include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>
 #include <sycl/ext/intel/experimental/kernel_execution_properties.hpp>
@@ -261,7 +262,8 @@ struct KernelLaunchPropertyWrapper {
     if constexpr (ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<const KernelType &>::value) {
 
-      h->template processProperties<detail::isKernelESIMD<KernelName>()>(
+      h->template processProperties<
+          detail::CompileTimeKernelInfo<KernelName>.IsESIMD>(
           KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
     }
 #endif

sycl/include/sycl/handler.hpp

@@ -529,10 +529,8 @@ class __SYCL_EXPORT handler {
     // kernel. Else it is necessary use set_atg(s) for resolve the order and
     // values of arguments for the kernel.
     assert(MKernel && "MKernel is not initialized");
-    constexpr std::string_view LambdaName =
-        detail::getKernelName<LambdaNameT>();
     detail::ABINeutralKernelNameStrT KernelName = getKernelName();
-    return KernelName == LambdaName;
+    return KernelName == detail::CompileTimeKernelInfo<LambdaNameT>.Name;
   }
 
   /// Saves the location of user's code passed in \p CodeLoc for future usage in
@@ -823,20 +821,18 @@ class __SYCL_EXPORT handler {
         detail::GetInstantiateKernelOnHostPtr<KernelType, LambdaArgType,
                                               Dims>());
 #endif
+    constexpr auto Info = detail::CompileTimeKernelInfo<KernelName>;
 
-    constexpr bool KernelHasName =
-        detail::getKernelName<KernelName>() != nullptr &&
-        detail::getKernelName<KernelName>()[0] != '\0';
+    constexpr bool KernelHasName = (Info.Name != std::string_view{});
 
     // 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.
+    // 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) == detail::getKernelSize<KernelName>(),
+        !KernelHasName || 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."
@@ -854,15 +850,11 @@ class __SYCL_EXPORT handler {
       // TODO support ESIMD in no-integration-header case too.
 
       // Force hasSpecialCaptures to be evaluated at compile-time.
-      constexpr bool HasSpecialCapt = detail::hasSpecialCaptures<KernelName>();
-      setKernelInfo((void *)MHostKernel->getPtr(),
-                    detail::getKernelNumParams<KernelName>(),
-                    &(detail::getKernelParamDesc<KernelName>),
-                    detail::isKernelESIMD<KernelName>(), HasSpecialCapt);
-
-      constexpr std::string_view KernelNameStr =
-          detail::getKernelName<KernelName>();
-      MKernelName = KernelNameStr;
+      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
@@ -1240,6 +1232,12 @@ class __SYCL_EXPORT handler {
       typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
   void parallel_for_lambda_impl(range<Dims> UserRange, PropertiesT Props,
                                 const KernelType &KernelFunc) {
+    // TODO: Properties may change the kernel function, so in order to avoid
+    //       conflicts they should be included in the name.
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
+
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
     throwOnKernelParameterMisuse<KernelName, KernelType>();
@@ -1275,11 +1273,6 @@ class __SYCL_EXPORT handler {
         "SYCL kernel lambda/functor has an unexpected signature, it should be "
         "invocable with sycl::item and optionally sycl::kernel_handler");
 
-    // TODO: Properties may change the kernel function, so in order to avoid
-    //       conflicts they should be included in the name.
-    using NameT =
-        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-
     // Range rounding can be disabled by the user.
     // Range rounding is supported only for newer SYCL standards.
 #if !defined(__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__) &&                  \
@@ -1301,8 +1294,7 @@ class __SYCL_EXPORT handler {
       detail::KernelLaunchPropertyWrapper::parseProperties<KName>(this,
                                                                   Wrapper);
 #ifndef __SYCL_DEVICE_ONLY__
-      constexpr detail::string_view Name{detail::getKernelName<NameT>()};
-      verifyUsedKernelBundleInternal(Name);
+      verifyUsedKernelBundleInternal(Info.Name);
       // We are executing over the rounded range, but there are still
       // items/ids that are are constructed in ther range rounded
       // kernel use items/ids in the user range, which means that
@@ -1328,10 +1320,8 @@ class __SYCL_EXPORT handler {
       detail::KernelLaunchPropertyWrapper::parseProperties<NameT>(this,
                                                                   KernelFunc);
 #ifndef __SYCL_DEVICE_ONLY__
-      constexpr detail::string_view Name{detail::getKernelName<NameT>()};
-
-      verifyUsedKernelBundleInternal(Name);
-      processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
+      verifyUsedKernelBundleInternal(Info.Name);
+      processProperties<Info.IsESIMD, PropertiesT>(Props);
       detail::checkValueRange<Dims>(UserRange);
       setNDRangeDescriptor(std::move(UserRange));
       StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
@@ -1414,8 +1404,8 @@ class __SYCL_EXPORT handler {
       throwOnKernelParameterMisuse<KernelName, KernelType>();
     }
     throwIfActionIsCreated();
-    constexpr detail::string_view Name{detail::getKernelName<NameT>()};
-    verifyUsedKernelBundleInternal(Name);
+    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
+    verifyUsedKernelBundleInternal(Info.Name);
     setType(detail::CGType::Kernel);
 
     detail::checkValueRange<Dims>(params...);
@@ -1427,7 +1417,7 @@ class __SYCL_EXPORT handler {
     }
 
     StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
-    processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
+    processProperties<Info.IsESIMD, PropertiesT>(Props);
 #endif
   }
 
@@ -1460,8 +1450,8 @@ class __SYCL_EXPORT handler {
     // Ignore any set kernel bundles and use the one associated with the
     // kernel.
     setHandlerKernelBundle(Kernel);
-    constexpr detail::string_view Name{detail::getKernelName<NameT>()};
-    verifyUsedKernelBundleInternal(Name);
+    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
+    verifyUsedKernelBundleInternal(Info.Name);
     setType(detail::CGType::Kernel);
 
     detail::checkValueRange<Dims>(params...);
@@ -1479,7 +1469,7 @@ class __SYCL_EXPORT handler {
     } else {
       StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
     }
-    processProperties<detail::isKernelESIMD<NameT>(), PropertiesT>(Props);
+    processProperties<Info.IsESIMD, PropertiesT>(Props);
 #endif
   }
 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
@@ -1931,8 +1921,8 @@ class __SYCL_EXPORT handler {
     detail::KernelWrapperHelperFuncs::kernel_single_task<NameT>(KernelFunc);
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
-    constexpr detail::string_view Name{detail::getKernelName<NameT>()};
-    verifyUsedKernelBundleInternal(Name);
+    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
+    verifyUsedKernelBundleInternal(Info.Name);
     // No need to check if range is out of INT_MAX limits as it's compile-time
     // known constant
     setNDRangeDescriptor(range<1>{1});
@@ -3560,8 +3550,8 @@ class __SYCL_EXPORT handler {
   void throwOnKernelParameterMisuse() const {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
-    throwOnKernelParameterMisuseHelper(detail::getKernelNumParams<NameT>(),
-                                       &detail::getKernelParamDesc<NameT>);
+    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
+    throwOnKernelParameterMisuseHelper(Info.NumParams, Info.ParamDescGetter);
   }
 
   template <typename T, int Dims, access::mode AccessMode,

sycl/include/sycl/kernel_bundle.hpp

@@ -624,7 +624,7 @@ template <typename KernelName> kernel_id get_kernel_id() {
   // FIXME: This must fail at link-time if KernelName not in any available
   // translation units.
   return detail::get_kernel_id_impl(
-      detail::string_view{detail::getKernelName<KernelName>()});
+      detail::CompileTimeKernelInfo<KernelName>.Name);
 }
 
 /// \returns a vector with all kernel_id's defined in the application