Skip to content

Fix incorrect results from cumulative_sum and cumulative_prod with dtype=bool #2018

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 6 commits into from
Feb 28, 2025
Merged

Fix incorrect results from cumulative_sum and cumulative_prod with dtype=bool #2018

Show file tree
Hide file tree
Changes from 5 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
e2d2b3f
Add missing kernels to arithmetic accumulators
ndgrigorian Feb 25, 2025
47239cd
Add tests based on gh-2017
ndgrigorian Feb 25, 2025
48b0e66
Add comment specifying need for work-around in accumulators
ndgrigorian Feb 26, 2025
c58d775
Reuse SYCL utils in workaround for logical operators in accumulators
ndgrigorian Feb 26, 2025
35d46d3
Alias conditional binary op type for cumulative_sum and cumulative_prod
ndgrigorian Feb 26, 2025
fa525a7
Add comment to needs_workaround in accumulators.hpp
ndgrigorian Feb 28, 2025
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
14 changes: 11 additions & 3 deletions dpctl/tensor/libtensor/include/kernels/accumulators.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,8 @@ namespace kernels
namespace accumulators
{

namespace su_ns = dpctl::tensor::sycl_utils;

using dpctl::tensor::ssize_t;
using namespace dpctl::tensor::offset_utils;

Expand Down Expand Up @@ -84,9 +86,17 @@ template <typename srcTy, typename dstTy> struct CastTransformer
}
};

template <typename ScanOpT, typename T> struct needs_workaround
{
// work-around needed due to crash in JITing on CPU
static constexpr bool value = su_ns::IsSyclLogicalAnd<T, ScanOpT>::value ||
su_ns::IsSyclLogicalOr<T, ScanOpT>::value;
};

template <typename BinOpT, typename T> struct can_use_inclusive_scan_over_group
{
static constexpr bool value = sycl::has_known_identity<BinOpT, T>::value;
static constexpr bool value = sycl::has_known_identity<BinOpT, T>::value &&
!needs_workaround<BinOpT, T>::value;
};

namespace detail
Expand Down Expand Up @@ -144,8 +154,6 @@ template <typename T> class stack_strided_t

// Iterative cumulative summation

namespace su_ns = dpctl::tensor::sycl_utils;

using nwiT = std::uint32_t;

template <typename inputT,
Expand Down
15 changes: 11 additions & 4 deletions dpctl/tensor/libtensor/source/accumulators/cumulative_prod.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,10 +70,12 @@ template <typename argTy, typename outTy>
struct TypePairSupportDataForProdAccumulation
{
static constexpr bool is_defined = std::disjunction<
td_ns::TypePairDefinedEntry<argTy, bool, outTy, bool>,
td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int32_t>,
td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int64_t>,

// input int8_t
td_ns::TypePairDefinedEntry<argTy, std::int8_t, outTy, std::int8_t>,
td_ns::TypePairDefinedEntry<argTy, std::int8_t, outTy, std::int32_t>,
td_ns::TypePairDefinedEntry<argTy, std::int8_t, outTy, std::int64_t>,

Expand Down Expand Up @@ -130,6 +132,11 @@ struct TypePairSupportDataForProdAccumulation
td_ns::NotDefinedEntry>::is_defined;
};

template <typename T>
using CumProdScanOpT = std::conditional_t<std::is_same_v<T, bool>,
sycl::logical_and<T>,
sycl::multiplies<T>>;

template <typename fnT, typename srcTy, typename dstTy>
struct CumProd1DContigFactory
{
Expand All @@ -138,7 +145,7 @@ struct CumProd1DContigFactory
if constexpr (TypePairSupportDataForProdAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::multiplies<dstTy>;
using ScanOpT = CumProdScanOpT<dstTy>;
constexpr bool include_initial = false;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down Expand Up @@ -171,7 +178,7 @@ struct CumProd1DIncludeInitialContigFactory
if constexpr (TypePairSupportDataForProdAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::multiplies<dstTy>;
using ScanOpT = CumProdScanOpT<dstTy>;
constexpr bool include_initial = true;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down Expand Up @@ -204,7 +211,7 @@ struct CumProdStridedFactory
if constexpr (TypePairSupportDataForProdAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::multiplies<dstTy>;
using ScanOpT = CumProdScanOpT<dstTy>;
constexpr bool include_initial = false;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down Expand Up @@ -237,7 +244,7 @@ struct CumProdIncludeInitialStridedFactory
if constexpr (TypePairSupportDataForProdAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::multiplies<dstTy>;
using ScanOpT = CumProdScanOpT<dstTy>;
constexpr bool include_initial = true;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down
14 changes: 10 additions & 4 deletions dpctl/tensor/libtensor/source/accumulators/cumulative_sum.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -70,10 +70,12 @@ template <typename argTy, typename outTy>
struct TypePairSupportDataForSumAccumulation
{
static constexpr bool is_defined = std::disjunction<
td_ns::TypePairDefinedEntry<argTy, bool, outTy, bool>,
td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int32_t>,
td_ns::TypePairDefinedEntry<argTy, bool, outTy, std::int64_t>,

// input int8_t
td_ns::TypePairDefinedEntry<argTy, std::int8_t, outTy, std::int8_t>,
td_ns::TypePairDefinedEntry<argTy, std::int8_t, outTy, std::int32_t>,
td_ns::TypePairDefinedEntry<argTy, std::int8_t, outTy, std::int64_t>,

Expand Down Expand Up @@ -130,6 +132,10 @@ struct TypePairSupportDataForSumAccumulation
td_ns::NotDefinedEntry>::is_defined;
};

template <typename T>
using CumSumScanOpT = std::
conditional_t<std::is_same_v<T, bool>, sycl::logical_or<T>, sycl::plus<T>>;

template <typename fnT, typename srcTy, typename dstTy>
struct CumSum1DContigFactory
{
Expand All @@ -138,7 +144,7 @@ struct CumSum1DContigFactory
if constexpr (TypePairSupportDataForSumAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::plus<dstTy>;
using ScanOpT = CumSumScanOpT<dstTy>;
constexpr bool include_initial = false;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down Expand Up @@ -171,7 +177,7 @@ struct CumSum1DIncludeInitialContigFactory
if constexpr (TypePairSupportDataForSumAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::plus<dstTy>;
using ScanOpT = CumSumScanOpT<dstTy>;
constexpr bool include_initial = true;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down Expand Up @@ -204,7 +210,7 @@ struct CumSumStridedFactory
if constexpr (TypePairSupportDataForSumAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::plus<dstTy>;
using ScanOpT = CumSumScanOpT<dstTy>;
constexpr bool include_initial = false;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down Expand Up @@ -237,7 +243,7 @@ struct CumSumIncludeInitialStridedFactory
if constexpr (TypePairSupportDataForSumAccumulation<srcTy,
dstTy>::is_defined)
{
using ScanOpT = sycl::plus<dstTy>;
using ScanOpT = CumSumScanOpT<dstTy>;
constexpr bool include_initial = true;
if constexpr (std::is_same_v<srcTy, dstTy>) {
using dpctl::tensor::kernels::accumulators::NoOpTransformer;
Expand Down
12 changes: 12 additions & 0 deletions dpctl/tests/test_tensor_accumulation.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -421,3 +421,15 @@ def test_cumulative_sum_gh_1901(p):
inp = dpt.ones(n, dtype=dt)
r = dpt.cumulative_sum(inp, dtype=dt)
assert dpt.all(r == dpt.arange(1, n + 1, dtype=dt))


@pytest.mark.parametrize(
"dt", ["i1", "i2", "i4", "i8", "f2", "f4", "f8", "c8", "c16"]
)
def test_gh_2017(dt):
"See https://github.com/IntelPython/dpctl/issues/2017"
q = get_queue_or_skip()
skip_if_dtype_not_supported(dt, q)
x = dpt.asarray([-1, 1], dtype=dpt.dtype(dt), sycl_queue=q)
r = dpt.cumulative_sum(x, dtype="?")
assert dpt.all(r)
Loading