Do not use sg.get_local_range

Use sg.get_max_local_range instead. The `sg.get_local_range` must perform lots of checks to determine if this is the last trailing sub-group in the work-group and its actual size may be smaller. We set the local work-group size to be 128, which is a multiple of any sub-group size, and hence get_local_range() always equals to get_max_local_raneg(). The size of the work-groups was increated from 128 to 256, which is chosen so that all 8 threads of single vector with simd32 are used. Set vec_sz and n_vecs in implementations of contig_impl for each support function Make local work-groups size dependent on number of elements to process Fixes for type dispatching utils 1. Add missing include <type_traits> needed for std::true_type, and std::disjunction, std::conjunction 2. Replace std::bool_constant<std::same_v<T1, T2>> with direct and simpler std::same<T1, T2> in couple of instances Hide hyperparameter selection struct in anonymous namespace
ndgrigorian · Nov 22, 2024 · Oct 21, 2024 · Oct 21, 2024 · Oct 28, 2024 · Nov 12, 2024
commit 02a7c5ab2b68f71c93107553d8b499ddd89cf2e8
@@ -32,9 +32,11 @@
 #include <type_traits>
 
 #include "cabs_impl.hpp"
-#include "kernels/elementwise_functions/common.hpp"
+#include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
+#include "kernels/elementwise_functions/common.hpp"
+
 #include "utils/offset_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 #include "utils/type_utils.hpp"
@@ -89,8 +91,8 @@ template <typename argT, typename resT> struct AbsFunctor
 
 template <typename argT,
           typename resT = argT,
-          unsigned int vec_sz = 4,
-          unsigned int n_vecs = 2,
+          unsigned int vec_sz = 4u,
+          unsigned int n_vecs = 2u,
           bool enable_sg_loadstore = true>
 using AbsContigFunctor =
     elementwise_common::UnaryContigFunctor<argT,
@@ -122,6 +124,24 @@ template <typename T> struct AbsOutputType
     static constexpr bool is_defined = !std::is_same_v<value_type, void>;
 };
 
+namespace
+{
+
+namespace vsu_ns = dpctl::tensor::kernels::vec_size_utils;
+
+using vsu_ns::ContigHyperparameterSetDefault;
+
+template <typename argTy> struct AbsContigHyperparameterSet
+{
+    using value_type =
+        typename std::disjunction<ContigHyperparameterSetDefault<4u, 2u>>;
+
+    constexpr static auto vec_sz = value_type::vec_sz;
+    constexpr static auto n_vecs = value_type::n_vecs;
+};
+
+} // namespace
+
 template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
 class abs_contig_kernel;
 
@@ -132,9 +152,12 @@ sycl::event abs_contig_impl(sycl::queue &exec_q,
                             char *res_p,
                             const std::vector<sycl::event> &depends = {})
 {
+    constexpr unsigned int vec_sz = AbsContigHyperparameterSet<argTy>::vec_sz;
+    constexpr unsigned int n_vec = AbsContigHyperparameterSet<argTy>::n_vecs;
+
     return elementwise_common::unary_contig_impl<
-        argTy, AbsOutputType, AbsContigFunctor, abs_contig_kernel>(
-        exec_q, nelems, arg_p, res_p, depends);
+        argTy, AbsOutputType, AbsContigFunctor, abs_contig_kernel, vec_sz,
+        n_vec>(exec_q, nelems, arg_p, res_p, depends);
 }
 
 template <typename fnT, typename T> struct AbsContigFactory

@@ -29,10 +29,12 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "kernels/elementwise_functions/common.hpp"
 #include "sycl_complex.hpp"
+#include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
+#include "kernels/elementwise_functions/common.hpp"
+
 #include "utils/offset_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 #include "utils/type_utils.hpp"
@@ -128,8 +130,8 @@ template <typename argT, typename resT> struct AcosFunctor
 
 template <typename argTy,
           typename resTy = argTy,
-          unsigned int vec_sz = 4,
-          unsigned int n_vecs = 2,
+          unsigned int vec_sz = 4u,
+          unsigned int n_vecs = 2u,
           bool enable_sg_loadstore = true>
 using AcosContigFunctor =
     elementwise_common::UnaryContigFunctor<argTy,
@@ -156,6 +158,24 @@ template <typename T> struct AcosOutputType
     static constexpr bool is_defined = !std::is_same_v<value_type, void>;
 };
 
+namespace
+{
+
+namespace vsu_ns = dpctl::tensor::kernels::vec_size_utils;
+
+using vsu_ns::ContigHyperparameterSetDefault;
+
+template <typename argTy> struct AcosContigHyperparameterSet
+{
+    using value_type =
+        typename std::disjunction<ContigHyperparameterSetDefault<4u, 2u>>;
+
+    constexpr static auto vec_sz = value_type::vec_sz;
+    constexpr static auto n_vecs = value_type::n_vecs;
+};
+
+} // namespace
+
 template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
 class acos_contig_kernel;
 
@@ -166,9 +186,12 @@ sycl::event acos_contig_impl(sycl::queue &exec_q,
                              char *res_p,
                              const std::vector<sycl::event> &depends = {})
 {
+    constexpr unsigned int vec_sz = AcosContigHyperparameterSet<argTy>::vec_sz;
+    constexpr unsigned int n_vec = AcosContigHyperparameterSet<argTy>::n_vecs;
+
     return elementwise_common::unary_contig_impl<
-        argTy, AcosOutputType, AcosContigFunctor, acos_contig_kernel>(
-        exec_q, nelems, arg_p, res_p, depends);
+        argTy, AcosOutputType, AcosContigFunctor, acos_contig_kernel, vec_sz,
+        n_vec>(exec_q, nelems, arg_p, res_p, depends);
 }
 
 template <typename fnT, typename T> struct AcosContigFactory

@@ -29,10 +29,12 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "kernels/elementwise_functions/common.hpp"
 #include "sycl_complex.hpp"
+#include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
+#include "kernels/elementwise_functions/common.hpp"
+
 #include "utils/offset_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 #include "utils/type_utils.hpp"
@@ -155,8 +157,8 @@ template <typename argT, typename resT> struct AcoshFunctor
 
 template <typename argTy,
           typename resTy = argTy,
-          unsigned int vec_sz = 4,
-          unsigned int n_vecs = 2,
+          unsigned int vec_sz = 4u,
+          unsigned int n_vecs = 2u,
           bool enable_sg_loadstore = true>
 using AcoshContigFunctor =
     elementwise_common::UnaryContigFunctor<argTy,
@@ -183,6 +185,25 @@ template <typename T> struct AcoshOutputType
     static constexpr bool is_defined = !std::is_same_v<value_type, void>;
 };
 
+namespace
+{
+
+namespace vsu_ns = dpctl::tensor::kernels::vec_size_utils;
+
+using vsu_ns::ContigHyperparameterSetDefault;
+using vsu_ns::UnaryContigHyperparameterSetEntry;
+
+template <typename argTy> struct AcoshContigHyperparameterSet
+{
+    using value_type =
+        typename std::disjunction<ContigHyperparameterSetDefault<4u, 2u>>;
+
+    constexpr static auto vec_sz = value_type::vec_sz;
+    constexpr static auto n_vecs = value_type::n_vecs;
+};
+
+} // namespace
+
 template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
 class acosh_contig_kernel;
 
@@ -193,9 +214,12 @@ sycl::event acosh_contig_impl(sycl::queue &exec_q,
                               char *res_p,
                               const std::vector<sycl::event> &depends = {})
 {
+    constexpr unsigned int vec_sz = AcoshContigHyperparameterSet<argTy>::vec_sz;
+    constexpr unsigned int n_vec = AcoshContigHyperparameterSet<argTy>::n_vecs;
+
     return elementwise_common::unary_contig_impl<
-        argTy, AcoshOutputType, AcoshContigFunctor, acosh_contig_kernel>(
-        exec_q, nelems, arg_p, res_p, depends);
+        argTy, AcoshOutputType, AcoshContigFunctor, acosh_contig_kernel, vec_sz,
+        n_vec>(exec_q, nelems, arg_p, res_p, depends);
 }
 
 template <typename fnT, typename T> struct AcoshContigFactory

@@ -30,6 +30,8 @@
 #include <type_traits>
 
 #include "sycl_complex.hpp"
+#include "vec_size_util.hpp"
+
 #include "utils/offset_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 #include "utils/type_utils.hpp"
@@ -110,8 +112,8 @@ template <typename argT1, typename argT2, typename resT> struct AddFunctor
 template <typename argT1,
           typename argT2,
           typename resT,
-          unsigned int vec_sz = 4,
-          unsigned int n_vecs = 2,
+          unsigned int vec_sz = 4u,
+          unsigned int n_vecs = 2u,
           bool enable_sg_loadstore = true>
 using AddContigFunctor =
     elementwise_common::BinaryContigFunctor<argT1,
@@ -196,6 +198,43 @@ template <typename T1, typename T2> struct AddOutputType
     static constexpr bool is_defined = !std::is_same_v<value_type, void>;
 };
 
+namespace
+{
+
+namespace vsu_ns = dpctl::tensor::kernels::vec_size_utils;
+
+using vsu_ns::BinaryContigHyperparameterSetEntry;
+using vsu_ns::ContigHyperparameterSetDefault;
+
+template <typename argTy1, typename argTy2> struct AddContigHyperparameterSet
+{
+    using value_type = typename std::disjunction<
+        BinaryContigHyperparameterSetEntry<argTy1,
+                                           std::int64_t,
+                                           argTy2,
+                                           std::int64_t,
+                                           1u,
+                                           2u>,
+        BinaryContigHyperparameterSetEntry<argTy1,
+                                           std::uint64_t,
+                                           argTy2,
+                                           std::uint64_t,
+                                           1u,
+                                           2u>,
+        BinaryContigHyperparameterSetEntry<argTy1,
+                                           double,
+                                           argTy2,
+                                           double,
+                                           1u,
+                                           2u>,
+        ContigHyperparameterSetDefault<4u, 2u>>;
+
+    constexpr static auto vec_sz = value_type::vec_sz;
+    constexpr static auto n_vecs = value_type::n_vecs;
+};
+
+} // end of anonymous namespace
+
 template <typename argT1,
           typename argT2,
           typename resT,
@@ -214,10 +253,13 @@ sycl::event add_contig_impl(sycl::queue &exec_q,
                             ssize_t res_offset,
                             const std::vector<sycl::event> &depends = {})
 {
+    constexpr auto vec_sz = AddContigHyperparameterSet<argTy1, argTy2>::vec_sz;
+    constexpr auto n_vecs = AddContigHyperparameterSet<argTy1, argTy2>::n_vecs;
+
     return elementwise_common::binary_contig_impl<
-        argTy1, argTy2, AddOutputType, AddContigFunctor, add_contig_kernel>(
-        exec_q, nelems, arg1_p, arg1_offset, arg2_p, arg2_offset, res_p,
-        res_offset, depends);
+        argTy1, argTy2, AddOutputType, AddContigFunctor, add_contig_kernel,
+        vec_sz, n_vecs>(exec_q, nelems, arg1_p, arg1_offset, arg2_p,
+                        arg2_offset, res_p, res_offset, depends);
 }
 
 template <typename fnT, typename T1, typename T2> struct AddContigFactory
@@ -410,8 +452,8 @@ template <typename argT, typename resT> struct AddInplaceFunctor
 
 template <typename argT,
           typename resT,
-          unsigned int vec_sz = 4,
-          unsigned int n_vecs = 2,
+          unsigned int vec_sz = 4u,
+          unsigned int n_vecs = 2u,
           bool enable_sg_loadstore = true>
 using AddInplaceContigFunctor = elementwise_common::BinaryInplaceContigFunctor<
     argT,
@@ -489,9 +531,13 @@ add_inplace_contig_impl(sycl::queue &exec_q,
                         ssize_t res_offset,
                         const std::vector<sycl::event> &depends = {})
 {
+    constexpr auto vec_sz = AddContigHyperparameterSet<resTy, argTy>::vec_sz;
+    constexpr auto n_vecs = AddContigHyperparameterSet<resTy, argTy>::n_vecs;
+
     return elementwise_common::binary_inplace_contig_impl<
-        argTy, resTy, AddInplaceContigFunctor, add_inplace_contig_kernel>(
-        exec_q, nelems, arg_p, arg_offset, res_p, res_offset, depends);
+        argTy, resTy, AddInplaceContigFunctor, add_inplace_contig_kernel,
+        vec_sz, n_vecs>(exec_q, nelems, arg_p, arg_offset, res_p, res_offset,
+                        depends);
 }
 
 template <typename fnT, typename T1, typename T2> struct AddInplaceContigFactory

@@ -30,10 +30,12 @@
 #include <sycl/sycl.hpp>
 #include <type_traits>
 
-#include "kernels/elementwise_functions/common.hpp"
 #include "sycl_complex.hpp"
+#include "vec_size_util.hpp"
 
 #include "kernels/dpctl_tensor_types.hpp"
+#include "kernels/elementwise_functions/common.hpp"
+
 #include "utils/offset_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 #include "utils/type_utils.hpp"
@@ -74,8 +76,8 @@ template <typename argT, typename resT> struct AngleFunctor
 
 template <typename argTy,
           typename resTy = argTy,
-          unsigned int vec_sz = 4,
-          unsigned int n_vecs = 2,
+          unsigned int vec_sz = 4u,
+          unsigned int n_vecs = 2u,
           bool enable_sg_loadstore = true>
 using AngleContigFunctor =
     elementwise_common::UnaryContigFunctor<argTy,
@@ -99,6 +101,25 @@ template <typename T> struct AngleOutputType
     static constexpr bool is_defined = !std::is_same_v<value_type, void>;
 };
 
+namespace
+{
+
+namespace vsu_ns = dpctl::tensor::kernels::vec_size_utils;
+
+using vsu_ns::ContigHyperparameterSetDefault;
+using vsu_ns::UnaryContigHyperparameterSetEntry;
+
+template <typename argTy> struct AngleContigHyperparameterSet
+{
+    using value_type =
+        typename std::disjunction<ContigHyperparameterSetDefault<4u, 2u>>;
+
+    constexpr static auto vec_sz = value_type::vec_sz;
+    constexpr static auto n_vecs = value_type::n_vecs;
+};
+
+} // end of anonymous namespace
+
 template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
 class angle_contig_kernel;
 
@@ -109,9 +130,12 @@ sycl::event angle_contig_impl(sycl::queue &exec_q,
                               char *res_p,
                               const std::vector<sycl::event> &depends = {})
 {
+    constexpr unsigned int vec_sz = AngleContigHyperparameterSet<argTy>::vec_sz;
+    constexpr unsigned int n_vec = AngleContigHyperparameterSet<argTy>::n_vecs;
+
     return elementwise_common::unary_contig_impl<
-        argTy, AngleOutputType, AngleContigFunctor, angle_contig_kernel>(
-        exec_q, nelems, arg_p, res_p, depends);
+        argTy, AngleOutputType, AngleContigFunctor, angle_contig_kernel, vec_sz,
+        n_vec>(exec_q, nelems, arg_p, res_p, depends);
 }
 
 template <typename fnT, typename T> struct AngleContigFactory