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.

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

Add queue::submit and handler::parallel_for overloads accepting property_list #602

Open
al42and opened this issue Aug 5, 2024 · 0 comments

Comments

@al42and
Copy link

al42and commented Aug 5, 2024

Properties is SYCL serve as a generic way to expanding the behavior of a class or function, either with standardized properties or with vendor-specific ones.

Submitting tasks to a queue is, in a way, the most important operation in heterogeneous programming, so it seems reasonable that there could be need to customize it with various runtime properties (compile-time properties are a separate issue; I am primarily concerned with run-time ones here). However, neither queue::submit nor handler::parallel_for (and it's variations) allow specifying a property list in SYCL 2020.

The need to provide properties is definitely there, as both major implementations have to work around that:

  • AdaptiveCpp has ACPP_EXT_CG_PROPERTY_ extension family, which introduce queue::submit(const property_list& prop_list, T cgf) function. It can be used not only for kernels, but for other submissions as well.
  • DPC++ has sycl_ext_oneapi_kernel_properties which introduces handler::single_task(PropertyList properties, const KernelType &kernelFunc); and similar overloads for parallel_for and their shorthands in sycl::queue. The extension is focused on compile-time properties but notes that "an implementation may support additional properties which could have run-time values".

From a user perspective, the difference in function signatures mean that if extensions from both vendors are used, the kernel invocation code starts looking like:

#if DPCPP
#  define DPCPP_PROPERTY sycl::ext::oneapi::experimental::properties { /*...*/ },
#  define ACPP_PROPERTY
#elif ACPP
#  define DPCPP_PROPERTY 
#  define ACPP_PROPERTY sycl::property_list { /*...*/ },
#else
#  define DPCPP_PROPERTY
#  define ACPP_PROPERTY
#endif
    q.submit([&](ACPP_PROPERTY sycl::handler& cgh) {
        cgh.parallel_for(range, DPCPP_PROPERTY [=](/*...*/) { /*...*/ });
    });

In particular, the comma at the end of the macro is very nasty. A wrapper functions like my_parallel_for(sycl::queue, sycl::range<Dim>, Kernel&& f) could make the code prettier, but this still is not optimal from the user perspective, given that SYCL standard has a much better solution for this problem in other places: optional sycl::property_list parameter.

In SYCL 2020, sycl::property_list is mainly used in constructors, but there are cases where it is used in free functions, such as sycl::malloc_device &co, where property_list is explicitly said to be added for future extensibility, with no suitable properties defined by the standard. I see little reason to limit this extensibility to memory allocation alone.

The suggestion is to add overloads to queue::submit, queue::single_task, queue::parallel_for, queue::memcpy, queue::copy, queue::memset, queue::fill, handler::single_task, handler::parallel_for, handler::parallel_for_work_group, handler::memcpy, handler::copy, handler::memset, and handler::fill functions with a property_list parameter (right before the kernel parameter) to allow users to pass run-time properties. Even if the construction of property_lists relies on macros to construct vendor-specific property objects, the kernel invocation code will be much cleaner and more aligned with how the rest of SYCL API is organized.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

1 participant