diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a76d6002d9d87..8765198f0bc0f 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2324,49 +2324,69 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - single_task(PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t::value> single_task(PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { single_task_lambda_impl(Props, KernelFunc); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<1> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t::value> parallel_for(range<1> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<2> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t::value> parallel_for(range<2> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<3> NumWorkItems, PropertiesT Props, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t::value> parallel_for(range<3> NumWorkItems, + PropertiesT Props, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } template - std::enable_if_t< - ext::oneapi::experimental::is_property_list::value> - parallel_for(nd_range Range, PropertiesT Properties, - _KERNELFUNCPARAM(KernelFunc)) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t::value> parallel_for(nd_range Range, + PropertiesT Properties, + _KERNELFUNCPARAM(KernelFunc)) { parallel_for_impl(Range, Properties, std::move(KernelFunc)); } @@ -2374,11 +2394,16 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<1> Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(range<1> Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2389,11 +2414,16 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<2> Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(range<2> Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2404,11 +2434,16 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(range<3> Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(range<3> Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2443,11 +2478,16 @@ class __SYCL_EXPORT handler { template - std::enable_if_t< - (sizeof...(RestT) > 1) && - detail::AreAllButLastReductions::value && - ext::oneapi::experimental::is_property_list::value> - parallel_for(nd_range Range, PropertiesT Properties, RestT &&...Rest) { + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t<(sizeof...(RestT) > 1) && + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list< + PropertiesT>::value> parallel_for(nd_range Range, + PropertiesT Properties, + RestT &&...Rest) { #ifndef __SYCL_DEVICE_ONLY__ throwIfGraphAssociated(); @@ -2469,6 +2509,10 @@ class __SYCL_EXPORT handler { template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") void parallel_for_work_group(range NumWorkGroups, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { parallel_for_work_group_lambda_impl + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 9c83bc3cac5ba..82a5d17f1cac2 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -2599,11 +2599,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - single_task( - PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), - const detail::code_location &CodeLoc = detail::code_location::current()) { + ext::oneapi::experimental::is_property_list::value, + event> single_task(PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), + const detail::code_location &CodeLoc = + detail::code_location::current()) { static_assert( (detail::check_fn_signature, void()>::value || @@ -2641,11 +2645,16 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - single_task( - event DepEvent, PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), - const detail::code_location &CodeLoc = detail::code_location::current()) { + ext::oneapi::experimental::is_property_list::value, + event> single_task(event DepEvent, PropertiesT Properties, + _KERNELFUNCPARAM(KernelFunc), + const detail::code_location &CodeLoc = + detail::code_location::current()) { static_assert( (detail::check_fn_signature, void()>::value || @@ -2687,12 +2696,16 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// \param CodeLoc contains the code location of user code template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") std::enable_if_t< - ext::oneapi::experimental::is_property_list::value, event> - single_task( - const std::vector &DepEvents, PropertiesT Properties, - _KERNELFUNCPARAM(KernelFunc), - const detail::code_location &CodeLoc = detail::code_location::current()) { + ext::oneapi::experimental::is_property_list::value, + event> single_task(const std::vector &DepEvents, + PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc), + const detail::code_location &CodeLoc = + detail::code_location::current()) { static_assert( (detail::check_fn_signature, void()>::value || @@ -2726,6 +2739,26 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { CodeLoc); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<1> Range, PropertiesT Properties, + RestT &&...Rest) { + return parallel_for_impl(Range, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2737,6 +2770,26 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<2> Range, PropertiesT Properties, + RestT &&...Rest) { + return parallel_for_impl(Range, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2748,6 +2801,26 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<3> Range, PropertiesT Properties, + RestT &&...Rest) { + return parallel_for_impl(Range, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2759,6 +2832,27 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<1> Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvent, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2771,6 +2865,27 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvent, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<2> Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvent, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2783,6 +2898,27 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvent, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<3> Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvent, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2795,6 +2931,28 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvent, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<1> Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvents, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2809,6 +2967,28 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvents, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<2> Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvents, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2823,6 +3003,28 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { return parallel_for_impl(Range, DepEvents, Rest...); } + /// parallel_for version with a kernel represented as a lambda + range that + /// specifies global size only. + /// + /// \param Range specifies the global work space of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(range<3> Range, const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + return parallel_for_impl(Range, DepEvents, Properties, Rest...); + } + /// parallel_for version with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -2938,11 +3140,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") std::enable_if_t< detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value, - event> - parallel_for(nd_range Range, PropertiesT Properties, RestT &&...Rest) { + event> parallel_for(nd_range Range, PropertiesT Properties, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2962,8 +3168,42 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { typename... RestT> std::enable_if_t::value, event> parallel_for(nd_range Range, RestT &&...Rest) { - return parallel_for( - Range, ext::oneapi::experimental::empty_properties_t{}, Rest...); + constexpr detail::code_location CodeLoc = getCodeLocation(); + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.template parallel_for(Range, Rest...); + }, + TlsCodeLocCapture.query()); + } + + /// parallel_for version with a kernel represented as a lambda + nd_range that + /// specifies global, local sizes and offset. + /// + /// \param Range specifies the global and local work spaces of the kernel + /// \param DepEvent is an event that specifies the kernel dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(nd_range Range, event DepEvent, + PropertiesT Properties, RestT &&...Rest) { + constexpr detail::code_location CodeLoc = getCodeLocation(); + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.template parallel_for(Range, Properties, Rest...); + }, + TlsCodeLocCapture.query()); } /// parallel_for version with a kernel represented as a lambda + nd_range that @@ -2975,7 +3215,8 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template - event parallel_for(nd_range Range, event DepEvent, RestT &&...Rest) { + std::enable_if_t::value, event> + parallel_for(nd_range Range, event DepEvent, RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( @@ -2986,6 +3227,37 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { TlsCodeLocCapture.query()); } + /// parallel_for version with a kernel represented as a lambda + nd_range that + /// specifies global, local sizes and offset. + /// + /// \param Range specifies the global and local work spaces of the kernel + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies + /// \param Properties is the kernel properties. + /// \param Rest acts as-if: "ReductionTypes&&... Reductions, + /// const KernelType &KernelFunc". + template + __SYCL_DEPRECATED("To specify properties, use a launch configuration object " + "of type launch_config or a kernel functor with a " + "get(sycl::ext::oneapi::experimental::properties_tag) " + "member function instead.") + std::enable_if_t< + detail::AreAllButLastReductions::value && + ext::oneapi::experimental::is_property_list::value, + event> parallel_for(nd_range Range, + const std::vector &DepEvents, + PropertiesT Properties, RestT &&...Rest) { + constexpr detail::code_location CodeLoc = getCodeLocation(); + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.template parallel_for(Range, Properties, Rest...); + }, + TlsCodeLocCapture.query()); + } + /// parallel_for version with a kernel represented as a lambda + nd_range that /// specifies global, local sizes and offset. /// @@ -2996,8 +3268,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { /// const KernelType &KernelFunc". template - event parallel_for(nd_range Range, const std::vector &DepEvents, - RestT &&...Rest) { + std::enable_if_t::value, event> + parallel_for(nd_range Range, const std::vector &DepEvents, + RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); return submit( diff --git a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp index 9ce59931405d6..4ea4e7cf125b7 100644 --- a/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp +++ b/sycl/test-e2e/VirtualFunctions/multiple-translation-units/Inputs/call.cpp @@ -1,17 +1,28 @@ #include "declarations.hpp" +template struct KernelFunctor { + T1 mDeviceStorage; + T2 mDataAcc; + KernelFunctor(T1 &DeviceStorage, T2 &DataAcc) + : mDeviceStorage(DeviceStorage), mDataAcc(DataAcc) {} + + void operator()() const { + auto *Ptr = mDeviceStorage->template getAs(); + Ptr->increment( + mDataAcc.template get_multi_ptr().get()); + } + auto get(oneapi::properties_tag) const { + return oneapi::properties{oneapi::assume_indirect_calls}; + } +}; + int call(sycl::queue Q, storage_t *DeviceStorage, int Init) { int Data = Init; { sycl::buffer DataStorage(&Data, sycl::range{1}); - constexpr oneapi::properties props{oneapi::assume_indirect_calls}; Q.submit([&](sycl::handler &CGH) { sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only); - CGH.single_task(props, [=]() { - auto *Ptr = DeviceStorage->getAs(); - Ptr->increment( - DataAcc.get_multi_ptr().get()); - }); + CGH.single_task(KernelFunctor(DeviceStorage, DataAcc)); }); } diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp index 055b25b920b8b..3d1c528744afd 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp @@ -1,5 +1,8 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp index a179c134749e9..ad81d1db1fe0b 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_sub_group_size.cpp @@ -1,5 +1,8 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp index 932b92fab9009..63280fcc638f3 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size.cpp @@ -1,5 +1,8 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp index a844b484b8b51..a0bae31ad8004 100644 --- a/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp +++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_work_group_size_hint.cpp @@ -1,5 +1,8 @@ -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// TODO: Currently using the -Wno-deprecated-declarations flag due to issue +// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is +// resolved. +// RUN: %clangxx -fsycl-device-only -S -Wno-deprecated-declarations -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations -Xclang -verify %s // expected-no-diagnostics #include diff --git a/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp new file mode 100644 index 0000000000000..033c062d6b538 --- /dev/null +++ b/sycl/test/warnings/deprecated_single_task_parallel_for_with_props.cpp @@ -0,0 +1,65 @@ +// Ignore unexpected warnings because for some reason the warnings are emitted +// twice, e.g. once for `single_task`, then for `single_task>>`. +// RUN: %clangxx -fsycl -sycl-std=2020 -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning -Xclang -verify-ignore-unexpected=note %s -fsyntax-only -Wall -Wextra +#include + +using namespace sycl; +int main() { + queue Q; + event Ev; + range<1> R1{1}; + range<2> R2(1, 1); + range<3> R3(1, 1, 1); + nd_range<1> NDR1{R1, R1}; + constexpr auto Props = sycl::ext::oneapi::experimental::properties{}; + + // expected-warning@+1{{'single_task' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.single_task(Props, []() {}); + // expected-warning@+1{{'single_task' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.single_task(Ev, Props, []() {}); + // expected-warning@+1{{'single_task' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.single_task({Ev}, Props, []() {}); + + // expected-warning@+1{{'parallel_for' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.parallel_for(NDR1, Props, [](nd_item<1>) {}); + + // expected-warning@+2{{'single_task' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.single_task(Props, []() {}); + }); + + // expected-warning@+2{{'parallel_for_work_group' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R1, Props, + [](sycl::group<1>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R2, Props, + [](sycl::group<2>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R3, Props, + [](sycl::group<3>) {}); + }); + + // expected-warning@+2{{'parallel_for_work_group' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R1, R1, Props, + [](sycl::group<1>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R2, R2, Props, + [](sycl::group<2>) {}); + }); + // expected-warning@+2{{'parallel_for_work_group' is deprecated: To specify properties, use a launch configuration object of type launch_config or a kernel functor with a get(sycl::ext::oneapi::experimental::properties_tag) member function instead.}} + Q.submit([&](handler &CGH) { + CGH.parallel_for_work_group(R3, R3, Props, + [](sycl::group<3>) {}); + }); + return 0; +}