Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Document issue with Kokkos::sort and comparators #504

Merged
merged 4 commits into from
Apr 15, 2024
Merged

Document issue with Kokkos::sort and comparators #504

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
b036e87
Document issue with Kokkos::sort and comparators
masterleinad Feb 28, 2024
25284b5
Provide example with error
masterleinad Apr 15, 2024
4e016a2
Multiple -> Several
masterleinad Apr 15, 2024
5f4bda3
Demonstrate sycl::is_device_copyable
masterleinad Apr 15, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
52 changes: 52 additions & 0 deletions docs/source/known-issues.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,58 @@ HIP

gcc 7, 9, and later do not have this issue.

SYCL
====

- Several of the Kokkos algorithm functions use third-party libraries like oneDPL.
When using these, Kokkos doesn't control the kernel launch and thus the user has to make sure that all arguments
that are forwarded to the TPL satisfy the sycl::is_device_copyable trait to avoid compiler errors. This holds true in particular
for comparators used with Kokkos::sort in Kokkos versions prior to 4.3. The best advice to give is to make sure the respective
parameters are trivially-copyable. If this isn't possible, sycl::is_device_copyable should be specialized and users should make
sure to use raw pointers instead of Kokkos::Views.
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you add a code example and include the compile error so that it can be easily found by users?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more. 

Kokkos::sort(exec, values, KOKKOS_LAMBDA(int i, int j) { return keys(i) < keys(j); });  // error blah with SYCL

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

25284b5


.. code-block:: cpp

MyComparator my_comparator;
Kokkos::sort(exec, values, my_comparator);

would give errors similar to

.. code-block:: console

/usr/bin/compiler/../../include/sycl/types.hpp:2572:17: error: static assertion failed due to requirement 'is_device_copyable_v<(lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1816:20)> || detail::IsDeprecatedDeviceCopyable<(lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1816:20), void>::value': The specified type is not device copyable
2572 | static_assert(is_device_copyable_v<FieldT> ||
| ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
2573 | detail::IsDeprecatedDeviceCopyable<FieldT>::value,
| ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/usr/bin/compiler/../../include/sycl/types.hpp:2605:7: note: in instantiation of template class 'sycl::detail::CheckFieldsAreDeviceCopyable<(lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83), 4>' requested here
2605 | : CheckFieldsAreDeviceCopyable<FuncT, __builtin_num_fields(FuncT)>,
| ^
/usr/bin/compiler/../../include/sycl/types.hpp:2613:7: note: in instantiation of template class 'sycl::detail::CheckDeviceCopyable<(lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83)>' requested here
2613 | : CheckDeviceCopyable<KernelType> {};
| ^
/usr/bin/compiler/../../include/sycl/handler.hpp:1652:5: note: in instantiation of template class 'sycl::detail::CheckDeviceCopyable<sycl::detail::RoundedRangeKernel<sycl::item<1, true>, 1, (lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83)>>' requested here
1652 | detail::CheckDeviceCopyable<KernelType>();
| ^
/usr/bin/compiler/../../include/sycl/handler.hpp:1694:5: note: in instantiation of function template specialization 'sycl::handler::unpack<sycl::detail::RoundedRangeKernel<sycl::item<1, true>, 1, (lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83)>, sycl::detail::RoundedRangeKernel<sycl::item<1, true>, 1, (lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83)>, sycl::ext::oneapi::experimental::properties<std::tuple<>>, false, (lambda at /usr/bin/compiler/../../include/sycl/handler.hpp:1697:21)>' requested here
1694 | unpack<KernelName, KernelType, PropertiesT,
| ^
/usr/bin/compiler/../../include/sycl/handler.hpp:1293:7: note: in instantiation of function template specialization 'sycl::handler::kernel_parallel_for_wrapper<sycl::detail::RoundedRangeKernel<sycl::item<1, true>, 1, (lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83)>, sycl::item<1, true>, sycl::detail::RoundedRangeKernel<sycl::item<1, true>, 1, (lambda at /usr/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:1578:83)>, sycl::ext::oneapi::experimental::properties<std::tuple<>>>' requested here
1293 | kernel_parallel_for_wrapper<KName, TransformedArgType, decltype(Wrapper),
| ^
/usr/bin/compiler/../../include/sycl/handler.hpp:2332:5: note: (skipping 7 contexts in backtrace; use -ftemplate-backtrace-limit=0 to see all)
2332 | parallel_for_lambda_impl<KernelName, KernelType, 1, PropertiesT>(
| ^
[...]
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add a Functor definition followed by a sycl::is_device_copyable traits specialization to showcase what the user is expected to do.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

5f4bda3

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I meant also adding a definition for MyComparator and showing the actual Kokkos::sort invocation.
Do you think it is too obvious?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That commit changes the example into

MyComparator my_comparator;
     Kokkos::sort(exec, values, my_comparator);

I find this already to be pretty verbose and don't think we need to provide details on the potential implementation of MyComparator.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Oh I thought you'd transform the lambda into that comparator thing.
I suppose that is good enough.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The error you posted was with a lambda though right? (Not necessarily asking you to update)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, it was with a lambda but I printed only those errors that are independent of the example. The lambda referenced in it comes from oneDPL.


this is fixed by

.. code-block:: cpp

struct sycl::is_device_copyable<MyComparator>
: std::true_type {};


Mathematical functions
======================

Expand Down
Loading