diff --git a/adoc/extensions/sycl_khr_free_function_commands.adoc b/adoc/extensions/sycl_khr_free_function_commands.adoc index 0efb0ee2..4406e187 100644 --- a/adoc/extensions/sycl_khr_free_function_commands.adoc +++ b/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 -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 -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...); })#. '''