From afcb732244ed55ff5fd903b87f49099b1a54103f Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 3 Sep 2025 12:02:39 -0700 Subject: [PATCH 1/4] [NFC][SYCL] Use `CompileTimeKernelInfo` instead of individual properties Factored out from https://github.com/intel/llvm/pull/19929. --- .../sycl/detail/compile_time_kernel_info.hpp | 75 +++++++++++++--- sycl/include/sycl/detail/kernel_desc.hpp | 90 ------------------- .../sycl/detail/kernel_launch_helper.hpp | 4 +- sycl/include/sycl/handler.hpp | 70 +++++++-------- sycl/include/sycl/kernel_bundle.hpp | 2 +- sycl/include/sycl/queue.hpp | 7 +- .../KernelParams/has-special-captures.cpp | 4 +- sycl/unittests/helpers/MockKernelInfo.hpp | 5 +- .../passing_link_and_compile_options.cpp | 6 +- 9 files changed, 106 insertions(+), 157 deletions(-) diff --git a/sycl/include/sycl/detail/compile_time_kernel_info.hpp b/sycl/include/sycl/detail/compile_time_kernel_info.hpp index f2eb59e874cd8..f5f7bb6d5b3d1 100644 --- a/sycl/include/sycl/detail/compile_time_kernel_info.hpp +++ b/sycl/include/sycl/detail/compile_time_kernel_info.hpp @@ -13,15 +13,35 @@ namespace sycl { inline namespace _V1 { namespace detail { -inline namespace compile_time_kernel_info_v1 { +template +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(), Idx); + ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor + ? __builtin_sycl_kernel_param_access_target( + KernelIdentity(), Idx) + : __builtin_sycl_kernel_param_size( + KernelIdentity(), Idx); + ParamDesc.offset = + __builtin_sycl_kernel_param_offset(KernelIdentity(), Idx); + return ParamDesc; +#else + return KernelInfo::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; @@ -29,22 +49,49 @@ struct CompileTimeKernelInfoTy { 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 inline constexpr CompileTimeKernelInfoTy CompileTimeKernelInfo{ - std::string_view(getKernelName()), - getKernelNumParams(), - isKernelESIMD(), - std::string_view(getKernelFileName()), - std::string_view(getKernelFunctionName()), - getKernelLineNumber(), - getKernelColumnNumber(), - getKernelSize(), - &getKernelParamDesc, - hasSpecialCaptures()}; - +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS + __builtin_sycl_kernel_name(KernelIdentity()), + __builtin_sycl_kernel_param_count(KernelIdentity()), + false /*IsESIMD*/, // TODO needs a builtin counterpart + __builtin_sycl_kernel_file_name(KernelIdentity()), + __builtin_sycl_kernel_function_name(KernelIdentity()), + __builtin_sycl_kernel_line_number(KernelIdentity()), + __builtin_sycl_kernel_column_number(KernelIdentity()), + // TODO needs a builtin counterpart, but is currently only used for checking + // cases with external host compiler, which use integration headers. + 0 /* KernelSize */, &getKernelParamDesc +#else + detail::string_view{KernelInfo::getName()}, + KernelInfo::getNumParams(), KernelInfo::isESIMD(), + detail::string_view{KernelInfo::getFileName()}, + detail::string_view{KernelInfo::getFunctionName()}, + KernelInfo::getLineNumber(), KernelInfo::getColumnNumber(), + KernelInfo::getKernelSize(), + // Can't use KernelInfo::getParamDesc due to different return type (const + // ref vs. by val): + &getKernelParamDesc +#endif +}; } // namespace compile_time_kernel_info_v1 } // namespace detail } // namespace _V1 diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index ae01f46c57052..2e6f5fdad5f80 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -187,96 +187,6 @@ template struct KernelIdentity { using type = KNT; }; -template constexpr unsigned getKernelNumParams() { -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - return __builtin_sycl_kernel_param_count(KernelIdentity()); -#else - return KernelInfo::getNumParams(); -#endif -} - -template -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(), Idx); - ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor - ? __builtin_sycl_kernel_param_access_target( - KernelIdentity(), Idx) - : __builtin_sycl_kernel_param_size( - KernelIdentity(), Idx); - ParamDesc.offset = - __builtin_sycl_kernel_param_offset(KernelIdentity(), Idx); - return ParamDesc; -#else - return KernelInfo::getParamDesc(Idx); -#endif -} - -template constexpr const char *getKernelName() { -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - return __builtin_sycl_kernel_name(KernelIdentity()); -#else - return KernelInfo::getName(); -#endif -} - -template constexpr bool isKernelESIMD() { - // TODO Needs a builtin counterpart - return KernelInfo::isESIMD(); -} - -template constexpr const char *getKernelFileName() { -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - return __builtin_sycl_kernel_file_name(KernelIdentity()); -#else - return KernelInfo::getFileName(); -#endif -} - -template -constexpr const char *getKernelFunctionName() { -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - return __builtin_sycl_kernel_function_name(KernelIdentity()); -#else - return KernelInfo::getFunctionName(); -#endif -} - -template constexpr unsigned getKernelLineNumber() { -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - return __builtin_sycl_kernel_line_number(KernelIdentity()); -#else - return KernelInfo::getLineNumber(); -#endif -} - -template constexpr unsigned getKernelColumnNumber() { -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - return __builtin_sycl_kernel_column_number(KernelIdentity()); -#else - return KernelInfo::getColumnNumber(); -#endif -} - -template 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::getKernelSize(); -} - -template constexpr bool hasSpecialCaptures() { - bool FoundSpecialCapture = false; - for (unsigned I = 0; I < getKernelNumParams(); ++I) { - auto ParamDesc = getKernelParamDesc(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 diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp index f5af740ad8751..254084a36adbb 100644 --- a/sycl/include/sycl/detail/kernel_launch_helper.hpp +++ b/sycl/include/sycl/detail/kernel_launch_helper.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -261,7 +262,8 @@ struct KernelLaunchPropertyWrapper { if constexpr (ext::oneapi::experimental::detail:: HasKernelPropertiesGetMethod::value) { - h->template processProperties()>( + h->template processProperties< + detail::CompileTimeKernelInfo.IsESIMD>( KernelFunc.get(ext::oneapi::experimental::properties_tag{})); } #endif diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a17358775e812..03f8b5cee1a0a 100644 --- a/sycl/include/sycl/handler.hpp +++ b/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(); detail::ABINeutralKernelNameStrT KernelName = getKernelName(); - return KernelName == LambdaName; + return KernelName == detail::CompileTimeKernelInfo.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()); #endif + constexpr auto Info = detail::CompileTimeKernelInfo; - constexpr bool KernelHasName = - detail::getKernelName() != nullptr && - detail::getKernelName()[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(), + !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(); - setKernelInfo((void *)MHostKernel->getPtr(), - detail::getKernelNumParams(), - &(detail::getKernelParamDesc), - detail::isKernelESIMD(), HasSpecialCapt); - - constexpr std::string_view KernelNameStr = - detail::getKernelName(); - MKernelName = KernelNameStr; + setKernelInfo((void *)MHostKernel->getPtr(), Info.NumParams, + Info.ParamDescGetter, Info.IsESIMD, + Info.HasSpecialCaptures); + + MKernelName = Info.Name; setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo()); } 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 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::name; + constexpr auto Info = detail::CompileTimeKernelInfo; + #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); throwOnKernelParameterMisuse(); @@ -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::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(this, Wrapper); #ifndef __SYCL_DEVICE_ONLY__ - constexpr detail::string_view Name{detail::getKernelName()}; - 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(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ - constexpr detail::string_view Name{detail::getKernelName()}; - - verifyUsedKernelBundleInternal(Name); - processProperties(), PropertiesT>(Props); + verifyUsedKernelBundleInternal(Info.Name); + processProperties(Props); detail::checkValueRange(UserRange); setNDRangeDescriptor(std::move(UserRange)); StoreLambda( @@ -1414,8 +1404,8 @@ class __SYCL_EXPORT handler { throwOnKernelParameterMisuse(); } throwIfActionIsCreated(); - constexpr detail::string_view Name{detail::getKernelName()}; - verifyUsedKernelBundleInternal(Name); + constexpr auto Info = detail::CompileTimeKernelInfo; + verifyUsedKernelBundleInternal(Info.Name); setType(detail::CGType::Kernel); detail::checkValueRange(params...); @@ -1427,7 +1417,7 @@ class __SYCL_EXPORT handler { } StoreLambda(std::move(KernelFunc)); - processProperties(), PropertiesT>(Props); + processProperties(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()}; - verifyUsedKernelBundleInternal(Name); + constexpr auto Info = detail::CompileTimeKernelInfo; + verifyUsedKernelBundleInternal(Info.Name); setType(detail::CGType::Kernel); detail::checkValueRange(params...); @@ -1479,7 +1469,7 @@ class __SYCL_EXPORT handler { } else { StoreLambda(std::move(KernelFunc)); } - processProperties(), PropertiesT>(Props); + processProperties(Props); #endif } #endif // __INTEL_PREVIEW_BREAKING_CHANGES @@ -1931,8 +1921,8 @@ class __SYCL_EXPORT handler { detail::KernelWrapperHelperFuncs::kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); - constexpr detail::string_view Name{detail::getKernelName()}; - verifyUsedKernelBundleInternal(Name); + constexpr auto Info = detail::CompileTimeKernelInfo; + 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::name; - throwOnKernelParameterMisuseHelper(detail::getKernelNumParams(), - &detail::getKernelParamDesc); + constexpr auto Info = detail::CompileTimeKernelInfo; + throwOnKernelParameterMisuseHelper(Info.NumParams, Info.ParamDescGetter); } template 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()}); + detail::CompileTimeKernelInfo.Name); } /// \returns a vector with all kernel_id's defined in the application diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 36006ac341bb6..ac0e5f477a0ad 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -3876,10 +3876,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { template static constexpr detail::code_location getCodeLocation() { - return {detail::getKernelFileName(), - detail::getKernelFunctionName(), - detail::getKernelLineNumber(), - detail::getKernelColumnNumber()}; + constexpr auto Info = detail::CompileTimeKernelInfo; + return {Info.FileName.data(), Info.FunctionName.data(), Info.LineNumber, + Info.ColumnNumber}; } }; diff --git a/sycl/test-e2e/KernelParams/has-special-captures.cpp b/sycl/test-e2e/KernelParams/has-special-captures.cpp index 94aa7b0644061..b7cc01e73b42e 100644 --- a/sycl/test-e2e/KernelParams/has-special-captures.cpp +++ b/sycl/test-e2e/KernelParams/has-special-captures.cpp @@ -18,7 +18,7 @@ int main() { Queue.parallel_for(nd_range<1>{1, 1}, [=](nd_item<1> Item) { *Pointer += Value; }); #ifndef __SYCL_DEVICE_ONLY__ - static_assert(!detail::hasSpecialCaptures()); + static_assert(!detail::CompileTimeKernelInfo.HasSpecialCaptures); #endif // An accessor is a special capture. @@ -28,6 +28,6 @@ int main() { Accessor[0] += Value; }); #ifndef __SYCL_DEVICE_ONLY__ - static_assert(detail::hasSpecialCaptures()); + static_assert(detail::CompileTimeKernelInfo.HasSpecialCaptures); #endif } diff --git a/sycl/unittests/helpers/MockKernelInfo.hpp b/sycl/unittests/helpers/MockKernelInfo.hpp index fd1b1ed9435a9..b65cf679450e4 100644 --- a/sycl/unittests/helpers/MockKernelInfo.hpp +++ b/sycl/unittests/helpers/MockKernelInfo.hpp @@ -15,9 +15,10 @@ namespace sycl { inline namespace _V1 { namespace unittest { struct MockKernelInfoBase { + static constexpr detail::kernel_param_desc_t Dummy{}; + static constexpr unsigned getNumParams() { return 0; } - static const detail::kernel_param_desc_t &getParamDesc(int) { - static detail::kernel_param_desc_t Dummy; + static constexpr const detail::kernel_param_desc_t &getParamDesc(int) { return Dummy; } static constexpr bool isESIMD() { return false; } diff --git a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp index 4deac120a133e..0ac843e753d41 100644 --- a/sycl/unittests/program_manager/passing_link_and_compile_options.cpp +++ b/sycl/unittests/program_manager/passing_link_and_compile_options.cpp @@ -19,15 +19,15 @@ std::string current_link_options, current_compile_options, current_build_opts; class EAMTestKernel1; -const char EAMTestKernelName1[] = "LinkCompileTestKernel1"; +constexpr const char EAMTestKernelName1[] = "LinkCompileTestKernel1"; constexpr unsigned EAMTestKernelNumArgs1 = 4; class EAMTestKernel2; -const char EAMTestKernelName2[] = "LinkCompileTestKernel2"; +constexpr const char EAMTestKernelName2[] = "LinkCompileTestKernel2"; constexpr unsigned EAMTestKernelNumArgs2 = 4; class EAMTestKernel3; -const char EAMTestKernelName3[] = "LinkCompileTestKernel3"; +constexpr const char EAMTestKernelName3[] = "LinkCompileTestKernel3"; constexpr unsigned EAMTestKernelNumArgs3 = 4; namespace sycl { From fa571ade569e5173e7476a6397288ab5532d144a Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 3 Sep 2025 13:11:08 -0700 Subject: [PATCH 2/4] `throwOnKernelParameterMisuse` changes --- sycl/include/sycl/handler.hpp | 35 ++++++++++++++++++++--------------- 1 file changed, 20 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 03f8b5cee1a0a..3c805a9e86462 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1240,7 +1240,7 @@ class __SYCL_EXPORT handler { #ifndef __SYCL_DEVICE_ONLY__ throwIfActionIsCreated(); - throwOnKernelParameterMisuse(); + throwOnKernelParameterMisuse(Info); if (!range_size_fits_in_size_t(UserRange)) throw sycl::exception(make_error_code(errc::runtime), "The total number of work-items in " @@ -1400,11 +1400,11 @@ class __SYCL_EXPORT handler { detail::KernelLaunchPropertyWrapper::parseProperties(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ + constexpr auto Info = detail::CompileTimeKernelInfo; if constexpr (WrapAsVal == detail::WrapAs::single_task) { - throwOnKernelParameterMisuse(); + throwOnKernelParameterMisuse(Info); } throwIfActionIsCreated(); - constexpr auto Info = detail::CompileTimeKernelInfo; verifyUsedKernelBundleInternal(Info.Name); setType(detail::CGType::Kernel); @@ -1443,14 +1443,14 @@ class __SYCL_EXPORT handler { detail::KernelLaunchPropertyWrapper::parseProperties(this, KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ + constexpr auto Info = detail::CompileTimeKernelInfo; if constexpr (WrapAsVal == detail::WrapAs::single_task) { - throwOnKernelParameterMisuse(); + throwOnKernelParameterMisuse(Info); } throwIfActionIsCreated(); // Ignore any set kernel bundles and use the one associated with the // kernel. setHandlerKernelBundle(Kernel); - constexpr auto Info = detail::CompileTimeKernelInfo; verifyUsedKernelBundleInternal(Info.Name); setType(detail::CGType::Kernel); @@ -3523,10 +3523,10 @@ class __SYCL_EXPORT handler { // // 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); + void throwOnKernelParameterMisuse( + const detail::CompileTimeKernelInfoTy &Info) const { + for (size_t I = 0; I < Info.NumParams; ++I) { + detail::kernel_param_desc_t ParamDesc = (*Info.ParamDescGetter)(I); const detail::kernel_param_kind_t &Kind = ParamDesc.kind; const access::target AccTarget = static_cast(ParamDesc.info & AccessTargetMask); @@ -3546,13 +3546,18 @@ class __SYCL_EXPORT handler { "of parallel_for that takes a range parameter."); } } - template - void throwOnKernelParameterMisuse() const { - using NameT = - typename detail::get_kernel_name_t::name; - constexpr auto Info = detail::CompileTimeKernelInfo; - throwOnKernelParameterMisuseHelper(Info.NumParams, Info.ParamDescGetter); + +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // Exported on Windows for some reason, have to keep for backward ABI + // compatibility, at least formally. + void throwOnKernelParameterMisuseHelper( + int N, detail::kernel_param_desc_t (*f)(int)) const { + detail::CompileTimeKernelInfoTy Info{}; + Info.NumParams = N; + Info.ParamDescGetter = f; + throwOnKernelParameterMisuse(Info); } +#endif template Date: Wed, 3 Sep 2025 13:51:31 -0700 Subject: [PATCH 3/4] Update win symbols --- sycl/test/abi/sycl_symbols_windows.dump | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index eae017c88eac8..bdea6ef268fc3 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4488,6 +4488,7 @@ ?supports_importing_handle_type@experimental@oneapi@ext@_V1@sycl@@YA_NW4external_mem_handle_type@12345@AEBVdevice@45@@Z ?sycl_category@_V1@sycl@@YAAEBVerror_category@std@@XZ ?throwIfActionIsCreated@handler@_V1@sycl@@AEAAXXZ +?throwOnKernelParameterMisuse@handler@_V1@sycl@@AEBAXAEBUCompileTimeKernelInfoTy@compile_time_kernel_info_v1@detail@23@@Z ?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 From 7072c71923851ac4c4784bb1a3d622d26fe34ed3 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 3 Sep 2025 19:51:38 -0700 Subject: [PATCH 4/4] No `operator==(detail::string_view, detail::string_view)` --- sycl/include/sycl/handler.hpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 3c805a9e86462..dc9f8fcf1146e 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -529,8 +529,10 @@ 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::CompileTimeKernelInfo.Name; detail::ABINeutralKernelNameStrT KernelName = getKernelName(); - return KernelName == detail::CompileTimeKernelInfo.Name; + return KernelName == LambdaName; } /// Saves the location of user's code passed in \p CodeLoc for future usage in