RFC for "Passed Directly" Customization Point

rfcs/proposed/iterators_pass_directly_customization/README.md

@@ -0,0 +1,201 @@
+# Passed Directly Customization Point for User Defined Types
+
+## Introduction
+oneDPL processes data from some `std::vector::iterator` types automatically as input to its SYCL-based device backend by
+copying data between host and device, as described
+[here](https://uxlfoundation.github.io/oneDPL/parallel_api/pass_data_algorithms.html). Other iterator types, like
+Unified Shared Memory (USM) pointers, refer to data inherently accessible by the device, so no data transfers are
+required when passed into oneDPL APIs with a device policy. Iterators that do not require such data transfers are said
+to have the "passed directly" trait.
+
+oneDPL also defines rules for its provided
+[iterator types](https://uxlfoundation.github.io/oneDPL/parallel_api/iterators.html) to be "passed directly" under
+certain conditions, often based on their base types. When iterators lacking the "passed directly" trait are passed into
+oneDPL APIs with a device policy, additional overhead is incurred to transfer data to and from the device. Properly
+marking iterator types as "passed directly" is essential to avoid this overhead.
+
+Internally, these rules are defined using the trait `oneapi::dpl::__ranges::is_passed_directly<T>`, which evaluates to
+`std::true_type` or `std::false_type` to indicate whether the iterator type `T` should be passed directly to SYCL
+kernels, rather than copied to a SYCL buffer first to enable a transfer to and from the device. A legacy
+`is_passed_directly` trait exists, allowing types to define `using is_passed_directly = std::true_type;`. This legacy
+method is used for helper types in SYCLomatic compatibility headers (`device_pointer`, `device_iterator`,
+`tagged_pointer`, `constant_iterator`, `iterator_adaptor`). However, there is no official public API for users to define
+custom iterator types that can be passed directly to SYCL kernels.
+
+This gap should be addressed with an official public API. Without it, users must rely on oneDPL's provided iterator
+types or access implementation details outside the specified interface.
+
+## Proposal
+### High Level Proposal
+
+We will create an Argument-Dependent Lookup (ADL) customization point which defines if iterators are "passed directly"
+`is_passed_directly_in_onedpl_device_policies`. Users may define if their iterator types should be "passed directly" by
+oneDPL by defining a `constexpr` function in the same namespace where the iterator is defined.
+`is_passed_directly_in_onedpl_device_policies` must accept a const lvalue reference to the iterator type being
+specialized and return `std::true_type{}` if the iterator type is "passed directly" and `std::false_type{}` otherwise.
+
+Additionally, oneDPL will provide a public trait,
+```cpp
+namespace oneapi
+{
+namespace dpl
+{
+template <typename T>
+struct is_passed_directly_to_device; // std::true_type or std::false_type
+
+template <typename T>
+inline constexpr bool is_passed_directly_to_device_v<T> = oneapi::dpl::is_passed_directly_to_device::value;
+} // dpl
+} // oneapi
+```
+indicating if the iterator type `T` is "passed directly". This public trait is intended to be used to help define
+`is_passed_directly_in_onedpl_device_policies` for wrapper iterator types which depend on the "passed directly" status
+of their base iterator(s) using only the base iterator type(s) rather than a named instance. It may also be used to
+confirm user iterator types "passed directly" traits are as intended to prevent unnecessary overhead in oneDPL calls.
+
+### Additional Information
+
+The default implementation of the customization point `is_passed_directly_in_onedpl_device_policies` will be used to
+explicitly mark the following iterators as "passed directly":
+* Pointers (assumes USM pointers)
+* Iterators containing the legacy `using is_passed_directly = true` trait defined within the type definition
+* Iterators to USM shared allocated `std::vectors` (when knowable)
+* `std::reverse_iterator<Iter>` when `Iter` is also "passed directly"
+
+oneDPL will define the "passed directly" definitions of its custom iterators as follows:
+* `zip_iterator` is "passed directly" when all base iterators are "passed directly"
+* `counting_iterator` and `discard_iterator` are always "passed directly"
+* `transform_iterator` is "passed directly" if its source iterator is "passed directly"
+* `permutation_iterator` is "passed directly" if both its source iterator and its index map are "passed directly"
+
+### Implementation Details
+
+When using device policies, oneDPL will run compile-time checks on argument iterator types by using `decltype` to
+determine the return type of `is_passed_directly_in_onedpl_device_policies` when called with iterator as argument. If
+`std::true_type{}` is returned, oneDPL will pass the iterator directly to SYCL kernels rather than copying the data into
+`sycl::buffers` and using those buffers to transfer data to SYCL kernels.
+
+The specification and implementation will be prepared once this RFC is accepted as "proposed". We do not intend to offer
+this first as experimental. This RFC will target "Supported" once the specification and implementation are accepted.
+
+### Examples
+
+Below is a simple example of an iterator and ADL customization point definition which is always "passed directly".
+
+```cpp
+namespace user
+{
+
+    struct my_passed_directly_iterator
+    {
+        /* unspecified user definition */
+    };
+
+    constexpr
+    auto
+    is_passed_directly_in_onedpl_device_policies(const my_passed_directly_iterator&)
+    {
+        return std::true_type{};
+    }
+} //namespace user
+```
+
+Users can use any logic based on their iterator type to determine if it can be "passed directly", but the logic must be
+`constexpr`. Commonly, this will involve some check of base template types and their iterator type's "passed directly"
+status. Below is an example of a type that contains a pair of iterators and should be treated as "passed directly" if
+and only if both base iterators are also "passed directly". Note that the use of the public trait enables the use of the
+base iterator type alone without creating a named instance of the base iterator to pass into the ADL.
+
+```cpp
+namespace user
+{
+    template <typename It1, typename It2>
+    struct iterator_pair
+    {
+        It1 first;
+        It2 second;
+    };
+
+    template <typename It1, typename It2>
+    constexpr 
+    auto is_passed_directly_in_onedpl_device_policies(const iterator_pair<It1, It2>&)
+    {
+        if constexpr (oneapi::dpl::is_passed_directly_to_device_v<It1> &&
+                      oneapi::dpl::is_passed_directly_to_device_v<It2>)
+            return std::true_type{};
+        else 
+            return std::false_type{}; 
+    }
+} //namespace user
+```
+It is also possible to write overloads without a body using an `auto` trailing return type. The following is an
+alternative written in this way for the previous example:
+```cpp
+    template <typename It1, typename It2>
+    constexpr 
+    auto is_passed_directly_in_onedpl_device_policies(const iterator_pair<It1, It2>&) ->
+        std::bool_constant<(oneapi::dpl::is_passed_directly_to_device_v<It1> &&
+                            oneapi::dpl::is_passed_directly_to_device_v<It2>)>;
+```
+
+## Alternatives Considered
+### Public Trait Struct Explicit Specialization
+oneDPL could make public our internal structure `oneapi::dpl::__ranges::is_passed_directly` as
+`oneapi::dpl::is_passed_directly` for users to specialize to define rules for their types. This would be a similar
+mechanism to `sycl::is_device_copyable`. The implementation details of this option should avoid some complexity required
+ to properly implement the customization point.
+
+However, as we have learned from experience within oneDPL, explicit specialization of a structure in another library's
+namespace creates maintenance problems. It either requires lots of closing of nested namespaces, opening of the external
+ library's namespace for the specialization, or it requires separating these specializations to a separate location
+ removed from the types they are specializing for. oneDPL has chosen to use the latter, which can be seen in
+ `include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h`. This has led to several errors where changes to structures should
+  have included changes to sycl_traits but did not, and needed to be fixed later.
+
+In an effort to avoid this same issue for our users, we propose a similar method but instead with a constexpr
+customization point, allowing the user to override that customization point within their own namespace as a free
+function.
+
+### Require Specifically Named Typedef / Using in User's Type
+oneDPL could make official our requirements for users' types to include a typedef or using statement to define if the
+type is "passed directly" like `using is_passed_directly = std::true_type;`, where the absence of this would be
+equivalent to a `std::false_type`.
+
+However, this clutters the user type definitions with specifics of oneDPL. It also may not be as clear what this
+signifies for maintenance of user code without appropriate comments describing the details of oneDPL and SYCL. Users
+have expressed that this is undesirable.
+
+### Wrapper Class
+oneDPL could provide some wrapper iterator `direct_iterator` which wraps an arbitrary base iterator and marks it as
+"passed directly". `direct_iterator` could utilize either of the above alternatives to accomplish this and signal that
+the iterator should be "passed directly". It would need to pass through all operations to the wrapped base iterator and
+make sure no overhead is added in its usage.
+
+There is some complexity in adding such a wrapper iterator, and it would need to be considered carefully to make sure no
+ problems would be introduced. This wrapper class may obfuscate users' types and make them more unwieldy to use. It is
+ also less expressive than the other options in that it only has the ability to unilaterally mark a type as "passed
+ directly". There is no logic that can be used to express some iterator type which may be conditionally "passed
+ directly", other than to have logic to conditionally apply the wrapper in the first place. This option seems less clear
+ and has more opportunity to cause problems.
+
+## Testing
+We will need a detailed test checking both positive and negative responses to
+`oneapi::dpl::is_passed_directly_to_device_v` to ensure they have the expected result. This should
+include tests with custom types, combinations of iterators, USM pointers, etc.
+
+## Open Questions
+
+* How should we name the (1) ADL customization point and (2) public trait and value in a way that is concise and
+properly conveys the meaning to the users?
+    * Other names proposed:
+        * `oneapi::dpl::onedpl_is_iterator_device_ready`
+        * `oneapi::dpl::is_passed_directly_to_sycl_backend`
+        * `oneapi::dpl::requires_explicit_data_transfer_onedpl_device_policies` (inverted)
+        * `oneapi::dpl::is_passed_directly_to_device`
+        * `oneapi::dpl::is_passed_directly`
+    * This RFC provides different names for the (1) and (2). (1) will be used in the user's namespace, so it needs to
+      identify onedpl with its name, but (2) will be used from `oneapi::dpl` so including onedpl in the name is
+      redundant. Additionally, separating the names of (1) and (2) allows us to provide both a struct and value version
+      of the public trait, which is the norm. These decisions could be reconsidered in the specification PR.
+* Where should this be located?
+    * Possible options include `oneapi/dpl/iterator`, `oneapi/dpl/type_traits`.