RFC for "Passed Directly" Customization Point

.gitignore

@@ -1,5 +1,5 @@
 # -------- CMake --------
-build/*
+build*/*
 
 # -------- IDE --------
 .vscode/*

rfcs/proposed/iterators_pass_directly_customization/README.md

@@ -0,0 +1,151 @@
+# Passed Directly Customization Point for User Defined Types
+
+## Introduction
+
+oneDPL handles some types of input data automatically as input to its dpcpp (sycl-based) backend as described
+[here](https://uxlfoundation.github.io/oneDPL/parallel_api/pass_data_algorithms.html). Unified Shared Memory (USM)
+pointers refer to data which is device accessible inherently, so no processing is required to pass this type of input
+data to SYCL kernels, we refer to this trait as "passed directly". oneDPL also defines some rules for its provided
+[iterator types](https://uxlfoundation.github.io/oneDPL/parallel_api/iterators.html) to be passed directly to SYCL
+under some circumstances (based usually on their base types).
+
+Internally, these rules are defined with a 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. There exists an unofficial legacy `is_passed_directly` trait which types can define like this:
+`using is_passed_directly = std::true_type;` which is supported within oneDPL. This method is currently used for a
+number of helper types within the SYCLomatic compatibility headers, (`device_pointer`, `device_iterator`,
+`tagged_pointer`, `constant_iterator`, `iterator_adaptor`). There is no official public API for users who want to
+create their own iterator types which could be passed directly to SYCL kernels, this is a gap that should be filled 
+with an official public API.
+
+Without something like this users are forced to only rely upon our provided iterator types, or reach into implementation
+details which are not part of oneDPL's specified interface.
+
+## Proposal
+
+Create a customization point `oneapi::dpl::is_passed_directly_in_onedpl_device_policies` free function which allows
+users to mark their types as passed directly:
+
+```
+template <typename T>
+constexpr bool is_passed_directly_in_onedpl_device_policies(const T&);
+```
+
+oneDPL will provide a default implementation which will defer to the existing trait:
+
+```
+template <typename T>
+constexpr
+bool
+is_passed_directly_in_onedpl_device_policies(const T&)
+{
+	return oneapi::dpl::__ranges::is_passed_directly_v<T>;
+}
+```
+
+Below is a simple example of a type and customization point definition which is always passed directly.
+
+```
+namespace user
+{
+
+    struct my_passed_directly_type
+    {
+        /* unspecified user definition */
+    };
+
+    constexpr
+    bool
+    is_passed_directly_in_onedpl_device_policies(const my_passed_directly_type&)
+    {
+        return true;
+    }
+} //namespace user
+```
+
+Users can use any constexpr logic based on their type to determine if the type can be passed directly into a SYCL kernel
+without any processing. Below is an example of a type which contains a pair of iterators, and should be treated as
+passed directly if and only if both base iterators are also passed directly. oneDPL will use this customization point
+internally when determining how to handle incoming data, picking up any user defined customizations.
+
+When using device policies, oneDPL will check iterator types by calling `is_passed_directly_in_onedpl_device_policies`
+and if `true` is returned, will pass the iterator directly to sycl kernels rather than the data into sycl buffers, and
+then using the sycl buffer to handle the device accessibility of the data. Users may also call
+`oneapi::dpl::is_passed_directly_in_onedpl_device_policies` themselves to check how the oneDPL internals will treat any
+iterator types. This may be useful to ensure that no extra overhead occurs in device policy calls.
+
+```
+namespace user
+{
+    template <typename It1, typename It2>
+    struct iterator_pair 
+    {
+        It1 first;
+        It2 second;
+    };
+
+    template <typename It1, typename It2>
+    constexpr
+    bool
+    is_passed_directly_in_onedpl_device_policies(const iterator_pair<It1, It2>& pair)
+    {
+        return oneapi::dpl::is_passed_directly_in_onedpl_device_policies(pair.first) &&
+               oneapi::dpl::is_passed_directly_in_onedpl_device_policies(pair.second);
+    }
+} //namespace user
+```
+
+This allows the user to provide rules for their types next to their implementation, without cluttering the
+implementation of the type itself with extra typedefs, etc.
+
+This option can exist in concert with existing methods, the legacy `is_passed_directly` typedef in types, the internal
+`oneapi::dpl::__ranges::is_passed_directly` trait specializations. It would be possible to simplify the internal
+implementation away from explicit specializations of the trait to the customization point, but that is not required
+at first implementation.
+
+### Implementation Details
+To make this robust, we will follow a C++17 updated version of what is discussed in
+[Eric Niebler's Post](https://ericniebler.com/2014/10/21/customization-point-design-in-c11-and-beyond/), using a
+callable, and using an `inline constexpr` to avoid issues with ODR and to avoid issues with resolving customization
+points when not separating the call to two steps with a `using` statement first. Using his proposed method will allow
+both qualified and unqualified calls to `is_passed_directly_in_onedpl_device_policies` to pick up the default
+implementation provided by oneDPL.
+
+## Alternatives considered
+### Public Trait Struct Explicit Specialization
+We could simply 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 makes for 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 later, which can be seen in
+`include/oneapi/dpl/pstl/hetero/dpcpp/sycl_traits.h`. This has made for 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
+We could simply make official our requirements for user's 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.
+
+## Testing
+We will need a detailed test checking both positive and negative responses to
+`is_passed_directly_in_onedpl_device_policies` have the expected result, with custom types and combinations of
+iterators, usm pointers etc.
+
+## Open Questions
+
+* Is there a better, more concise name than `is_passed_directly_in_onedpl_device_policies` that properly conveys the
+meaning to the users?
+* Should we be targeting Experimental or fully Supported with this proposal?
+(Do we think user feedback is required to solidify an interface / experience?)