diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp index fe5e52f0d2ab9..74fa2b7be6a98 100644 --- a/sycl/include/CL/sycl/detail/cg_types.hpp +++ b/sycl/include/CL/sycl/detail/cg_types.hpp @@ -242,7 +242,7 @@ class HostTask { }; // Class which stores specific lambda object. -template +template class HostKernel : public HostKernelBase { using IDBuilder = sycl::detail::Builder; KernelType MKernel; @@ -290,9 +290,6 @@ class HostKernel : public HostKernelBase { template typename detail::enable_if_t>::value> runOnHost(const NDRDescT &NDRDesc) { - using KI = detail::KernelInfo; - constexpr bool StoreLocation = KI::callsAnyThisFreeFunction(); - sycl::range Range(InitializedVal::template get<0>()); sycl::id Offset; sycl::range Stride( @@ -323,9 +320,6 @@ class HostKernel : public HostKernelBase { typename detail::enable_if_t< std::is_same>::value> runOnHost(const NDRDescT &NDRDesc) { - using KI = detail::KernelInfo; - constexpr bool StoreLocation = KI::callsAnyThisFreeFunction(); - sycl::id ID; sycl::range Range(InitializedVal::template get<0>()); for (int I = 0; I < Dims; ++I) @@ -348,9 +342,6 @@ class HostKernel : public HostKernelBase { typename detail::enable_if_t< std::is_same>::value> runOnHost(const NDRDescT &NDRDesc) { - using KI = detail::KernelInfo; - constexpr bool StoreLocation = KI::callsAnyThisFreeFunction(); - sycl::range Range(InitializedVal::template get<0>()); sycl::id Offset; sycl::range Stride( @@ -380,9 +371,6 @@ class HostKernel : public HostKernelBase { template typename detail::enable_if_t>::value> runOnHost(const NDRDescT &NDRDesc) { - using KI = detail::KernelInfo; - constexpr bool StoreLocation = KI::callsAnyThisFreeFunction(); - sycl::range GroupSize(InitializedVal::template get<0>()); for (int I = 0; I < Dims; ++I) { if (NDRDesc.LocalSize[I] == 0 || diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index c3ddbc1930258..1f0170ef8a455 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -576,14 +576,14 @@ class __SYCL_EXPORT handler { // For 'id, item w/wo offset, nd_item' kernel arguments template + bool StoreLocation> KernelType *ResetHostKernelHelper(const KernelType &KernelFunc) { NormalizedKernelType NormalizedKernel(KernelFunc); auto NormalizedKernelFunc = std::function &)>(NormalizedKernel); auto HostKernelPtr = new detail::HostKernel, Dims, KernelName>( + sycl::nd_item, Dims, StoreLocation>( NormalizedKernelFunc); MHostKernel.reset(HostKernelPtr); return &HostKernelPtr->MKernel.template target() @@ -591,7 +591,7 @@ class __SYCL_EXPORT handler { } // For 'sycl::id' kernel argument - template + template typename std::enable_if>::value, KernelType *>::type ResetHostKernel(const KernelType &KernelFunc) { @@ -604,11 +604,11 @@ class __SYCL_EXPORT handler { } }; return ResetHostKernelHelper(KernelFunc); + StoreLocation>(KernelFunc); } // For 'sycl::nd_item' kernel argument - template + template typename std::enable_if>::value, KernelType *>::type ResetHostKernel(const KernelType &KernelFunc) { @@ -621,11 +621,11 @@ class __SYCL_EXPORT handler { } }; return ResetHostKernelHelper(KernelFunc); + StoreLocation>(KernelFunc); } // For 'sycl::item' kernel argument - template + template typename std::enable_if>::value, KernelType *>::type ResetHostKernel(const KernelType &KernelFunc) { @@ -640,11 +640,11 @@ class __SYCL_EXPORT handler { } }; return ResetHostKernelHelper(KernelFunc); + StoreLocation>(KernelFunc); } // For 'sycl::item' kernel argument - template + template typename std::enable_if>::value, KernelType *>::type ResetHostKernel(const KernelType &KernelFunc) { @@ -659,7 +659,7 @@ class __SYCL_EXPORT handler { } }; return ResetHostKernelHelper(KernelFunc); + StoreLocation>(KernelFunc); } /* 'wrapper'-based approach using 'NormalizedKernelType' struct is @@ -669,13 +669,14 @@ class __SYCL_EXPORT handler { * not supported in ESIMD. */ // For 'void' and 'sycl::group' kernel argument - template + template typename std::enable_if::value || std::is_same>::value, KernelType *>::type ResetHostKernel(const KernelType &KernelFunc) { MHostKernel.reset( - new detail::HostKernel(KernelFunc)); + new detail::HostKernel( + KernelFunc)); return (KernelType *)(MHostKernel->getPtr()); } @@ -697,6 +698,9 @@ class __SYCL_EXPORT handler { template void StoreLambda(KernelType KernelFunc) { + using KI = detail::KernelInfo; + constexpr bool StoreLocation = KI::callsAnyThisFreeFunction(); + constexpr bool IsCallableWithKernelHandler = detail::KernelLambdaHasKernelHandlerArgT::value; @@ -707,7 +711,7 @@ class __SYCL_EXPORT handler { PI_INVALID_OPERATION); } KernelType *KernelPtr = - ResetHostKernel( + ResetHostKernel( KernelFunc); using KI = sycl::detail::KernelInfo; @@ -1481,7 +1485,7 @@ class __SYCL_EXPORT handler { MArgs = std::move(MAssociatedAccesors); MHostKernel.reset( - new detail::HostKernel(std::move(Func))); + new detail::HostKernel(std::move(Func))); setType(detail::CG::RunOnHostIntel); } diff --git a/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp b/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp index 395b7e2846f34..e1e87d8464f57 100644 --- a/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp +++ b/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp @@ -26,8 +26,7 @@ class MockHandler : public sycl::handler { typename KernelName> void setHostKernel(KernelType Kernel) { static_cast(this)->MHostKernel.reset( - new sycl::detail::HostKernel( - Kernel)); + new sycl::detail::HostKernel(Kernel)); } template void setNDRangeDesc(sycl::nd_range Range) {