Fix for kernel names check (#133)

2 files changed, 21 insertions, 4 deletions

diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
index ee275c252881..81595360d37c 100644
--- a/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
+++ b/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
@@ -112,11 +112,24 @@ make_wrapped_policy(_Policy&& __policy)
 namespace __internal
 {
 
+// extract the deepest kernel name when we have a policy wrapper that might hide the default name
+template <typename _CustomName>
+struct _HasDefaultName
+{
+    static constexpr bool value = std::is_same<_CustomName, oneapi::dpl::execution::DefaultKernelName>::value;
+};
+
+template <template <typename...> class _ExternalName, typename _InternalName>
+struct _HasDefaultName<_ExternalName<_InternalName>>
+{
+    static constexpr bool value = _HasDefaultName<_InternalName>::value;
+};
+
 template <template <typename...> class _BaseName, typename _CustomName, typename... _Args>
 using _KernelName_t =
 #if __SYCL_UNNAMED_LAMBDA__
-    typename std::conditional<std::is_same<_CustomName, oneapi::dpl::execution::DefaultKernelName>::value,
-                              _BaseName<_CustomName, _Args...>, _BaseName<_CustomName>>::type;
+    typename std::conditional<_HasDefaultName<_CustomName>::value, _BaseName<_CustomName, _Args...>,
+                              _BaseName<_CustomName>>::type;
 #else
     _BaseName<_CustomName>;
 #endif
diff --git a/test/general/lambda_naming.pass.cpp b/test/general/lambda_naming.pass.cpp
index b3040432bc56..98802322d2ac 100644
--- a/test/general/lambda_naming.pass.cpp
+++ b/test/general/lambda_naming.pass.cpp
@@ -32,8 +32,8 @@ using namespace TestUtils;
 int main() {
 #if _ONEDPL_BACKEND_SYCL
     const int n = 1000;
-    sycl::buffer<int, 1> buf{ sycl::range<1>(n) };
-    sycl::buffer<int, 1> out_buf{ sycl::range<1>(n) };
+    sycl::buffer<int> buf{ sycl::range<1>(n) };
+    sycl::buffer<int> out_buf{ sycl::range<1>(n) };
     auto buf_begin = oneapi::dpl::begin(buf);
     auto buf_end = buf_begin + n;
 
@@ -52,6 +52,10 @@ int main() {
     ::std::for_each(policy, buf_begin, buf_end, [](int& x) { x += 41; });
 
 #if !_ONEDPL_FPGA_DEVICE
+    sycl::buffer<float> out_buf_2{ sycl::range<1>(n) };
+    auto buf_out_begin_2 = oneapi::dpl::begin(out_buf_2);
+    ::std::copy(policy, buf_begin, buf_end, buf_out_begin_2);
+    ::std::copy(policy, buf_out_begin_2, buf_out_begin_2 + n, buf_begin);
     ::std::inplace_merge(policy, buf_begin, buf_begin + n / 2, buf_end);
     auto red_val = ::std::reduce(policy, buf_begin, buf_end, 1);
     EXPECT_TRUE(red_val == 42001, "wrong return value from reduce");