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..dc9f8fcf1146e 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -530,7 +530,7 @@ class __SYCL_EXPORT handler { // values of arguments for the kernel. assert(MKernel && "MKernel is not initialized"); constexpr std::string_view LambdaName = - detail::getKernelName(); + detail::CompileTimeKernelInfo.Name; detail::ABINeutralKernelNameStrT KernelName = getKernelName(); return KernelName == LambdaName; } @@ -823,20 +823,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 +852,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,9 +1234,15 @@ 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(); + 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 " @@ -1275,11 +1275,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 +1296,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 +1322,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( @@ -1410,12 +1402,12 @@ 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 detail::string_view Name{detail::getKernelName()}; - verifyUsedKernelBundleInternal(Name); + verifyUsedKernelBundleInternal(Info.Name); setType(detail::CGType::Kernel); detail::checkValueRange(params...); @@ -1427,7 +1419,7 @@ class __SYCL_EXPORT handler { } StoreLambda(std::move(KernelFunc)); - processProperties(), PropertiesT>(Props); + processProperties(Props); #endif } @@ -1453,15 +1445,15 @@ 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 detail::string_view Name{detail::getKernelName()}; - verifyUsedKernelBundleInternal(Name); + verifyUsedKernelBundleInternal(Info.Name); setType(detail::CGType::Kernel); detail::checkValueRange(params...); @@ -1479,7 +1471,7 @@ class __SYCL_EXPORT handler { } else { StoreLambda(std::move(KernelFunc)); } - processProperties(), PropertiesT>(Props); + processProperties(Props); #endif } #endif // __INTEL_PREVIEW_BREAKING_CHANGES @@ -1931,8 +1923,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}); @@ -3533,10 +3525,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); @@ -3556,13 +3548,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; - throwOnKernelParameterMisuseHelper(detail::getKernelNumParams(), - &detail::getKernelParamDesc); + +#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 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/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 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 {