Merge pull request #1802 from IntelPython/clean-up-reduction-headers

oleksandr-pavlyk · web-flow · commit 63ffaba70ac5 · 2024-08-19T11:16:51.000-05:00
Explicitly include headers used in the c++ file pertaining to reduction
diff --git a/dpctl/tensor/libtensor/include/kernels/reductions.hpp b/dpctl/tensor/libtensor/include/kernels/reductions.hpp
@@ -39,16 +39,16 @@
 #include "utils/type_dispatch_building.hpp"
 #include "utils/type_utils.hpp"
 
-namespace td_ns = dpctl::tensor::type_dispatch;
-namespace su_ns = dpctl::tensor::sycl_utils;
-
 namespace dpctl
 {
 namespace tensor
 {
 namespace kernels
 {
 
+namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
+
 namespace reduction_detail
 {
 
diff --git a/dpctl/tensor/libtensor/include/utils/math_utils.hpp b/dpctl/tensor/libtensor/include/utils/math_utils.hpp
@@ -122,17 +122,32 @@ template <typename T> T logaddexp(T x, T y)
         return x + log2;
     }
     else {
-        // FIXME: switch to `sycl::log1p` when
-        // compiler segfault in CUDA build is fixed
         const T tmp = x - y;
-        if (tmp > 0) {
-            return x + std::log1p(sycl::exp(-tmp));
-        }
-        else if (tmp <= 0) {
-            return y + std::log1p(sycl::exp(tmp));
+        constexpr T zero(0);
+
+        if constexpr (std::is_same_v<T, sycl::half>) {
+            return (tmp > zero)
+                       ? (x + sycl::log1p(sycl::exp(-tmp)))
+                       : ((tmp <= zero) ? y + sycl::log1p(sycl::exp(tmp))
+                                        : std::numeric_limits<T>::quiet_NaN());
         }
         else {
-            return std::numeric_limits<T>::quiet_NaN();
+            if constexpr (std::is_same_v<T, double>) {
+                // FIXME: switch to `sycl::log1p` when
+                // compiler segfault in CUDA build is fixed
+                return (tmp > zero)
+                           ? (x + std::log1p(sycl::exp(-tmp)))
+                           : ((tmp <= zero)
+                                  ? y + std::log1p(sycl::exp(tmp))
+                                  : std::numeric_limits<T>::quiet_NaN());
+            }
+            else {
+                return (tmp > zero)
+                           ? (x + sycl::log1p(sycl::exp(-tmp)))
+                           : ((tmp <= zero)
+                                  ? y + sycl::log1p(sycl::exp(tmp))
+                                  : std::numeric_limits<T>::quiet_NaN());
+            }
         }
     }
 }
diff --git a/dpctl/tensor/libtensor/source/linalg_functions/dot_dispatch.hpp b/dpctl/tensor/libtensor/source/linalg_functions/dot_dispatch.hpp
@@ -30,6 +30,7 @@
 
 #include "kernels/linalg_functions/dot_product.hpp"
 #include "kernels/linalg_functions/gemm.hpp"
+#include "utils/type_dispatch_building.hpp"
 
 namespace dpctl
 {
@@ -38,6 +39,8 @@ namespace tensor
 namespace py_internal
 {
 
+namespace td_ns = dpctl::tensor::type_dispatch;
+
 template <typename T1, typename T2> struct DotAtomicOutputType
 {
     using value_type = typename std::disjunction< // disjunction is C++17
diff --git a/dpctl/tensor/libtensor/source/reductions/argmax.cpp b/dpctl/tensor/libtensor/source/reductions/argmax.cpp
@@ -32,6 +32,7 @@
 
 #include "kernels/reductions.hpp"
 #include "reduction_over_axis.hpp"
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
 namespace py = pybind11;
@@ -44,6 +45,7 @@ namespace py_internal
 {
 
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
 
 namespace impl
 {
diff --git a/dpctl/tensor/libtensor/source/reductions/argmin.cpp b/dpctl/tensor/libtensor/source/reductions/argmin.cpp
@@ -32,6 +32,8 @@
 
 #include "kernels/reductions.hpp"
 #include "reduction_over_axis.hpp"
+
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
 namespace py = pybind11;
@@ -44,6 +46,7 @@ namespace py_internal
 {
 
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
 
 namespace impl
 {
diff --git a/dpctl/tensor/libtensor/source/reductions/logsumexp.cpp b/dpctl/tensor/libtensor/source/reductions/logsumexp.cpp
@@ -32,6 +32,7 @@
 
 #include "kernels/reductions.hpp"
 #include "reduction_over_axis.hpp"
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
 namespace py = pybind11;
@@ -44,6 +45,7 @@ namespace py_internal
 {
 
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
 
 namespace impl
 {
@@ -68,6 +70,7 @@ struct TypePairSupportDataForLogSumExpReductionTemps
     static constexpr bool is_defined = std::disjunction< // disjunction is C++17
                                                          // feature, supported
                                                          // by DPC++ input bool
+#if 1
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, sycl::half>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, float>,
         td_ns::TypePairDefinedEntry<argTy, bool, outTy, double>,
@@ -105,7 +108,6 @@ struct TypePairSupportDataForLogSumExpReductionTemps
         // input uint64_t
         td_ns::TypePairDefinedEntry<argTy, std::uint64_t, outTy, float>,
         td_ns::TypePairDefinedEntry<argTy, std::uint64_t, outTy, double>,
-
         // input half
         td_ns::TypePairDefinedEntry<argTy, sycl::half, outTy, sycl::half>,
         td_ns::TypePairDefinedEntry<argTy, sycl::half, outTy, float>,
@@ -117,6 +119,7 @@ struct TypePairSupportDataForLogSumExpReductionTemps
 
         // input double
         td_ns::TypePairDefinedEntry<argTy, double, outTy, double>,
+#endif
 
         // fall-through
         td_ns::NotDefinedEntry>::is_defined;
diff --git a/dpctl/tensor/libtensor/source/reductions/max.cpp b/dpctl/tensor/libtensor/source/reductions/max.cpp
@@ -31,6 +31,7 @@
 #include <vector>
 
 #include "kernels/reductions.hpp"
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
 #include "reduction_atomic_support.hpp"
@@ -46,6 +47,7 @@ namespace py_internal
 {
 
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
 
 namespace impl
 {
diff --git a/dpctl/tensor/libtensor/source/reductions/min.cpp b/dpctl/tensor/libtensor/source/reductions/min.cpp
@@ -31,6 +31,7 @@
 #include <vector>
 
 #include "kernels/reductions.hpp"
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
 #include "reduction_atomic_support.hpp"
@@ -46,6 +47,7 @@ namespace py_internal
 {
 
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
 
 namespace impl
 {
diff --git a/dpctl/tensor/libtensor/source/reductions/reduce_hypot.cpp b/dpctl/tensor/libtensor/source/reductions/reduce_hypot.cpp
@@ -32,6 +32,7 @@
 
 #include "kernels/reductions.hpp"
 #include "reduction_over_axis.hpp"
+#include "utils/sycl_utils.hpp"
 #include "utils/type_dispatch_building.hpp"
 
 namespace py = pybind11;
@@ -44,6 +45,7 @@ namespace py_internal
 {
 
 namespace td_ns = dpctl::tensor::type_dispatch;
+namespace su_ns = dpctl::tensor::sycl_utils;
 
 namespace impl
 {
diff --git a/dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp b/dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp
@@ -52,6 +52,8 @@ namespace tensor
 namespace py_internal
 {
 
+namespace td_ns = dpctl::tensor::type_dispatch;
+
 /* ====================== dtype supported ======================== */
 
 /*! @brief Template implementing Python API for querying type support by

Original file line number	Diff line number	Diff line change
`@@ -39,16 +39,16 @@`
`39`	`39`	`#include "utils/type_dispatch_building.hpp"`
`40`	`40`	`#include "utils/type_utils.hpp"`
`41`	`41`
`42`		`-namespace td_ns = dpctl::tensor::type_dispatch;`
`43`		`-namespace su_ns = dpctl::tensor::sycl_utils;`
`44`		`-`
`45`	`42`	`namespace dpctl`
`46`	`43`	`{`
`47`	`44`	`namespace tensor`
`48`	`45`	`{`
`49`	`46`	`namespace kernels`
`50`	`47`	`{`
`51`	`48`
	`49`	`+namespace td_ns = dpctl::tensor::type_dispatch;`
	`50`	`+namespace su_ns = dpctl::tensor::sycl_utils;`
	`51`	`+`
`52`	`52`	`namespace reduction_detail`
`53`	`53`	`{`
`54`	`54`
Original file line number	Diff line number	Diff line change
`@@ -30,6 +30,7 @@`
`30`	`30`
`31`	`31`	`#include "kernels/linalg_functions/dot_product.hpp"`
`32`	`32`	`#include "kernels/linalg_functions/gemm.hpp"`
	`33`	`+#include "utils/type_dispatch_building.hpp"`
`33`	`34`
`34`	`35`	`namespace dpctl`
`35`	`36`	`{`
`@@ -38,6 +39,8 @@ namespace tensor`
`38`	`39`	`namespace py_internal`
`39`	`40`	`{`
`40`	`41`
	`42`	`+namespace td_ns = dpctl::tensor::type_dispatch;`
	`43`	`+`
`41`	`44`	`template <typename T1, typename T2> struct DotAtomicOutputType`
`42`	`45`	`{`
`43`	`46`	`using value_type = typename std::disjunction< // disjunction is C++17`
Original file line number	Diff line number	Diff line change
`@@ -32,6 +32,7 @@`
`32`	`32`
`33`	`33`	`#include "kernels/reductions.hpp"`
`34`	`34`	`#include "reduction_over_axis.hpp"`
	`35`	`+#include "utils/sycl_utils.hpp"`
`35`	`36`	`#include "utils/type_dispatch_building.hpp"`
`36`	`37`
`37`	`38`	`namespace py = pybind11;`
`@@ -44,6 +45,7 @@ namespace py_internal`
`44`	`45`	`{`
`45`	`46`
`46`	`47`	`namespace td_ns = dpctl::tensor::type_dispatch;`
	`48`	`+namespace su_ns = dpctl::tensor::sycl_utils;`
`47`	`49`
`48`	`50`	`namespace impl`
`49`	`51`	`{`
Original file line number	Diff line number	Diff line change
`@@ -32,6 +32,8 @@`
`32`	`32`
`33`	`33`	`#include "kernels/reductions.hpp"`
`34`	`34`	`#include "reduction_over_axis.hpp"`
	`35`	`+`
	`36`	`+#include "utils/sycl_utils.hpp"`
`35`	`37`	`#include "utils/type_dispatch_building.hpp"`
`36`	`38`
`37`	`39`	`namespace py = pybind11;`
`@@ -44,6 +46,7 @@ namespace py_internal`
`44`	`46`	`{`
`45`	`47`
`46`	`48`	`namespace td_ns = dpctl::tensor::type_dispatch;`
	`49`	`+namespace su_ns = dpctl::tensor::sycl_utils;`
`47`	`50`
`48`	`51`	`namespace impl`
`49`	`52`	`{`
Original file line number	Diff line number	Diff line change
`@@ -31,6 +31,7 @@`
`31`	`31`	`#include <vector>`
`32`	`32`
`33`	`33`	`#include "kernels/reductions.hpp"`
	`34`	`+#include "utils/sycl_utils.hpp"`
`34`	`35`	`#include "utils/type_dispatch_building.hpp"`
`35`	`36`
`36`	`37`	`#include "reduction_atomic_support.hpp"`
`@@ -46,6 +47,7 @@ namespace py_internal`
`46`	`47`	`{`
`47`	`48`
`48`	`49`	`namespace td_ns = dpctl::tensor::type_dispatch;`
	`50`	`+namespace su_ns = dpctl::tensor::sycl_utils;`
`49`	`51`
`50`	`52`	`namespace impl`
`51`	`53`	`{`
Original file line number	Diff line number	Diff line change
`@@ -52,6 +52,8 @@ namespace tensor`
`52`	`52`	`namespace py_internal`
`53`	`53`	`{`
`54`	`54`
	`55`	`+namespace td_ns = dpctl::tensor::type_dispatch;`
	`56`	`+`
`55`	`57`	`/* ====================== dtype supported ======================== */`
`56`	`58`
`57`	`59`	`/*! @brief Template implementing Python API for querying type support by`