diff options
author | 2021-02-23 18:10:50 +0300 | |
---|---|---|
committer | 2021-02-23 18:10:50 +0300 | |
commit | 6b28a2efa1afa233b0edeb21be04c048820a15c1 (patch) | |
tree | 8a56ad25ff96f096587442d26dff54984aa46d48 | |
parent | Avoid divergence of work items in the same SIMD before calling collectives (#... (diff) | |
download | llvm-project-6b28a2efa1afa233b0edeb21be04c048820a15c1.tar.gz llvm-project-6b28a2efa1afa233b0edeb21be04c048820a15c1.tar.bz2 llvm-project-6b28a2efa1afa233b0edeb21be04c048820a15c1.zip |
Fix for kernel names check (#133)
-rw-r--r-- | include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h | 17 | ||||
-rw-r--r-- | test/general/lambda_naming.pass.cpp | 8 |
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"); |