From 49a29300b1ad0d21a401ea19afcc8d51b90ded9e Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Mon, 14 Apr 2025 17:48:24 -0700 Subject: [PATCH 1/5] [SYCL][COMPAT] Add launch lambda overload --- sycl/include/syclcompat/launch.hpp | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/sycl/include/syclcompat/launch.hpp b/sycl/include/syclcompat/launch.hpp index eb5d774bc12d3..62cea8f1600c0 100644 --- a/sycl/include/syclcompat/launch.hpp +++ b/sycl/include/syclcompat/launch.hpp @@ -112,12 +112,27 @@ launch(const dim3 &grid, const dim3 &threads, sycl::queue q, Args... args) { return launch(sycl::nd_range<3>{grid * threads, threads}, q, args...); } + + template std::enable_if_t, sycl::event> launch(const dim3 &grid, const dim3 &threads, Args... args) { return launch(grid, threads, get_default_queue(), args...); } +// Overload taking zero-argument function object +template +std::enable_if_t, sycl::event> +launch(const F& f, const sycl::nd_range &range, sycl::queue q=get_default_queue()) { + return q.parallel_for(detail::transform_nd_range(range), [=](sycl::nd_item) { f(); }); +} +// Alternative launch through dim3 objects +template +std::enable_if_t, sycl::event> +launch(const F& f, const dim3 &grid, const dim3 &threads, sycl::queue q=get_default_queue()) { + return launch(f, sycl::nd_range<3>{grid * threads, threads}, q); +} + } // namespace syclcompat namespace syclcompat::experimental { From e33b16bd4e10a6f7e38f853ede2f298be5b3241c Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Mon, 14 Apr 2025 17:53:05 -0700 Subject: [PATCH 2/5] format --- sycl/include/syclcompat/launch.hpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/include/syclcompat/launch.hpp b/sycl/include/syclcompat/launch.hpp index 62cea8f1600c0..e2030cbd67bb2 100644 --- a/sycl/include/syclcompat/launch.hpp +++ b/sycl/include/syclcompat/launch.hpp @@ -112,8 +112,6 @@ launch(const dim3 &grid, const dim3 &threads, sycl::queue q, Args... args) { return launch(sycl::nd_range<3>{grid * threads, threads}, q, args...); } - - template std::enable_if_t, sycl::event> launch(const dim3 &grid, const dim3 &threads, Args... args) { From d9667ef996fbdf5617b92efca568ab8794ccd797 Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Mon, 14 Apr 2025 20:07:48 -0700 Subject: [PATCH 3/5] format --- sycl/include/syclcompat/launch.hpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/include/syclcompat/launch.hpp b/sycl/include/syclcompat/launch.hpp index e2030cbd67bb2..909bb5d18e752 100644 --- a/sycl/include/syclcompat/launch.hpp +++ b/sycl/include/syclcompat/launch.hpp @@ -121,13 +121,15 @@ launch(const dim3 &grid, const dim3 &threads, Args... args) { // Overload taking zero-argument function object template std::enable_if_t, sycl::event> -launch(const F& f, const sycl::nd_range &range, sycl::queue q=get_default_queue()) { +launch(const F& f, const sycl::nd_range &range, + sycl::queue q=get_default_queue()) { return q.parallel_for(detail::transform_nd_range(range), [=](sycl::nd_item) { f(); }); } // Alternative launch through dim3 objects template std::enable_if_t, sycl::event> -launch(const F& f, const dim3 &grid, const dim3 &threads, sycl::queue q=get_default_queue()) { +launch(const F& f, const dim3 &grid, const dim3 &threads, + sycl::queue q=get_default_queue()) { return launch(f, sycl::nd_range<3>{grid * threads, threads}, q); } From 8e4f15c74836aab1b36d9fc5e3b49785c2293e52 Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Mon, 14 Apr 2025 20:09:02 -0700 Subject: [PATCH 4/5] format --- sycl/include/syclcompat/launch.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/syclcompat/launch.hpp b/sycl/include/syclcompat/launch.hpp index 909bb5d18e752..62a95be9e602a 100644 --- a/sycl/include/syclcompat/launch.hpp +++ b/sycl/include/syclcompat/launch.hpp @@ -123,7 +123,8 @@ template std::enable_if_t, sycl::event> launch(const F& f, const sycl::nd_range &range, sycl::queue q=get_default_queue()) { - return q.parallel_for(detail::transform_nd_range(range), [=](sycl::nd_item) { f(); }); + return q.parallel_for(detail::transform_nd_range(range), + [=](sycl::nd_item) { f(); }); } // Alternative launch through dim3 objects template From 197ab960e947dabb1872f3005d3f35cdee9739b4 Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Mon, 14 Apr 2025 20:19:38 -0700 Subject: [PATCH 5/5] format --- sycl/include/syclcompat/launch.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/include/syclcompat/launch.hpp b/sycl/include/syclcompat/launch.hpp index 62a95be9e602a..47e17b899f924 100644 --- a/sycl/include/syclcompat/launch.hpp +++ b/sycl/include/syclcompat/launch.hpp @@ -120,17 +120,17 @@ launch(const dim3 &grid, const dim3 &threads, Args... args) { // Overload taking zero-argument function object template -std::enable_if_t, sycl::event> -launch(const F& f, const sycl::nd_range &range, - sycl::queue q=get_default_queue()) { +std::enable_if_t, sycl::event> +launch(const F &f, const sycl::nd_range &range, + sycl::queue q = get_default_queue()) { return q.parallel_for(detail::transform_nd_range(range), [=](sycl::nd_item) { f(); }); } // Alternative launch through dim3 objects template -std::enable_if_t, sycl::event> -launch(const F& f, const dim3 &grid, const dim3 &threads, - sycl::queue q=get_default_queue()) { +std::enable_if_t, sycl::event> +launch(const F &f, const dim3 &grid, const dim3 &threads, + sycl::queue q = get_default_queue()) { return launch(f, sycl::nd_range<3>{grid * threads, threads}, q); }