[SYCL] sycl_khr_free_function_commands extension

sycl/include/sycl/ext/khr/free_function_commands.hpp

@@ -0,0 +1,304 @@
+#pragma once
+
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace khr {
+
+namespace detail {
+
+using empty_properties_t = ext::oneapi::experimental::empty_properties_t;
+
+template <typename CommandGroupFunc, typename PropertiesT>
+void submit(queue Q, PropertiesT Props, CommandGroupFunc &&CGF,
+            const sycl::detail::code_location &CodeLoc =
+                sycl::detail::code_location::current()) {
+  sycl::ext::oneapi::experimental::detail::submit_impl(
+      Q, Props, std::forward<CommandGroupFunc>(CGF), CodeLoc);
+}
+
+template <typename CommandGroupFunc, typename PropertiesT>
+event submit_tracked(queue Q, PropertiesT Props, CommandGroupFunc &&CGF,
+                     const sycl::detail::code_location &CodeLoc =
+                         sycl::detail::code_location::current()) {
+  return sycl::ext::oneapi::experimental::detail::submit_with_event_impl(
+      Q, Props, std::forward<CommandGroupFunc>(CGF), CodeLoc);
+}
+
+} // namespace detail
+
+template <typename CommandGroupFunc>
+void submit(queue q, CommandGroupFunc &&cgf,
+            const sycl::detail::code_location &code_loc =
+                sycl::detail::code_location::current()) {
+  detail::submit(q, detail::empty_properties_t{},
+                 std::forward<CommandGroupFunc>(cgf), code_loc);
+}
+
+template <typename CommandGroupFunc>
+event submit_tracked(queue q, CommandGroupFunc &&cgf,
+                     const sycl::detail::code_location &code_loc =
+                         sycl::detail::code_location::current()) {
+  return detail::submit_tracked(q, detail::empty_properties_t{},
+                                std::forward<CommandGroupFunc>(cgf), code_loc);
+}
+
+template <typename KernelType>
+void launch(handler &h, range<1> r, const KernelType &k) {
+  h.parallel_for<KernelType>(r, k);
+}
+template <typename KernelType>
+void launch(handler &h, range<2> r, const KernelType &k) {
+  h.parallel_for<KernelType>(r, k);
+}
+template <typename KernelType>
+void launch(handler &h, range<3> r, const KernelType &k) {
+  h.parallel_for<KernelType>(r, k);
+}
+
+template <typename KernelType>
+void launch(queue q, range<1> r, const KernelType &k) {
+  submit(q, [&](handler &h) { launch<KernelType>(h, r, k); });
+}
+template <typename KernelType>
+void launch(queue q, range<2> r, const KernelType &k) {
+  submit(q, [&](handler &h) { launch<KernelType>(h, r, k); });
+}
+template <typename KernelType>
+void launch(queue q, range<3> r, const KernelType &k) {
+  submit(q, [&](handler &h) { launch<KernelType>(h, r, k); });
+}
+
+template <typename... ArgsT>
+void launch(handler &h, range<1> r, const kernel &k, ArgsT &&...args) {
+  h.set_args<ArgsT...>(std::forward<ArgsT>(args)...);
+  h.parallel_for(r, k);
+}
+
+template <typename... ArgsT>
+void launch(handler &h, range<2> r, const kernel &k, ArgsT &&...args) {
+  h.set_args<ArgsT...>(std::forward<ArgsT>(args)...);
+  h.parallel_for(r, k);
+}
+
+template <typename... ArgsT>
+void launch(handler &h, range<3> r, const kernel &k, ArgsT &&...args) {
+  h.set_args<ArgsT...>(std::forward<ArgsT>(args)...);
+  h.parallel_for(r, k);
+}
+
+template <typename... ArgsT>
+void launch(queue q, range<1> r, const kernel &k, ArgsT &&...args) {
+  submit(q, [&](handler &h) {
+    parallel_for(h, r, k, std::forward<ArgsT>(args)...);
+  });
+}
+
+template <typename... ArgsT>
+void launch(queue q, range<2> r, const kernel &k, ArgsT &&...args) {
+  submit(q, [&](handler &h) {
+    parallel_for(h, r, k, std::forward<ArgsT>(args)...);
+  });
+}
+
+template <typename... ArgsT>
+void launch(queue q, range<3> r, const kernel &k, ArgsT &&...args) {
+  submit(q, [&](handler &h) {
+    parallel_for(h, r, k, std::forward<ArgsT>(args)...);
+  });
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_reduce(handler &h, range<1> r, const KernelType &k,
+                   Reductions &&...reductions) {
+  h.parallel_for<KernelType>(r, std::forward<Reductions>(reductions)..., k);
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_reduce(handler &h, range<2> r, const KernelType &k,
+                   Reductions &&...reductions) {
+  h.parallel_for<KernelType>(r, std::forward<Reductions>(reductions)..., k);
+}
+template <typename KernelType, typename... Reductions>
+void launch_reduce(handler &h, range<3> r, const KernelType &k,
+                   Reductions &&...reductions) {
+  h.parallel_for<KernelType>(r, std::forward<Reductions>(reductions)..., k);
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_reduce(queue q, range<1> r, const KernelType &k,
+                   Reductions &&...reductions) {
+  submit(q, [&](handler &h) {
+    launch_reduce<KernelType>(h, r, k, std::forward<Reductions>(reductions)...);
+  });
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_reduce(queue q, range<2> r, const KernelType &k,
+                   Reductions &&...reductions) {
+  submit(q, [&](handler &h) {
+    launch_reduce<KernelType>(h, r, k, std::forward<Reductions>(reductions)...);
+  });
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_reduce(queue q, range<3> r, const KernelType &k,
+                   Reductions &&...reductions) {
+  submit(q, [&](handler &h) {
+    launch_reduce<KernelType>(h, r, k, std::forward<Reductions>(reductions)...);
+  });
+}
+
+template <typename KernelType>
+void launch_grouped(handler &h, range<1> r, range<1> size,
+                    const KernelType &k) {
+  h.parallel_for<sycl::detail::auto_name>(nd_range<1>(r, size), k);
+}
+
+template <typename KernelType>
+void launch_grouped(handler &h, range<2> r, range<2> size,
+                    const KernelType &k) {
+  h.parallel_for<sycl::detail::auto_name>(nd_range<2>(r, size), k);
+}
+
+template <typename KernelType>
+void launch_grouped(handler &h, range<3> r, range<3> size,
+                    const KernelType &k) {
+  h.parallel_for<sycl::detail::auto_name>(nd_range<3>(r, size), k);
+}
+
+template <typename KernelType>
+void launch_grouped(queue q, range<1> r, range<1> size, const KernelType &k) {
+  submit(q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); });
+}
+template <typename KernelType>
+void launch_grouped(queue q, range<2> r, range<2> size, const KernelType &k) {
+  submit(q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); });
+}
+template <typename KernelType>
+void launch_grouped(queue q, range<3> r, range<3> size, const KernelType &k) {
+  submit(q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); });
+}
+
+template <typename... Args>
+void launch_grouped(sycl::handler &h, sycl::range<1> r, sycl::range<1> size,
+                    const sycl::kernel &k, Args &&...args) {
+  h.set_args<Args...>(std::forward<Args>(args)...);
+  h.parallel_for(nd_range<1>(r, size), k);
+}
+
+template <typename... Args>
+void launch_grouped(sycl::handler &h, sycl::range<2> r, sycl::range<2> size,
+                    const sycl::kernel &k, Args &&...args) {
+  h.set_args<Args...>(std::forward<Args>(args)...);
+  h.parallel_for(nd_range<2>(r, size), k);
+}
+
+template <typename... Args>
+void launch_grouped(sycl::handler &h, sycl::range<3> r, sycl::range<3> size,
+                    const sycl::kernel &k, Args &&...args) {
+  h.set_args<Args...>(std::forward<Args>(args)...);
+  h.parallel_for(nd_range<3>(r, size), k);
+}
+
+template <typename... Args>
+void launch_grouped(sycl::queue q, sycl::range<1> r, sycl::range<1> size,
+                    const sycl::kernel &k, Args &&...args) {
+  submit(q, [&](handler &h) {
+    launch_grouped(h, r, size, k, std::forward<Args>(args)...);
+  });
+}
+
+template <typename... Args>
+void launch_grouped(sycl::queue q, sycl::range<2> r, sycl::range<2> size,
+                    const sycl::kernel &k, Args &&...args) {
+  submit(q, [&](handler &h) {
+    launch_grouped(h, r, size, k, std::forward<Args>(args)...);
+  });
+}
+
+template <typename... Args>
+void launch_grouped(sycl::queue q, sycl::range<3> r, sycl::range<3> size,
+                    const sycl::kernel &k, Args &&...args) {
+  submit(q, [&](handler &h) {
+    launch_grouped(h, r, size, k, std::forward<Args>(args)...);
+  });
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_grouped_reduce(sycl::handler &h, sycl::range<1> r,
+                           sycl::range<1> size, const KernelType &k,
+                           Reductions &&...reductions) {
+  h.parallel_for<KernelType>(nd_range<1>(r, size),
+                             std::forward<Reductions>(reductions)..., k);
+}
+template <typename KernelType, typename... Reductions>
+void launch_grouped_reduce(sycl::handler &h, sycl::range<2> r,
+                           sycl::range<2> size, const KernelType &k,
+                           Reductions &&...reductions) {
+  h.parallel_for<KernelType>(nd_range<2>(r, size),
+                             std::forward<Reductions>(reductions)..., k);
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_grouped_reduce(sycl::handler &h, sycl::range<3> r,
+                           sycl::range<3> size, const KernelType &k,
+                           Reductions &&...reductions) {
+  h.parallel_for<KernelType>(nd_range<3>(r, size),
+                             std::forward<Reductions>(reductions)..., k);
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_grouped_reduce(sycl::queue q, sycl::range<1> r, sycl::range<1> size,
+                           const KernelType &k, Reductions &&...reductions) {
+  submit(q, [&](handler &h) {
+    launch_grouped_reduce<KernelType>(h, r, size, k,
+                                      std::forward<Reductions>(reductions)...);
+  });
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_grouped_reduce(sycl::queue q, sycl::range<2> r, sycl::range<2> size,
+                           const KernelType &k, Reductions &&...reductions) {
+  submit(q, [&](handler &h) {
+    launch_grouped_reduce<KernelType>(h, r, size, k,
+                                      std::forward<Reductions>(reductions)...);
+  });
+}
+
+template <typename KernelType, typename... Reductions>
+void launch_grouped_reduce(sycl::queue q, sycl::range<3> r, sycl::range<3> size,
+                           const KernelType &k, Reductions &&...reductions) {
+  submit(q, [&](handler &h) {
+    launch_grouped_reduce<KernelType>(h, r, size, k,
+                                      std::forward<Reductions>(reductions)...);
+  });
+}
+
+template <typename KernelType>
+void single_task(handler &h, const KernelType &k) {
+  h.single_task<KernelType>(k);
+}
+
+template <typename KernelType>
+void launch_task(sycl::queue q, const KernelType &k,
+                 const sycl::detail::code_location &code_loc =
+                     sycl::detail::code_location::current()) {
+  submit(q, [&](handler &h) { single_task<KernelType>(h, k); }, code_loc);
+}
+
+template <typename... Args>
+void launch_task(sycl::handler &h, const sycl::kernel &k, Args &&...args) {
+  h.set_args<Args...>(std::forward<Args>(args)...);
+  h.single_task(k);
+}
+
+template <typename... Args>
+void launch_task(queue q, const kernel &k, Args &&...args) {
+  submit(q,
+         [&](handler &h) { launch_task(h, k, std::forward<Args>(args)...); });
+}
+} // namespace khr
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/sycl.hpp

@@ -74,6 +74,7 @@
 #include <sycl/ext/intel/experimental/task_sequence_properties.hpp>
 #include <sycl/ext/intel/experimental/usm_properties.hpp>
 #include <sycl/ext/intel/usm_pointers.hpp>
+#include <sycl/ext/khr/free_function_commands.hpp>
 #include <sycl/ext/oneapi/bfloat16.hpp>
 #include <sycl/ext/oneapi/bindless_images.hpp>
 #include <sycl/ext/oneapi/device_global/device_global.hpp>

sycl/source/feature_test.hpp.in

@@ -112,6 +112,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_WORK_GROUP_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_SCRATCH_MEMORY 1
 #define SYCL_EXT_ONEAPI_WORK_GROUP_STATIC 1
+#define SYCL_KHR_FREE_FUNCTION_COMMANDS 1 
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0