From c87f6b216181846fd3614abcfcebdf92a26c0bb1 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Tue, 10 Sep 2024 03:01:21 -0700 Subject: [PATCH] [SYCL] Use built-ins to retrieve kernel information (#15070) Using built-ins is going to be the preferred way to fetch kernel information, while integration headers are still going to be used for cases where built-ins are unavailable (i.e., different host compiler). Additionally, switch to the new entry point attribute when using the built-ins. --- sycl/include/sycl/detail/kernel_desc.hpp | 104 +++++++++++++++++++- sycl/include/sycl/handler.hpp | 115 ++++++++++++++--------- sycl/include/sycl/kernel_bundle.hpp | 4 +- sycl/include/sycl/queue.hpp | 38 +++----- sycl/source/handler.cpp | 23 +++-- sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/test/abi/sycl_symbols_windows.dump | 2 +- 7 files changed, 207 insertions(+), 80 deletions(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index 85519c3388efd..1049c4d78aadd 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -10,10 +10,26 @@ // FIXME: include export.hpp because integration header emitted by the compiler // uses the macro defined in this header, but it doesn't explicitly include it. +#include #include - // This header file must not include any standard C++ header files. +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS +#if __has_builtin(__builtin_sycl_kernel_name) +static_assert(__has_builtin(__builtin_sycl_kernel_param_count) && + __has_builtin(__builtin_sycl_kernel_name) && + __has_builtin(__builtin_sycl_kernel_param_access_target) && + __has_builtin(__builtin_sycl_kernel_param_size) && + __has_builtin(__builtin_sycl_kernel_param_offset) && + __has_builtin(__builtin_sycl_kernel_file_name) && + __has_builtin(__builtin_sycl_kernel_function_name) && + __has_builtin(__builtin_sycl_kernel_line_number) && + __has_builtin(__builtin_sycl_kernel_column_number)); +#else +#define __INTEL_SYCL_USE_INTEGRATION_HEADERS 1 +#endif +#endif + namespace sycl { inline namespace _V1 { namespace detail { @@ -151,6 +167,92 @@ template struct KernelInfo { }; #endif //__SYCL_UNNAMED_LAMBDA__ +// Built-ins accept an object due to lacking infrastructure support for +// accepting types. The kernel name type itself isn't used because it might be +// incomplete, cv-qualified, or not default constructible. Passing an object +// also allows future extension for SYCL kernels defined as free functions. +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 +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(); +} } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 6f2e9f9fc19b7..6181a41e6ef8c 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -429,6 +429,17 @@ template bool range_size_fits_in_size_t(const range &r) { } return true; } + +template +std::vector getKernelParamDescs() { + std::vector Result; + int NumParams = getKernelNumParams(); + Result.reserve(NumParams); + for (int I = 0; I < NumParams; ++I) { + Result.push_back(getKernelParamDesc(I)); + } + return Result; +} } // namespace detail /// Command group handler class. @@ -528,14 +539,12 @@ class __SYCL_EXPORT handler { void throwOnLocalAccessorMisuse() const { using NameT = typename detail::get_kernel_name_t::name; - using KI = sycl::detail::KernelInfo; - - auto *KernelArgs = &KI::getParamDesc(0); - - for (unsigned I = 0; I < KI::getNumParams(); ++I) { - const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind; + for (unsigned I = 0; I < detail::getKernelNumParams(); ++I) { + const detail::kernel_param_desc_t ParamDesc = + detail::getKernelParamDesc(I); + const detail::kernel_param_kind_t &Kind = ParamDesc.kind; const access::target AccTarget = - static_cast(KernelArgs[I].info & AccessTargetMask); + static_cast(ParamDesc.info & AccessTargetMask); if ((Kind == detail::kernel_param_kind_t::kind_accessor) && (AccTarget == target::local)) throw sycl::exception( @@ -546,8 +555,12 @@ class __SYCL_EXPORT handler { } } - /// Extracts and prepares kernel arguments from the lambda using integration - /// header. + /// Extracts and prepares kernel arguments from the lambda using information + /// from the built-ins or integration header. + void extractArgsAndReqsFromLambda( + char *LambdaPtr, + const std::vector &ParamDescs, bool IsESIMD); + // TODO Unused, remove during ABI breaking window void extractArgsAndReqsFromLambda(char *LambdaPtr, size_t KernelArgsNum, const detail::kernel_param_desc_t *KernelArgs, @@ -570,7 +583,7 @@ 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"); - const std::string LambdaName = detail::KernelInfo::getName(); + const std::string LambdaName = detail::getKernelName(); detail::string KernelName = getKernelName(); return KernelName == LambdaName; } @@ -885,21 +898,22 @@ class __SYCL_EXPORT handler { /// /// \param KernelName is the name of the SYCL kernel to check that the used /// kernel bundle contains. - void verifyUsedKernelBundle(const std::string &KernelName) { - verifyUsedKernelBundleInternal(detail::string_view{KernelName}); + template void verifyUsedKernelBundle() { + verifyUsedKernelBundleInternal( + detail::string_view{detail::getKernelName()}); } void verifyUsedKernelBundleInternal(detail::string_view KernelName); /// Stores lambda to the template-free object /// /// Also initializes kernel name, list of arguments and requirements using - /// information from the integration header. + /// information from the integration header/built-ins. /// - /// \param KernelFunc is a SYCL kernel function. + /// \param KernelFunc is a SYCL kernel function + /// \param ParamDescs is the vector of kernel parameter descriptors. template void StoreLambda(KernelType KernelFunc) { - using KI = detail::KernelInfo; constexpr bool IsCallableWithKernelHandler = detail::KernelLambdaHasKernelHandlerArgT::value; @@ -908,13 +922,18 @@ class __SYCL_EXPORT handler { ResetHostKernel(KernelFunc); constexpr bool KernelHasName = - KI::getName() != nullptr && KI::getName()[0] != '\0'; + detail::getKernelName() != nullptr && + detail::getKernelName()[0] != '\0'; // 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(KernelFunc) == KI::getKernelSize(), + !KernelHasName || + sizeof(KernelFunc) == detail::getKernelSize(), "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." @@ -925,16 +944,16 @@ class __SYCL_EXPORT handler { "In case of MSVC, passing " "-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 (KernelHasName) { // TODO support ESIMD in no-integration-header case too. clearArgs(); extractArgsAndReqsFromLambda(reinterpret_cast(KernelPtr), - KI::getNumParams(), &KI::getParamDesc(0), - KI::isESIMD()); - MKernelName = KI::getName(); + detail::getKernelParamDescs(), + detail::isKernelESIMD()); + MKernelName = detail::getKernelName(); } else { // In case w/o the integration header it is necessary to process // accessors from the list(which are associated with this handler) as @@ -1031,7 +1050,6 @@ class __SYCL_EXPORT handler { typename KernelName, typename PropertiesT = ext::oneapi::experimental::empty_properties_t> void processProperties(PropertiesT Props) { - using KI = detail::KernelInfo; static_assert( ext::oneapi::experimental::is_property_list::value, "Template type is not a property list."); @@ -1040,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>() && - KI::isESIMD()), + detail::isKernelESIMD()), "Floating point control property is supported for ESIMD kernels only."); static_assert( !PropertiesT::template has_property< @@ -1334,8 +1352,7 @@ class __SYCL_EXPORT handler { // conflicts they should be included in the name. using NameT = typename detail::get_kernel_name_t::name; - - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); // Range rounding can be disabled by the user. // Range rounding is not done on the host device. @@ -1417,7 +1434,7 @@ class __SYCL_EXPORT handler { // conflicts they should be included in the name. using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; static_assert( @@ -1507,7 +1524,7 @@ class __SYCL_EXPORT handler { // conflicts they should be included in the name. using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)NumWorkGroups; @@ -1548,7 +1565,7 @@ class __SYCL_EXPORT handler { // conflicts they should be included in the name. using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)NumWorkGroups; @@ -1568,10 +1585,14 @@ class __SYCL_EXPORT handler { } #ifdef SYCL_LANGUAGE_VERSION +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS +#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel_entry_point(KernelName)]] +#else #define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]] +#endif // __INTEL_SYCL_USE_INTEGRATION_HEADERS #else #define __SYCL_KERNEL_ATTR__ -#endif +#endif // SYCL_LANGUAGE_VERSION // NOTE: the name of this function - "kernel_single_task" - is used by the // Front End to determine kernel invocation kind. @@ -1583,7 +1604,9 @@ class __SYCL_EXPORT handler { nullptr, ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { + + __SYCL_KERNEL_ATTR__ static void + kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(); #else @@ -1601,8 +1624,8 @@ class __SYCL_EXPORT handler { nullptr, ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void kernel_single_task(_KERNELFUNCPARAM(KernelFunc), - kernel_handler KH) { + __SYCL_KERNEL_ATTR__ static void + kernel_single_task(_KERNELFUNCPARAM(KernelFunc), kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(KH); #else @@ -1620,7 +1643,8 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::name..., ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc)) { + __SYCL_KERNEL_ATTR__ static void + kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc)) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(detail::Builder::getElement(detail::declptr())); #else @@ -1637,8 +1661,8 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::name..., ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc), - kernel_handler KH) { + __SYCL_KERNEL_ATTR__ static void + kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc), kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(detail::Builder::getElement(detail::declptr()), KH); #else @@ -1656,7 +1680,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::name..., ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void + __SYCL_KERNEL_ATTR__ static void kernel_parallel_for_work_group(_KERNELFUNCPARAM(KernelFunc)) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(detail::Builder::getElement(detail::declptr())); @@ -1674,7 +1698,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::name..., ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void + __SYCL_KERNEL_ATTR__ static void kernel_parallel_for_work_group(_KERNELFUNCPARAM(KernelFunc), kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ @@ -1822,7 +1846,8 @@ class __SYCL_EXPORT handler { // conflicts they should be included in the name. using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + + verifyUsedKernelBundle(); kernel_single_task_wrapper(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ // No need to check if range is out of INT_MAX limits as it's compile-time @@ -2118,7 +2143,7 @@ class __SYCL_EXPORT handler { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; using TransformedArgType = std::conditional_t< std::is_integral::value && Dims == 1, item, @@ -2259,7 +2284,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); (void)Kernel; kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ @@ -2294,7 +2319,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; (void)NumWorkItems; @@ -2333,7 +2358,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; (void)NumWorkItems; @@ -2372,7 +2397,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; @@ -2415,7 +2440,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; @@ -2455,7 +2480,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 036bb6a3afe6a..1237bc0651b40 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -496,8 +496,8 @@ __SYCL_EXPORT kernel_id get_kernel_id_impl(string_view KernelName); template kernel_id get_kernel_id() { // FIXME: This must fail at link-time if KernelName not in any available // translation units. - using KI = sycl::detail::KernelInfo; - return detail::get_kernel_id_impl(detail::string_view{KI::getName()}); + return detail::get_kernel_id_impl( + detail::string_view{detail::getKernelName()}); } /// \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 f07e09db7a8b3..39f69046ad2aa 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2357,10 +2357,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { ext::oneapi::experimental::is_property_list::value, event> parallel_for(nd_range Range, PropertiesT Properties, RestT &&...Rest) { - using KI = sycl::detail::KernelInfo; - constexpr detail::code_location CodeLoc( - KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(), - KI::getColumnNumber()); + constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( [&](handler &CGH) { @@ -2393,10 +2390,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { template event parallel_for(nd_range Range, event DepEvent, RestT &&...Rest) { - using KI = sycl::detail::KernelInfo; - constexpr detail::code_location CodeLoc( - KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(), - KI::getColumnNumber()); + constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( [&](handler &CGH) { @@ -2418,10 +2412,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { typename... RestT> event parallel_for(nd_range Range, const std::vector &DepEvents, RestT &&...Rest) { - using KI = sycl::detail::KernelInfo; - constexpr detail::code_location CodeLoc( - KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(), - KI::getColumnNumber()); + constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( [&](handler &CGH) { @@ -2754,10 +2745,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { event> parallel_for_impl(range Range, PropertiesT Properties, RestT &&...Rest) { - using KI = sycl::detail::KernelInfo; - constexpr detail::code_location CodeLoc( - KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(), - KI::getColumnNumber()); + constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( [&](handler &CGH) { @@ -2791,10 +2779,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { ext::oneapi::experimental::is_property_list::value, event> parallel_for_impl(range Range, event DepEvent, PropertiesT Properties, RestT &&...Rest) { - using KI = sycl::detail::KernelInfo; - constexpr detail::code_location CodeLoc( - KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(), - KI::getColumnNumber()); + constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( [&](handler &CGH) { @@ -2831,10 +2816,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { ext::oneapi::experimental::is_property_list::value, event> parallel_for_impl(range Range, const std::vector &DepEvents, PropertiesT Properties, RestT &&...Rest) { - using KI = sycl::detail::KernelInfo; - constexpr detail::code_location CodeLoc( - KI::getFileName(), KI::getFunctionName(), KI::getLineNumber(), - KI::getColumnNumber()); + constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( [&](handler &CGH) { @@ -2869,6 +2851,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { size_t Offset, const std::vector &DepEvents); const property_list &getPropList() const; + + template + static constexpr detail::code_location getCodeLocation() { + return {detail::getKernelFileName(), + detail::getKernelFunctionName(), + detail::getKernelLineNumber(), + detail::getKernelColumnNumber()}; + } }; } // namespace _V1 diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index db50068328854..50e7d007e537e 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -842,16 +842,16 @@ void handler::extractArgsAndReqs() { } void handler::extractArgsAndReqsFromLambda( - char *LambdaPtr, size_t KernelArgsNum, - const detail::kernel_param_desc_t *KernelArgs, bool IsESIMD) { + char *LambdaPtr, const std::vector &ParamDescs, + bool IsESIMD) { const bool IsKernelCreatedFromSource = false; size_t IndexShift = 0; - impl->MArgs.reserve(MaxNumAdditionalArgs * KernelArgsNum); + impl->MArgs.reserve(MaxNumAdditionalArgs * ParamDescs.size()); - for (size_t I = 0; I < KernelArgsNum; ++I) { - void *Ptr = LambdaPtr + KernelArgs[I].offset; - const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind; - const int &Size = KernelArgs[I].info; + for (size_t I = 0; I < ParamDescs.size(); ++I) { + void *Ptr = LambdaPtr + ParamDescs[I].offset; + const detail::kernel_param_kind_t &Kind = ParamDescs[I].kind; + const int &Size = ParamDescs[I].info; if (Kind == detail::kernel_param_kind_t::kind_accessor) { // For args kind of accessor Size is information about accessor. // The first 11 bits of Size encodes the accessor target. @@ -875,6 +875,15 @@ void handler::extractArgsAndReqsFromLambda( } } +// TODO Unused, remove during ABI breaking window +void handler::extractArgsAndReqsFromLambda( + char *LambdaPtr, size_t KernelArgsNum, + const detail::kernel_param_desc_t *KernelArgs, bool IsESIMD) { + std::vector ParamDescs( + KernelArgs, KernelArgs + KernelArgsNum); + extractArgsAndReqsFromLambda(LambdaPtr, ParamDescs, IsESIMD); +} + // Calling methods of kernel_impl requires knowledge of class layout. // As this is impossible in header, there's a function that calls necessary // method inside the library and returns the result. diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index d9c1b61f7f691..4c73f43ed6ba2 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3514,6 +3514,7 @@ _ZN4sycl3_V17handler26setNDRangeDescriptorPaddedENS0_5rangeILi3EEES3_NS0_2idILi3 _ZN4sycl3_V17handler26setNDRangeDescriptorPaddedENS0_5rangeILi3EEEbi _ZN4sycl3_V17handler27addLifetimeSharedPtrStorageESt10shared_ptrIKvE _ZN4sycl3_V17handler27computeFallbackKernelBoundsEmm +_ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcRKSt6vectorINS0_6detail19kernel_param_desc_tESaIS5_EEb _ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_param_desc_tEb _ZN4sycl3_V17handler28memcpyToHostOnlyDeviceGlobalEPKvS3_mbmm _ZN4sycl3_V17handler28setArgsToAssociatedAccessorsEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 6f55e82a9151a..e2c3643c557be 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3862,6 +3862,7 @@ ?ext_oneapi_wait_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uexternal_semaphore@experimental@oneapi@ext@23@_KAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z ?ext_oneapi_wait_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uexternal_semaphore@experimental@oneapi@ext@23@_KV423@AEBUcode_location@detail@23@@Z ?extractArgsAndReqs@handler@_V1@sycl@@AEAAXXZ +?extractArgsAndReqsFromLambda@handler@_V1@sycl@@AEAAXPEADAEBV?$vector@Ukernel_param_desc_t@detail@_V1@sycl@@V?$allocator@Ukernel_param_desc_t@detail@_V1@sycl@@@std@@@std@@_N@Z ?extractArgsAndReqsFromLambda@handler@_V1@sycl@@AEAAXPEAD_KPEBUkernel_param_desc_t@detail@23@_N@Z ?fill_impl@handler@_V1@sycl@@AEAAXPEAXPEBX_K2@Z ?finalize@handler@_V1@sycl@@AEAA?AVevent@23@XZ @@ -4269,7 +4270,6 @@ ?updateValue@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEBVraw_kernel_arg@34567@_K@Z ?use_kernel_bundle@handler@_V1@sycl@@QEAAXAEBV?$kernel_bundle@$01@23@@Z ?verifyDeviceHasProgressGuarantee@handler@_V1@sycl@@AEAAXW4forward_progress_guarantee@experimental@oneapi@ext@23@W4execution_scope@56723@1@Z -?verifyUsedKernelBundle@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?verifyUsedKernelBundleInternal@handler@_V1@sycl@@AEAAXVstring_view@detail@23@@Z ?wait@event@_V1@sycl@@QEAAXXZ ?wait@event@_V1@sycl@@SAXAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@@Z