Define kernel object overloads via equivalence

adoc/extensions/sycl_khr_free_function_commands.adoc

@@ -204,9 +204,16 @@ void launch(sycl::queue q, sycl::range<3> r,
 
 }
 ----
-_Effects_: Enqueues a kernel object that is invoked for every work-item in the
-specified [code]#sycl::range#.
-The arguments in [code]#args# are passed to the kernel in the same order.
+_Effects (1-3)_: Equivalent to:
+
+[source,c++]
+----
+h.set_args(args...);
+h.parallel_for(r, k);
+----
+
+_Effects (4-6)_: Equivalent to [code]#q.submit([&](handler& h) { launch(h, r, k,
+args...); })#.
 
 '''
 
@@ -283,10 +290,15 @@ void launch_reduce(sycl::queue q, sycl::range<3> r,
 
 }
 ----
-_Effects_: Enqueues a kernel object that is invoked for every work-item in the
-specified [code]#sycl::range#, where each work-item contributes to one or more
-reductions.
-The arguments in [code]#args# are passed to the kernel in the same order.
+_Effects (1-3)_: Equivalent to:
+[source,c++]
+----
+h.set_args(args...);
+h.parallel_for(r, k);
+----
+
+_Effects (4-6)_: Equivalent to [code]#q.submit([&](handler& h) {
+launch_reduce(h, r, k, args...); })#.
 
 '''
 
@@ -359,10 +371,15 @@ void launch_grouped(sycl::queue q, sycl::range<3> r, sycl::range<3> size,
 
 }
 ----
-_Effects_: Enqueues a kernel object that is invoked for every work-item in the
-specified [code]#sycl::range#.
-Work-items are grouped into work-groups of size [code]#size#.
-The arguments in [code]#args# are passed to the kernel in the same order.
+_Effects (1-3)_: Equivalent to:
+[source,c++]
+----
+h.set_args(args...);
+h.parallel_for(nd_range(r, size), k);
+----
+
+_Effects (4-6)_: Equivalent to [code]#q.submit([&](handler& h) {
+launch_grouped(h, r, size, k, args...); })#.
 
 '''
 
@@ -451,11 +468,15 @@ void launch_grouped_reduce(sycl::queue q, sycl::range<3> r,
 
 }
 ----
-_Effects_: Enqueues a kernel object that is invoked for every work-item in the
-specified [code]#sycl::range#, where each work-item contributes to one or more
-reductions.
-Work-items are grouped into work-groups of size [code]#size#.
-The arguments in [code]#args# are passed to the kernel in the same order.
+_Effects (1-3)_: Equivalent to:
+[source,c++]
+----
+h.set_args(args...);
+h.parallel_for(nd_range(r, size), k);
+----
+
+_Effects (4-6)_: Equivalent to [code]#q.submit([&](handler& h) {
+launch_grouped_reduce(h, r, size, k, args...); })#.
 
 '''
 
@@ -485,15 +506,22 @@ k); })#.
 namespace sycl::khr {
 
 template <typename Args...>
-void launch_task(sycl::queue q, const sycl::kernel& k, Args&&... args);    (1)
+void launch_task(sycl::handler& h, const sycl::kernel& k, Args&&... args); (1)
 
 template <typename Args...>
-void launch_task(sycl::handler& h, const sycl::kernel& k, Args&&... args); (2)
+void launch_task(sycl::queue q, const sycl::kernel& k, Args&&... args);    (2)
 
 }
 ----
-_Effects_: Enqueues a kernel object as a device task.
-The arguments in [code]#args# are passed to the kernel in the same order.
+_Effects (1)_: Equivalent to:
+[source,c++]
+----
+h.set_args(args...);
+h.parallel_for(k);
+----
+
+_Effects (2)_: Equivalent to [code]#q.submit([&](handler& h) { launch_task(h, k,
+args...); })#.
 
 '''