Add bool->bool loops to dpt.sum and dpt.prod

This is done to fix edge cases where the input type is not bool and the output type is bool, which ends up falling back on loops in the input data type, which are cast to bool

Leads to incorrect results in edge cases, i.e.,
```
import dpctl.tensor as dpt

a = dpt.asarray([-1, 1], dtype=dpt.int32)
dpt.sum(a, dtype=dpt.bool)
# usm_ndarray(False)
```

Author: ndgrigorian

dpctl/tensor/libtensor/source/reductions/prod.cpp

@@ -120,6 +120,7 @@ struct TypePairSupportDataForProductReductionTemps
 {
 
     static constexpr bool is_defined = std::disjunction<
+        td_ns::TypePairDefinedEntry<argTy, bool, outTy, bool>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int8_t>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::uint8_t>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int16_t>,
@@ -224,7 +225,7 @@ struct TypePairSupportDataForProductReductionTemps
                                     outTy,
                                     std::complex<double>>,
 
-        // fall-throug
+        // fall-through
         td_ns::NotDefinedEntry>::is_defined;
 };
 
@@ -255,7 +256,9 @@ struct ProductOverAxisTempsStridedFactory
         if constexpr (TypePairSupportDataForProductReductionTemps<
                           srcTy, dstTy>::is_defined)
         {
-            using ReductionOpT = sycl::multiplies<dstTy>;
+            using ReductionOpT = std::conditional_t<std::is_same_v<dstTy, bool>,
+                                                    sycl::logical_and<dstTy>,
+                                                    sycl::multiplies<dstTy>>;
             return dpctl::tensor::kernels::
                 reduction_over_group_temps_strided_impl<srcTy, dstTy,
                                                         ReductionOpT>;
@@ -312,7 +315,9 @@ struct ProductOverAxis1TempsContigFactory
         if constexpr (TypePairSupportDataForProductReductionTemps<
                           srcTy, dstTy>::is_defined)
         {
-            using ReductionOpT = sycl::multiplies<dstTy>;
+            using ReductionOpT = std::conditional_t<std::is_same_v<dstTy, bool>,
+                                                    sycl::logical_and<dstTy>,
+                                                    sycl::multiplies<dstTy>>;
             return dpctl::tensor::kernels::
                 reduction_axis1_over_group_temps_contig_impl<srcTy, dstTy,
                                                              ReductionOpT>;
@@ -331,7 +336,9 @@ struct ProductOverAxis0TempsContigFactory
         if constexpr (TypePairSupportDataForProductReductionTemps<
                           srcTy, dstTy>::is_defined)
         {
-            using ReductionOpT = sycl::multiplies<dstTy>;
+            using ReductionOpT = std::conditional_t<std::is_same_v<dstTy, bool>,
+                                                    sycl::logical_and<dstTy>,
+                                                    sycl::multiplies<dstTy>>;
             return dpctl::tensor::kernels::
                 reduction_axis0_over_group_temps_contig_impl<srcTy, dstTy,
                                                              ReductionOpT>;

dpctl/tensor/libtensor/source/reductions/sum.cpp

@@ -120,6 +120,7 @@ struct TypePairSupportDataForSumReductionTemps
 {
 
     static constexpr bool is_defined = std::disjunction<
+        td_ns::TypePairDefinedEntry<argTy, bool, outTy, bool>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int8_t>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::uint8_t>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int16_t>,
@@ -224,7 +225,7 @@ struct TypePairSupportDataForSumReductionTemps
                                     outTy,
                                     std::complex<double>>,
 
-        // fall-throug
+        // fall-through
         td_ns::NotDefinedEntry>::is_defined;
 };
 
@@ -255,7 +256,9 @@ struct SumOverAxisTempsStridedFactory
         if constexpr (TypePairSupportDataForSumReductionTemps<
                           srcTy, dstTy>::is_defined)
         {
-            using ReductionOpT = sycl::plus<dstTy>;
+            using ReductionOpT =
+                std::conditional_t<std::is_same_v<dstTy, bool>,
+                                   sycl::logical_or<dstTy>, sycl::plus<dstTy>>;
             return dpctl::tensor::kernels::
                 reduction_over_group_temps_strided_impl<srcTy, dstTy,
                                                         ReductionOpT>;
@@ -312,7 +315,9 @@ struct SumOverAxis1TempsContigFactory
         if constexpr (TypePairSupportDataForSumReductionTemps<
                           srcTy, dstTy>::is_defined)
         {
-            using ReductionOpT = sycl::plus<dstTy>;
+            using ReductionOpT =
+                std::conditional_t<std::is_same_v<dstTy, bool>,
+                                   sycl::logical_or<dstTy>, sycl::plus<dstTy>>;
             return dpctl::tensor::kernels::
                 reduction_axis1_over_group_temps_contig_impl<srcTy, dstTy,
                                                              ReductionOpT>;
@@ -331,7 +336,9 @@ struct SumOverAxis0TempsContigFactory
         if constexpr (TypePairSupportDataForSumReductionTemps<
                           srcTy, dstTy>::is_defined)
         {
-            using ReductionOpT = sycl::plus<dstTy>;
+            using ReductionOpT =
+                std::conditional_t<std::is_same_v<dstTy, bool>,
+                                   sycl::logical_or<dstTy>, sycl::plus<dstTy>>;
             return dpctl::tensor::kernels::
                 reduction_axis0_over_group_temps_contig_impl<srcTy, dstTy,
                                                              ReductionOpT>;