Skip to content

Explicitly include headers used in the c++ file pertaining to reduction #1802

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Aug 19, 2024
Merged

Explicitly include headers used in the c++ file pertaining to reduction #1802

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions dpctl/tensor/libtensor/include/kernels/reductions.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
{

Expand Down
31 changes: 23 additions & 8 deletions dpctl/tensor/libtensor/include/utils/math_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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());
}
}
}
}
Expand Down
3 changes: 3 additions & 0 deletions dpctl/tensor/libtensor/source/linalg_functions/dot_dispatch.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
{
Expand All @@ -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
Expand Down
2 changes: 2 additions & 0 deletions dpctl/tensor/libtensor/source/reductions/argmax.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -44,6 +45,7 @@ namespace py_internal
{

namespace td_ns = dpctl::tensor::type_dispatch;
namespace su_ns = dpctl::tensor::sycl_utils;

namespace impl
{
Expand Down
3 changes: 3 additions & 0 deletions dpctl/tensor/libtensor/source/reductions/argmin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -44,6 +46,7 @@ namespace py_internal
{

namespace td_ns = dpctl::tensor::type_dispatch;
namespace su_ns = dpctl::tensor::sycl_utils;

namespace impl
{
Expand Down
5 changes: 4 additions & 1 deletion dpctl/tensor/libtensor/source/reductions/logsumexp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -44,6 +45,7 @@ namespace py_internal
{

namespace td_ns = dpctl::tensor::type_dispatch;
namespace su_ns = dpctl::tensor::sycl_utils;

namespace impl
{
Expand All @@ -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>,
Expand Down Expand Up @@ -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>,
Expand All @@ -117,6 +119,7 @@ struct TypePairSupportDataForLogSumExpReductionTemps

// input double
td_ns::TypePairDefinedEntry<argTy, double, outTy, double>,
#endif

// fall-through
td_ns::NotDefinedEntry>::is_defined;
Expand Down
2 changes: 2 additions & 0 deletions dpctl/tensor/libtensor/source/reductions/max.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -46,6 +47,7 @@ namespace py_internal
{

namespace td_ns = dpctl::tensor::type_dispatch;
namespace su_ns = dpctl::tensor::sycl_utils;

namespace impl
{
Expand Down
2 changes: 2 additions & 0 deletions dpctl/tensor/libtensor/source/reductions/min.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -46,6 +47,7 @@ namespace py_internal
{

namespace td_ns = dpctl::tensor::type_dispatch;
namespace su_ns = dpctl::tensor::sycl_utils;

namespace impl
{
Expand Down
2 changes: 2 additions & 0 deletions dpctl/tensor/libtensor/source/reductions/reduce_hypot.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -44,6 +45,7 @@ namespace py_internal
{

namespace td_ns = dpctl::tensor::type_dispatch;
namespace su_ns = dpctl::tensor::sycl_utils;

namespace impl
{
Expand Down
2 changes: 2 additions & 0 deletions dpctl/tensor/libtensor/source/reductions/reduction_over_axis.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
Loading