From 558dfea81d816f9a1729c4c0966445a4a7e5bc1d Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Tue, 6 Aug 2024 02:38:54 -0700 Subject: [PATCH 01/18] [SYCL] Use built-ins to retrieve kernel information --- sycl/include/sycl/detail/kernel_desc.hpp | 103 ++++++++++++++++++++++- sycl/include/sycl/handler.hpp | 94 +++++++++++---------- sycl/include/sycl/kernel_bundle.hpp | 4 +- sycl/include/sycl/queue.hpp | 38 +++------ sycl/source/handler.cpp | 35 ++++++++ sycl/test/abi/sycl_symbols_linux.dump | 1 + 6 files changed, 203 insertions(+), 72 deletions(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index 85519c3388efd..b1feab6848468 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -11,7 +11,7 @@ // 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. namespace sycl { @@ -151,6 +151,107 @@ template struct KernelInfo { }; #endif //__SYCL_UNNAMED_LAMBDA__ +template struct KernelIdentity { + using type = KNT; +}; + +template constexpr unsigned getKernelNumParams() { +#if __has_builtin(__builtin_sycl_kernel_param_count) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + __builtin_sycl_kernel_param_count(KernelIdentity()); +#else + return KernelInfo::getNumParams(); +#endif +} + +template +std::vector getKernelParamDescs() { + std::vector Result; + int NumParams = getKernelNumParams(); + Result.reserve(NumParams); + for (int I = 0; I < NumParams; ++I) { +#if __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) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + kernel_param_desc_t ParamDesc; + ParamDesc.kind = + __builtin_sycl_kernel_param_kind(KernelIdentity(), I); + ParamDesc.info = Result.kind == kind_accessor + ? __builtin_sycl_kernel_param_access_target( + KernelIdentity(), I) + : __builtin_sycl_kernel_param_size( + KernelIdentity(), I); + ParamDesc.offset = + __builtin_sycl_kernel_param_offset(KernelIdentity(), I); + Result.push_back(ParamDesc); +#else + Result.push_back(KernelInfo::getParamDesc(I)); +#endif + } + return Result; +} + +template constexpr const char *getKernelName() { +#if __has_builtin(__builtin_sycl_kernel_name) && \ + !defined(__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() { +#if __has_builtin(__builtin_sycl_kernel_file_name) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + return __builtin_sycl_kernel_file_name(KernelIdentity()); +#else + return KernelInfo::getFileName(); +#endif +} + +template +constexpr const char *getKernelFunctionName() { +#if __has_builtin(__builtin_sycl_kernel_function_name) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + return __builtin_sycl_kernel_function_name(KernelIdentity()); +#else + return KernelInfo::getFunctionName(); +#endif +} + +template constexpr unsigned getKernelLineNumber() { +#if __has_builtin(__builtin_sycl_kernel_line_number) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + __builtin_sycl_kernel_line_number(KernelIdentity()); +#else + return KernelInfo::getLineNumber(); +#endif +} + +template constexpr unsigned getKernelColumnNumber() { +#if __has_builtin(__builtin_sycl_kernel_column_number) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + __builtin_sycl_kernel_column_number(KernelIdentity()); +#else + return KernelInfo::getColumnNumber(); +#endif +} + +template constexpr int64_t getKernelSize() { +#if __has_builtin(__builtin_sycl_kernel_size) && \ + !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) + __builtin_sycl_kernel_size(KernelIdentity()); +#else + return KernelInfo::getKernelSize(); +#endif +} } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 6f2e9f9fc19b7..87f8c6642820d 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -524,18 +524,12 @@ 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 - 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; + void throwOnLocalAccessorMisuse( + const std::vector &ParamDescs) const { + for (const auto &ParamDesc : ParamDescs) { + 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 +540,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 +568,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; } @@ -893,13 +891,15 @@ class __SYCL_EXPORT handler { /// 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; + void StoreLambda(KernelType KernelFunc, + const std::vector &ParamDescs = + detail::getKernelParamDescs()) { constexpr bool IsCallableWithKernelHandler = detail::KernelLambdaHasKernelHandlerArgT::value; @@ -908,13 +908,15 @@ 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. 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." @@ -932,9 +934,9 @@ class __SYCL_EXPORT handler { // TODO support ESIMD in no-integration-header case too. clearArgs(); extractArgsAndReqsFromLambda(reinterpret_cast(KernelPtr), - KI::getNumParams(), &KI::getParamDesc(0), - KI::isESIMD()); - MKernelName = KI::getName(); + ParamDescs, + 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 +1033,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 +1041,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< @@ -1298,7 +1299,13 @@ class __SYCL_EXPORT handler { void parallel_for_lambda_impl(range UserRange, PropertiesT Props, KernelType KernelFunc) { throwIfActionIsCreated(); - throwOnLocalAccessorMisuse(); + // 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; + std::vector ParamDescs = + detail::getKernelParamDescs(); + throwOnLocalAccessorMisuse(ParamDescs); if (!range_size_fits_in_size_t(UserRange)) throw sycl::exception(make_error_code(errc::runtime), "The total number of work-items in " @@ -1330,12 +1337,7 @@ 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; - - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); // Range rounding can be disabled by the user. // Range rounding is not done on the host device. @@ -1417,7 +1419,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(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; static_assert( @@ -1507,7 +1509,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(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)NumWorkGroups; @@ -1548,7 +1550,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(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)NumWorkGroups; @@ -1817,19 +1819,21 @@ class __SYCL_EXPORT handler { _KERNELFUNCPARAM(KernelFunc)) { (void)Props; throwIfActionIsCreated(); - throwOnLocalAccessorMisuse(); // 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; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + std::vector ParamDescs = + detail::getKernelParamDescs(); + throwOnLocalAccessorMisuse(ParamDescs); + verifyUsedKernelBundle(detail::getKernelName()); 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 // known constant. setNDRangeDescriptor(range<1>{1}); processProperties(Props); - StoreLambda(KernelFunc); + StoreLambda(KernelFunc, ParamDescs); setType(detail::CGType::Kernel); #endif } @@ -2118,7 +2122,7 @@ class __SYCL_EXPORT handler { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; using TransformedArgType = std::conditional_t< std::is_integral::value && Dims == 1, item, @@ -2259,7 +2263,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); (void)Kernel; kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ @@ -2294,7 +2298,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; (void)NumWorkItems; @@ -2333,7 +2337,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; (void)NumWorkItems; @@ -2372,7 +2376,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; @@ -2415,7 +2419,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; @@ -2455,7 +2459,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::KernelInfo::getName()); + verifyUsedKernelBundle(detail::getKernelName()); 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 fbab1e5ca9148..6b7e3aa3ee454 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) { @@ -2753,10 +2744,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) { @@ -2790,10 +2778,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) { @@ -2830,10 +2815,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) { @@ -2868,6 +2850,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 4d89d4826a722..8b247dfbee4fe 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -841,6 +841,41 @@ void handler::extractArgsAndReqs() { } } +void handler::extractArgsAndReqsFromLambda( + char *LambdaPtr, const std::vector &ParamDescs, + bool IsESIMD) { + const bool IsKernelCreatedFromSource = false; + size_t IndexShift = 0; + impl->MArgs.reserve(MaxNumAdditionalArgs * ParamDescs.size()); + + 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. + const access::target AccTarget = + static_cast(Size & AccessTargetMask); + if ((AccTarget == access::target::device || + AccTarget == access::target::constant_buffer) || + (AccTarget == access::target::image || + AccTarget == access::target::image_array)) { + detail::AccessorBaseHost *AccBase = + static_cast(Ptr); + Ptr = detail::getSyclObjImpl(*AccBase).get(); + } else if (AccTarget == access::target::local) { + detail::LocalAccessorBaseHost *LocalAccBase = + static_cast(Ptr); + Ptr = detail::getSyclObjImpl(*LocalAccBase).get(); + } + } + processArg(Ptr, Kind, Size, I, IndexShift, IsKernelCreatedFromSource, + IsESIMD); + } +} + +// TODO Unused, remove during ABI breaking window void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, size_t KernelArgsNum, const detail::kernel_param_desc_t *KernelArgs, bool IsESIMD) { diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 26822ac6e3bf2..c32a783f868da 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3513,6 +3513,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 From 8406f5f99300c619e22abc1a0c1f1d54bca3bcef Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Wed, 14 Aug 2024 06:41:05 -0700 Subject: [PATCH 02/18] Add missing include --- sycl/include/sycl/detail/kernel_desc.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index b1feab6848468..786e14d4986e4 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -10,6 +10,7 @@ // 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 #include // This header file must not include any standard C++ header files. From 1ca7e678e53c9d3ce0e680196367347c37ef8fe5 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Tue, 27 Aug 2024 07:45:31 -0700 Subject: [PATCH 03/18] Apply comments --- sycl/include/sycl/detail/kernel_desc.hpp | 57 +++++++++++++----------- sycl/include/sycl/handler.hpp | 5 ++- sycl/source/handler.cpp | 32 ++----------- 3 files changed, 38 insertions(+), 56 deletions(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index 786e14d4986e4..a0ec56cbd3446 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -15,6 +15,22 @@ #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 { @@ -152,14 +168,15 @@ template struct KernelInfo { }; #endif //__SYCL_UNNAMED_LAMBDA__ +// Built-ins use an object of this type to identify the kernel the information +// is requested for. template struct KernelIdentity { using type = KNT; }; template constexpr unsigned getKernelNumParams() { -#if __has_builtin(__builtin_sycl_kernel_param_count) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) - __builtin_sycl_kernel_param_count(KernelIdentity()); +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS + return __builtin_sycl_kernel_param_count(KernelIdentity()); #else return KernelInfo::getNumParams(); #endif @@ -171,15 +188,11 @@ std::vector getKernelParamDescs() { int NumParams = getKernelNumParams(); Result.reserve(NumParams); for (int I = 0; I < NumParams; ++I) { -#if __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) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS kernel_param_desc_t ParamDesc; ParamDesc.kind = __builtin_sycl_kernel_param_kind(KernelIdentity(), I); - ParamDesc.info = Result.kind == kind_accessor + ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor ? __builtin_sycl_kernel_param_access_target( KernelIdentity(), I) : __builtin_sycl_kernel_param_size( @@ -195,8 +208,7 @@ std::vector getKernelParamDescs() { } template constexpr const char *getKernelName() { -#if __has_builtin(__builtin_sycl_kernel_name) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS return __builtin_sycl_kernel_name(KernelIdentity()); #else return KernelInfo::getName(); @@ -209,8 +221,7 @@ template constexpr bool isKernelESIMD() { } template constexpr const char *getKernelFileName() { -#if __has_builtin(__builtin_sycl_kernel_file_name) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS return __builtin_sycl_kernel_file_name(KernelIdentity()); #else return KernelInfo::getFileName(); @@ -219,8 +230,7 @@ template constexpr const char *getKernelFileName() { template constexpr const char *getKernelFunctionName() { -#if __has_builtin(__builtin_sycl_kernel_function_name) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS return __builtin_sycl_kernel_function_name(KernelIdentity()); #else return KernelInfo::getFunctionName(); @@ -228,30 +238,25 @@ constexpr const char *getKernelFunctionName() { } template constexpr unsigned getKernelLineNumber() { -#if __has_builtin(__builtin_sycl_kernel_line_number) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) - __builtin_sycl_kernel_line_number(KernelIdentity()); +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS + return __builtin_sycl_kernel_line_number(KernelIdentity()); #else return KernelInfo::getLineNumber(); #endif } template constexpr unsigned getKernelColumnNumber() { -#if __has_builtin(__builtin_sycl_kernel_column_number) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) - __builtin_sycl_kernel_column_number(KernelIdentity()); +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS + return __builtin_sycl_kernel_column_number(KernelIdentity()); #else return KernelInfo::getColumnNumber(); #endif } template constexpr int64_t getKernelSize() { -#if __has_builtin(__builtin_sycl_kernel_size) && \ - !defined(__INTEL_SYCL_USE_INTEGRATION_HEADERS) - __builtin_sycl_kernel_size(KernelIdentity()); -#else + // TODO needs a builtin counterpart, but is currently only used for checking + // cases with external host compiler, which use integration headers. return KernelInfo::getKernelSize(); -#endif } } // namespace detail } // namespace _V1 diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 87f8c6642820d..d4cb2fa6d4578 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -914,6 +914,9 @@ class __SYCL_EXPORT handler { // 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) == detail::getKernelSize(), @@ -927,7 +930,7 @@ 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) { diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index ab0506695eb5e..c576d1e999f63 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -879,35 +879,9 @@ void handler::extractArgsAndReqsFromLambda( void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, size_t KernelArgsNum, const detail::kernel_param_desc_t *KernelArgs, bool IsESIMD) { - const bool IsKernelCreatedFromSource = false; - size_t IndexShift = 0; - impl->MArgs.reserve(MaxNumAdditionalArgs * KernelArgsNum); - - 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; - 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. - const access::target AccTarget = - static_cast(Size & AccessTargetMask); - if ((AccTarget == access::target::device || - AccTarget == access::target::constant_buffer) || - (AccTarget == access::target::image || - AccTarget == access::target::image_array)) { - detail::AccessorBaseHost *AccBase = - static_cast(Ptr); - Ptr = detail::getSyclObjImpl(*AccBase).get(); - } else if (AccTarget == access::target::local) { - detail::LocalAccessorBaseHost *LocalAccBase = - static_cast(Ptr); - Ptr = detail::getSyclObjImpl(*LocalAccBase).get(); - } - } - processArg(Ptr, Kind, Size, I, IndexShift, IsKernelCreatedFromSource, - IsESIMD); - } + std::vector ParamDescs( + KernelArgs, KernelArgs + KernelArgsNum); + extractArgsAndReqsFromLambda(LambdaPtr, ParamDescs, IsESIMD); } // Calling methods of kernel_impl requires knowledge of class layout. From dedc6383311933d3c6156a4f9ebfadf5c632d489 Mon Sep 17 00:00:00 2001 From: Semenov Date: Wed, 28 Aug 2024 05:11:26 -0700 Subject: [PATCH 04/18] Update Windows ABI dumps --- sycl/test/abi/sycl_symbols_windows.dump | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 910d966928f07..858769ce01ce9 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 @@ -4254,6 +4255,7 @@ ?supportsUSMMemset2D@handler@_V1@sycl@@AEAA_NXZ ?sycl_category@_V1@sycl@@YAAEBVerror_category@std@@XZ ?throwIfActionIsCreated@handler@_V1@sycl@@AEAAXXZ +?throwOnLocalAccessorMisuse@handler@_V1@sycl@@AEBAXAEBV?$vector@Ukernel_param_desc_t@detail@_V1@sycl@@V?$allocator@Ukernel_param_desc_t@detail@_V1@sycl@@@std@@@std@@@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 From 2330ade58661c233428c9ae3f9f972296301bee1 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Wed, 28 Aug 2024 05:50:56 -0700 Subject: [PATCH 05/18] Expand comment --- sycl/include/sycl/detail/kernel_desc.hpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index a0ec56cbd3446..cb8176281c44d 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -168,8 +168,10 @@ template struct KernelInfo { }; #endif //__SYCL_UNNAMED_LAMBDA__ -// Built-ins use an object of this type to identify the kernel the information -// is requested for. +// 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; }; From 327e88f3886beada01451c5f291ea7397193fe7e Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 2 Sep 2024 05:21:13 -0700 Subject: [PATCH 06/18] Remove C++ header from kernel_desc.hpp Who reads comments anyway? --- sycl/include/sycl/detail/kernel_desc.hpp | 23 ++++++++--------------- sycl/include/sycl/handler.hpp | 20 +++++++++++++++++++- 2 files changed, 27 insertions(+), 16 deletions(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index cb8176281c44d..103a574131353 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -12,7 +12,6 @@ // uses the macro defined in this header, but it doesn't explicitly include it. #include #include -#include // This header file must not include any standard C++ header files. #ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS @@ -185,28 +184,22 @@ template constexpr unsigned getKernelNumParams() { } template -std::vector getKernelParamDescs() { - std::vector Result; - int NumParams = getKernelNumParams(); - Result.reserve(NumParams); - for (int I = 0; I < NumParams; ++I) { +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(), I); + __builtin_sycl_kernel_param_kind(KernelIdentity(), Idx); ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor ? __builtin_sycl_kernel_param_access_target( - KernelIdentity(), I) + KernelIdentity(), Idx) : __builtin_sycl_kernel_param_size( - KernelIdentity(), I); - ParamDesc.offset = - __builtin_sycl_kernel_param_offset(KernelIdentity(), I); - Result.push_back(ParamDesc); + KernelIdentity(), Idx); + ParamDesc.offset = __builtin_sycl_kernel_param_offset( + KernelIdentity(), Idx); + return ParamDesc #else - Result.push_back(KernelInfo::getParamDesc(I)); + return KernelInfo::getParamDesc(Idx); #endif - } - return Result; } template constexpr const char *getKernelName() { diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index d4cb2fa6d4578..dcaf15cdc484b 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. @@ -1588,7 +1599,14 @@ class __SYCL_EXPORT handler { nullptr, ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { + +#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS + [[clang::sycl_kernel_entry_point(KernelName)]] +#else + __SYCL_KERNEL_ATTR__ +#endif + + void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(); #else From 1e6c0104537ca7e6ac176680416673e317e3a7f1 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 2 Sep 2024 05:28:44 -0700 Subject: [PATCH 07/18] Remove unrelated edit --- sycl/include/sycl/handler.hpp | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index dcaf15cdc484b..2c23a8400452f 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1600,13 +1600,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif -#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS - [[clang::sycl_kernel_entry_point(KernelName)]] -#else - __SYCL_KERNEL_ATTR__ -#endif - - void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { + __SYCL_KERNEL_ATTR__ void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(); #else From d0d8f93e1ad74aad9684872d2a4010b3f57d870c Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 2 Sep 2024 05:52:35 -0700 Subject: [PATCH 08/18] Appease clang-format --- sycl/include/sycl/detail/kernel_desc.hpp | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index 103a574131353..49106dc9fa6cd 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -186,17 +186,17 @@ template constexpr unsigned getKernelNumParams() { 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 + 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 From ae91935304fdd0950336ded00626a6b0a8f89b27 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Wed, 4 Sep 2024 04:34:41 -0700 Subject: [PATCH 09/18] Add missing semicolon --- sycl/include/sycl/detail/kernel_desc.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/detail/kernel_desc.hpp b/sycl/include/sycl/detail/kernel_desc.hpp index 49106dc9fa6cd..1049c4d78aadd 100644 --- a/sycl/include/sycl/detail/kernel_desc.hpp +++ b/sycl/include/sycl/detail/kernel_desc.hpp @@ -196,7 +196,7 @@ kernel_param_desc_t getKernelParamDesc(int Idx) { KernelIdentity(), Idx); ParamDesc.offset = __builtin_sycl_kernel_param_offset(KernelIdentity(), Idx); - return ParamDesc + return ParamDesc; #else return KernelInfo::getParamDesc(Idx); #endif From 9d21cb2ca28b23c17fd348e2d9e348367dc8d5d5 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Wed, 4 Sep 2024 06:46:35 -0700 Subject: [PATCH 10/18] Use the new entry point attribute with the built-ins --- sycl/include/sycl/handler.hpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 2c23a8400452f..a49adedcb75f6 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1584,10 +1584,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. From b36d3bef4e5c691ced9bede856fd50e6fe4114ff Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Wed, 4 Sep 2024 11:30:02 -0700 Subject: [PATCH 11/18] Pass ParamDescs to StoreLambda --- sycl/include/sycl/handler.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a49adedcb75f6..a8d74ad755aef 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1380,7 +1380,7 @@ class __SYCL_EXPORT handler { detail::checkValueRange(UserRange); setNDRangeDescriptor(RoundedRange); StoreLambda( - std::move(Wrapper)); + std::move(Wrapper), ParamDescs); setType(detail::CGType::Kernel); setNDRangeUsed(false); #endif From a21238fb65877518cce788304707987b1e9d015e Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Thu, 5 Sep 2024 07:31:56 -0700 Subject: [PATCH 12/18] Make entry points static functions --- sycl/include/sycl/handler.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a8d74ad755aef..b4113cfa0ae36 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1604,7 +1604,7 @@ class __SYCL_EXPORT handler { 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 @@ -1622,7 +1622,7 @@ 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), kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(KH); @@ -1641,7 +1641,7 @@ 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 @@ -1658,7 +1658,7 @@ 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), kernel_handler KH) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(detail::Builder::getElement(detail::declptr()), KH); @@ -1677,7 +1677,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())); @@ -1695,7 +1695,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__ From 01d651ad457b3496ada4fe25f1ef22f39d8c6ac2 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Thu, 5 Sep 2024 07:59:35 -0700 Subject: [PATCH 13/18] Apply clang-format --- sycl/include/sycl/handler.hpp | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index b4113cfa0ae36..c7ff0fbab5a05 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1604,7 +1604,8 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ static void kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { + __SYCL_KERNEL_ATTR__ static void + kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) { #ifdef __SYCL_DEVICE_ONLY__ KernelFunc(); #else @@ -1622,8 +1623,8 @@ class __SYCL_EXPORT handler { nullptr, ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ static 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 @@ -1641,7 +1642,8 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::name..., ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ static 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 @@ -1658,8 +1660,8 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::PropertyMetaInfo::name..., ext::oneapi::experimental::detail::PropertyMetaInfo::value...)]] #endif - __SYCL_KERNEL_ATTR__ static 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 From e0604b48fa95b855a663c8bbf8171ee58e3fd735 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Thu, 5 Sep 2024 11:26:40 -0700 Subject: [PATCH 14/18] Correct the StoreLambda call that reuses the descriptors --- sycl/include/sycl/handler.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index c7ff0fbab5a05..e2b364ecce390 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1380,7 +1380,7 @@ class __SYCL_EXPORT handler { detail::checkValueRange(UserRange); setNDRangeDescriptor(RoundedRange); StoreLambda( - std::move(Wrapper), ParamDescs); + std::move(Wrapper)); setType(detail::CGType::Kernel); setNDRangeUsed(false); #endif @@ -1401,7 +1401,7 @@ class __SYCL_EXPORT handler { detail::checkValueRange(UserRange); setNDRangeDescriptor(std::move(UserRange)); StoreLambda( - std::move(KernelFunc)); + std::move(KernelFunc), ParamDescs); setType(detail::CGType::Kernel); setNDRangeUsed(false); #endif From 969140447b16d2f69921b2593a5147c110186e13 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Fri, 6 Sep 2024 09:40:25 -0700 Subject: [PATCH 15/18] Address comment --- sycl/include/sycl/handler.hpp | 37 +++++++++++++++++------------------ 1 file changed, 18 insertions(+), 19 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index e2b364ecce390..cc06a6e00df6b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -535,9 +535,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. - void throwOnLocalAccessorMisuse( - const std::vector &ParamDescs) const { - for (const auto &ParamDesc : ParamDescs) { + template + void throwOnLocalAccessorMisuse() const { + using NameT = + typename detail::get_kernel_name_t::name; + 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(ParamDesc.info & AccessTargetMask); @@ -908,9 +912,7 @@ class __SYCL_EXPORT handler { /// \param ParamDescs is the vector of kernel parameter descriptors. template - void StoreLambda(KernelType KernelFunc, - const std::vector &ParamDescs = - detail::getKernelParamDescs()) { + void StoreLambda(KernelType KernelFunc) { constexpr bool IsCallableWithKernelHandler = detail::KernelLambdaHasKernelHandlerArgT::value; @@ -948,7 +950,7 @@ class __SYCL_EXPORT handler { // TODO support ESIMD in no-integration-header case too. clearArgs(); extractArgsAndReqsFromLambda(reinterpret_cast(KernelPtr), - ParamDescs, + detail::getKernelParamDescs(), detail::isKernelESIMD()); MKernelName = detail::getKernelName(); } else { @@ -1313,13 +1315,7 @@ class __SYCL_EXPORT handler { void parallel_for_lambda_impl(range UserRange, PropertiesT Props, KernelType KernelFunc) { throwIfActionIsCreated(); - // 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; - std::vector ParamDescs = - detail::getKernelParamDescs(); - throwOnLocalAccessorMisuse(ParamDescs); + throwOnLocalAccessorMisuse(); if (!range_size_fits_in_size_t(UserRange)) throw sycl::exception(make_error_code(errc::runtime), "The total number of work-items in " @@ -1351,6 +1347,10 @@ 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; verifyUsedKernelBundle(detail::getKernelName()); // Range rounding can be disabled by the user. @@ -1401,7 +1401,7 @@ class __SYCL_EXPORT handler { detail::checkValueRange(UserRange); setNDRangeDescriptor(std::move(UserRange)); StoreLambda( - std::move(KernelFunc), ParamDescs); + std::move(KernelFunc)); setType(detail::CGType::Kernel); setNDRangeUsed(false); #endif @@ -1840,13 +1840,12 @@ class __SYCL_EXPORT handler { _KERNELFUNCPARAM(KernelFunc)) { (void)Props; throwIfActionIsCreated(); + throwOnLocalAccessorMisuse(); // 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; - std::vector ParamDescs = - detail::getKernelParamDescs(); - throwOnLocalAccessorMisuse(ParamDescs); + verifyUsedKernelBundle(detail::getKernelName()); kernel_single_task_wrapper(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ @@ -1854,7 +1853,7 @@ class __SYCL_EXPORT handler { // known constant. setNDRangeDescriptor(range<1>{1}); processProperties(Props); - StoreLambda(KernelFunc, ParamDescs); + StoreLambda(KernelFunc); setType(detail::CGType::Kernel); #endif } From 54858dfb91f2e39f02fc53f8d314106ad549f807 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 9 Sep 2024 04:59:44 -0700 Subject: [PATCH 16/18] Update Windows ABI dump --- sycl/test/abi/sycl_symbols_windows.dump | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index a39145c63bf1b..c28723fc29245 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4256,7 +4256,6 @@ ?supportsUSMMemset2D@handler@_V1@sycl@@AEAA_NXZ ?sycl_category@_V1@sycl@@YAAEBVerror_category@std@@XZ ?throwIfActionIsCreated@handler@_V1@sycl@@AEAAXXZ -?throwOnLocalAccessorMisuse@handler@_V1@sycl@@AEBAXAEBV?$vector@Ukernel_param_desc_t@detail@_V1@sycl@@V?$allocator@Ukernel_param_desc_t@detail@_V1@sycl@@@std@@@std@@@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 From 4a89abad1240280427d146c593d9a88bfd3fe75d Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 9 Sep 2024 05:49:22 -0700 Subject: [PATCH 17/18] Apply comment --- sycl/include/sycl/handler.hpp | 29 +++++++++++++------------ sycl/test/abi/sycl_symbols_windows.dump | 1 - 2 files changed, 15 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index cc06a6e00df6b..c0c0e6bfe4c4c 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -898,8 +898,9 @@ 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); @@ -1351,7 +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::getKernelName()); + verifyUsedKernelBundle(); // Range rounding can be disabled by the user. // Range rounding is not done on the host device. @@ -1433,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::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; static_assert( @@ -1523,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::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)NumWorkGroups; @@ -1564,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::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)NumWorkGroups; @@ -1846,7 +1847,7 @@ class __SYCL_EXPORT handler { using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + 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 @@ -2142,7 +2143,7 @@ class __SYCL_EXPORT handler { throwIfActionIsCreated(); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; using TransformedArgType = std::conditional_t< std::is_integral::value && Dims == 1, item, @@ -2283,7 +2284,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); (void)Kernel; kernel_single_task(KernelFunc); #ifndef __SYCL_DEVICE_ONLY__ @@ -2318,7 +2319,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; (void)NumWorkItems; @@ -2357,7 +2358,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; (void)NumWorkItems; @@ -2396,7 +2397,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; @@ -2439,7 +2440,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; @@ -2479,7 +2480,7 @@ class __SYCL_EXPORT handler { setHandlerKernelBundle(Kernel); using NameT = typename detail::get_kernel_name_t::name; - verifyUsedKernelBundle(detail::getKernelName()); + verifyUsedKernelBundle(); using LambdaArgType = sycl::detail::lambda_arg_type>; (void)Kernel; diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index c28723fc29245..e2c3643c557be 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4270,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 From 74283ca9955e789b34f5b0b8eba8f5f64075cfa2 Mon Sep 17 00:00:00 2001 From: Sergey Semenov Date: Mon, 9 Sep 2024 05:50:03 -0700 Subject: [PATCH 18/18] Apply clang-format --- sycl/include/sycl/handler.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index c0c0e6bfe4c4c..6181a41e6ef8c 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -898,9 +898,9 @@ class __SYCL_EXPORT handler { /// /// \param KernelName is the name of the SYCL kernel to check that the used /// kernel bundle contains. - template - void verifyUsedKernelBundle() { - verifyUsedKernelBundleInternal(detail::string_view{detail::getKernelName()}); + template void verifyUsedKernelBundle() { + verifyUsedKernelBundleInternal( + detail::string_view{detail::getKernelName()}); } void verifyUsedKernelBundleInternal(detail::string_view KernelName);