12076405164

Committed 28 Nov 2024 11:09PM UTC coverage: 66.699% (-1.2%) from 67.852%

Build # 12076405164

Build Type

Pull #2180

github

Committed by

web-flow

Commit Message

Merge fa3af10fd into 4e021720c

Pull Request Pull Request #2180: Implementation of correlate

Run Details

4427 of 10378 branches covered (42.66%)

Branch coverage included in aggregate %.

282 of 590 new or added lines in 9 files covered. (47.8%)

58 existing lines in 3 files now uncovered.

16622 of 21180 relevant lines covered (78.48%)

20147.7 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

3.5

/dpnp/backend/extensions/statistics/sliding_window1d.hpp

//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// - Redistributions of source code must retain the above copyright notice,
//   this list of conditions and the following disclaimer.
// - Redistributions in binary form must reproduce the above copyright notice,
//   this list of conditions and the following disclaimer in the documentation
//   and/or other materials provided with the distribution.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
// THE POSSIBILITY OF SUCH DAMAGE.
//*****************************************************************************

#pragma once

#include "utils/math_utils.hpp"
#include <complex>
#include <sycl/sycl.hpp>
#include <tuple>
#include <type_traits>

#include <stdio.h>

#include "common.hpp"

namespace dpctl
{
namespace tensor
{
class usm_ndarray;
}
} // namespace dpctl

using dpctl::tensor::usm_ndarray;

namespace statistics
{
using common::Align;
using common::CeilDiv;

namespace sliding_window1d
{

template <typename T, uint32_t Size>
class _RegistryDataStorage
{
public:
    using ncT = typename std::remove_const_t<T>;
    using SizeT = decltype(Size);
    static constexpr SizeT _size = Size;

    _RegistryDataStorage(const sycl::nd_item<1> &item)
        : sbgroup(item.get_sub_group())
    {
    }

    template <typename yT>
    T &operator[](const yT &idx)
    {
        static_assert(std::is_integral_v<yT>,
                      "idx must be of an integral type");
        return data[idx];
    }

    template <typename yT>
    const T &operator[](const yT &idx) const
    {
        static_assert(std::is_integral_v<yT>,
                      "idx must be of an integral type");
        return data[idx];
    }

    T &value()
    {
        static_assert(Size == 1,
                      "Size is not equal to 1. Use value(idx) instead");
        return data[0];
    }

    const T &value() const
    {
        static_assert(Size == 1,
                      "Size is not equal to 1. Use value(idx) instead");
        return data[0];
    }

    template <typename yT, typename xT>
    T broadcast(const yT &y, const xT &x) const
    {
        static_assert(std::is_integral_v<std::remove_reference_t<yT>>,
                      "y must be of an integral type");
        static_assert(std::is_integral_v<std::remove_reference_t<xT>>,
                      "x must be of an integral type");

        return sycl::select_from_group(sbgroup, data[y], x);
    }

    template <typename iT>
    T broadcast(const iT &idx) const
    {
        if constexpr (Size == 1) {
            return broadcast(0, idx);
        }
        else {
            return broadcast(idx / size_x(), idx % size_x());
        }
    }

    template <typename yT, typename xT>
    T shift_left(const yT &y, const xT &x) const
    {
        static_assert(std::is_integral_v<yT>, "y must be of an integral type");
        static_assert(std::is_integral_v<xT>, "x must be of an integral type");

        return sycl::shift_group_left(sbgroup, data[y], x);
    }

    template <typename yT, typename xT>
    T shift_right(const yT &y, const xT &x) const
    {
        static_assert(std::is_integral_v<yT>, "y must be of an integral type");
        static_assert(std::is_integral_v<xT>, "x must be of an integral type");

        return sycl::shift_group_right(sbgroup, data[y], x);
    }

    constexpr SizeT size_y() const
    {
        return _size;
    }

    SizeT size_x() const
    {
        return sbgroup.get_max_local_range()[0];
    }

    SizeT total_size() const
    {
        return size_x() * size_y();
    }

    ncT *ptr()
    {
        return data;
    }

    SizeT x() const
    {
        return sbgroup.get_local_linear_id();
    }

protected:
    const sycl::sub_group sbgroup;
    ncT data[Size];
};

template <typename T, uint32_t Size = 1>
struct RegistryData : public _RegistryDataStorage<T, Size>
{
    using SizeT = typename _RegistryDataStorage<T, Size>::SizeT;

    using _RegistryDataStorage<T, Size>::_RegistryDataStorage;

    template <typename LaneIdT,
              typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, SizeT>>>
    void fill_lane(const LaneIdT &lane_id, const T &value, Condition &&mask)
    {
        static_assert(std::is_integral_v<LaneIdT>,
                      "lane_id must be of an integral type");
        if (mask(this->x())) {
            this->data[lane_id] = value;
        }
    }

    template <typename LaneIdT>
    void fill_lane(const LaneIdT &lane_id, const T &value, const bool &mask)
    {
        fill_lane(lane_id, value, [mask](auto &&) { return mask; });
    }

    template <typename LaneIdT>
    void fill_lane(const LaneIdT &lane_id, const T &value)
    {
        fill_lane(lane_id, value, true);
    }

    template <typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, SizeT, SizeT>>>
    void fill(const T &value, Condition &&mask)
    {
        for (SizeT i = 0; i < Size; ++i) {
            fill_lane(i, value, mask(i, this->x()));
        }
    }

    void fill(const T &value)
    {
        fill(value, [](auto &&, auto &&) { return true; });
    }

    template <typename LaneIdT,
              typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, const T *const>>>
    T *load_lane(const LaneIdT &lane_id,
                 const T *const data,
                 Condition &&mask,
                 const T &default_v)
    {
        static_assert(std::is_integral_v<LaneIdT>,
                      "lane_id must be of an integral type");
        this->data[lane_id] = mask(data) ? data[0] : default_v;

        return data + this->size_x();
    }

    template <typename LaneIdT>
    T *load_lane(const LaneIdT &laned_id,
                 const T *const data,
                 const bool &mask,
                 const T &default_v)
    {
        return load_lane(
            laned_id, data, [mask](auto &&) { return mask; }, default_v);
    }

    template <typename LaneIdT>
    T *load_lane(const LaneIdT &laned_id, const T *const data)
    {
        constexpr T default_v = 0;
        return load_lane(laned_id, data, true, default_v);
    }

    template <typename yStrideT,
              typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, const T *const>>>
    T *load(const T *const data,
            const yStrideT &y_stride,
            Condition &&mask,
            const T &default_v)
    {
        auto *it = data;
        for (SizeT i = 0; i < Size; ++i) {
            load_lane(i, it, mask, default_v);
            it += y_stride;
        }

        return it;
    }

    template <typename yStrideT>
    T *load(const T *const data,
            const yStrideT &y_stride,
            const bool &mask,
            const T &default_v)
    {
        return load(
            data, y_stride, [mask](auto &&) { return mask; }, default_v);
    }

    template <typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, const T *const>>>
    T *load(const T *const data, Condition &&mask, const T &default_v)
    {
        return load(data, this->size_x(), mask, default_v);
    }

    T *load(const T *const data, const bool &mask, const T &default_v)
    {
        return load(
            data, [mask](auto &&) { return mask; }, default_v);
    }

    T *load(const T *const data)
    {
        constexpr T default_v = 0;
        return load(data, true, default_v);
    }

    template <typename LaneIdT,
              typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, const T *const>>>
    T *store_lane(const LaneIdT &lane_id, T *const data, Condition &&mask)
    {
        static_assert(std::is_integral_v<LaneIdT>,
                      "lane_id must be of an integral type");

        if (mask(data)) {
            data[0] = this->data[lane_id];
        }

        return data + this->size_x();
    }

    template <typename LaneIdT>
    T *store_lane(const LaneIdT &lane_id, T *const data, const bool &mask)
    {
        return store_lane(lane_id, data, [mask](auto &&) { return mask; });
    }

    template <typename LaneIdT>
    T *store_lane(const LaneIdT &lane_id, T *const data)
    {
        return store_lane(lane_id, data, true);
    }

    template <typename yStrideT,
              typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, const T *const>>>
    T *store(T *const data, const yStrideT &y_stride, Condition &&condition)
    {
        auto *it = data;
        for (SizeT i = 0; i < Size; ++i) {
            store_lane(i, it, condition);
            it += y_stride;
        }

        return it;
    }

    template <typename yStrideT>
    T *store(T *const data, const yStrideT &y_stride, const bool &mask)
    {
        return store(data, y_stride, [mask](auto &&) { return mask; });
    }

    template <typename Condition,
              typename = std::enable_if_t<
                  std::is_invocable_r_v<bool, Condition, const T *const>>>
    T *store(T *const data, Condition &&condition)
    {
        return store(data, this->size_x(), condition);
    }

    T *store(T *const data, const bool &mask)
    {
        return store(data, [mask](auto &&) { return mask; });
    }

    T *store(T *const data)
    {
        return store(data, true);
    }
};

template <typename T, uint32_t Size>
struct RegistryWindow : public RegistryData<T, Size>
{
    using SizeT = typename RegistryData<T, Size>::SizeT;

    using RegistryData<T, Size>::RegistryData;

    template <typename shT>
    void advance_left(const shT &shift, const T &fill_value)
    {
        static_assert(std::is_integral_v<shT>,
                      "shift must be of an integral type");

        uint32_t shift_r = this->size_x() - shift;
        for (SizeT i = 0; i < Size; ++i) {
            this->data[i] = this->shift_left(i, shift);
            auto border =
                i < Size - 1 ? this->shift_right(i + 1, shift_r) : fill_value;
            if (this->x() >= shift_r) {
                this->data[i] = border;
            }
        }
    }

    void advance_left(const T &fill_value)
    {
        advance_left(1, fill_value);
    }

    void advance_left()
    {
        constexpr T fill_value = 0;
        advance_left(fill_value);
    }
};

template <typename T, typename SizeT = size_t>
class Span
{
public:
    using value_type = T;
    using size_type = SizeT;

    Span(T *const data, const SizeT size) : data_(data), size_(size) {}

    T *begin() const
    {
        return data();
    }

    T *end() const
    {
        return data() + size();
    }

    SizeT size() const
    {
        return size_;
    }

    T *data() const
    {
        return data_;
    }

protected:
    T *const data_;
    const SizeT size_;
};

template <typename T, typename SizeT = size_t>
Span<T, SizeT> make_span(T *const data, const SizeT size)
{
    return Span<T, SizeT>(data, size);
}

template <typename T, typename SizeT = size_t>
class PaddedSpan : public Span<T, SizeT>
{
public:
    using value_type = T;
    using size_type = SizeT;

    PaddedSpan(T *const data, const SizeT size, const SizeT pad)
        : Span<T>(data, size), pad_(pad)
    {
    }

    T *padded_begin() const
    {
        return this->begin() - pad();
    }

    SizeT pad() const
    {
        return pad_;
    }

protected:
    const SizeT pad_;
};

template <typename T, typename SizeT = size_t>
PaddedSpan<T, SizeT>
    make_padded_span(T *const data, const SizeT size, const SizeT offset)
{
    return PaddedSpan<T, SizeT>(data, size, offset);
}

template <typename Results,
          typename AData,
          typename VData,
          typename Op,
          typename Red>
void process_block(Results &results,
                   uint32_t r_size,
                   AData &a_data,
                   VData &v_data,
                   uint32_t block_size,
                   Op op,
                   Red red)
{
    for (uint32_t i = 0; i < block_size; ++i) {
        auto v_val = v_data.broadcast(i);
        for (uint32_t r = 0; r < r_size; ++r) {
            results[r] = red(results[r], op(a_data[r], v_val));
        }
        a_data.advance_left();
    }
}

template <typename SizeT>
SizeT get_global_linear_id(const uint32_t wpi, const sycl::nd_item<1> &item)
{
    auto sbgroup = item.get_sub_group();
    const auto sg_loc_id = sbgroup.get_local_linear_id();

    const SizeT sg_base_id = wpi * (item.get_global_linear_id() - sg_loc_id);
    const SizeT id = sg_base_id + sg_loc_id;

    return id;
}

template <typename SizeT>
uint32_t get_results_num(const uint32_t wpi,
                         const SizeT size,
                         const SizeT global_id,
                         const sycl::nd_item<1> &item)
{
    auto sbgroup = item.get_sub_group();

    const auto sbg_size = sbgroup.get_max_local_range()[0];
    const auto size_ = sycl::sub_sat(size, global_id);
    return std::min(SizeT(wpi), CeilDiv(size_, sbg_size));
}

template <uint32_t WorkPI,
          typename T,
          typename SizeT,
          typename Op,
          typename Red>
class sliding_window1d_kernel;

template <uint32_t WorkPI,
          typename T,
          typename SizeT,
          typename Op,
          typename Red>
void submit_sliding_window1d(const PaddedSpan<const T, SizeT> &a,
                             const Span<const T, SizeT> &v,
                             const Op &op,
                             const Red &red,
                             Span<T, SizeT> &out,
                             sycl::nd_range<1> nd_range,
                             sycl::handler &cgh)
{
    cgh.parallel_for<sliding_window1d_kernel<WorkPI, T, SizeT, Op, Red>>(
        nd_range, [=](sycl::nd_item<1> item) {
            auto glid = get_global_linear_id<SizeT>(WorkPI, item);

            auto results = RegistryData<T, WorkPI>(item);
            results.fill(0);

            auto results_num = get_results_num(WorkPI, out.size(), glid, item);

            const auto *a_begin = a.begin();
            const auto *a_end = a.end();

            auto sbgroup = item.get_sub_group();

            const auto chunks_count =
                CeilDiv(v.size(), sbgroup.get_max_local_range()[0]);

            const auto *a_ptr = &a.padded_begin()[glid];

            auto _a_load_cond = [a_begin, a_end](auto &&ptr) {
                return ptr >= a_begin && ptr < a_end;
            };

            auto a_data = RegistryWindow<const T, WorkPI + 1>(item);
            a_ptr = a_data.load(a_ptr, _a_load_cond, 0);

            const auto *v_ptr = &v.begin()[sbgroup.get_local_linear_id()];
            auto v_size = v.size();

            for (uint32_t b = 0; b < chunks_count; ++b) {
                auto v_data = RegistryData<const T>(item);
                v_ptr = v_data.load(v_ptr, v_data.x() < v_size, 0);

                uint32_t chunk_size_ =
                    std::min(v_size, SizeT(v_data.total_size()));
                process_block(results, results_num, a_data, v_data, chunk_size_,
                              op, red);

                if (b != chunks_count - 1) {
                    a_ptr = a_data.load_lane(a_data.size_y() - 1, a_ptr,
                                             _a_load_cond, 0);
                    v_size -= v_data.total_size();
                }
            }

            auto *const out_ptr = out.begin();
            auto *const out_end = out.end();
            results.store(&out_ptr[glid],
                          [out_end](auto &&ptr) { return ptr < out_end; });
        });
}

template <uint32_t WorkPI,
          typename T,
          typename SizeT,
          typename Op,
          typename Red>
class sliding_window1d_small_kernel;

template <uint32_t WorkPI,
          typename T,
          typename SizeT,
          typename Op,
          typename Red>
void submit_sliding_window1d_small_kernel(const PaddedSpan<const T, SizeT> &a,
                                          const Span<const T, SizeT> &v,
                                          const Op &op,
                                          const Red &red,
                                          Span<T, SizeT> &out,
                                          sycl::nd_range<1> nd_range,
                                          sycl::handler &cgh)
{
    cgh.parallel_for<sliding_window1d_small_kernel<WorkPI, T, SizeT, Op, Red>>(
        nd_range, [=](sycl::nd_item<1> item) {
            auto glid = get_global_linear_id<SizeT>(WorkPI, item);

            auto results = RegistryData<T, WorkPI>(item);
            results.fill(0);

            auto sbgroup = item.get_sub_group();
            auto sg_size = sbgroup.get_max_local_range()[0];

            const uint32_t to_read = WorkPI * sg_size + v.size();
            const auto *a_begin = a.begin();

            const auto *a_ptr = &a.padded_begin()[glid];
            const auto *a_end = std::min(a_ptr + to_read, a.end());

            auto _a_load_cond = [a_begin, a_end](auto &&ptr) {
                return ptr >= a_begin && ptr < a_end;
            };

            auto a_data = RegistryWindow<const T, WorkPI + 1>(item);
            a_data.load(a_ptr, _a_load_cond, 0);

            const auto *v_ptr = &v.begin()[sbgroup.get_local_linear_id()];
            auto v_size = v.size();

            auto v_data = RegistryData<const T>(item);
            v_ptr = v_data.load(v_ptr, v_data.x() < v_size, 0);

            auto results_num = get_results_num(WorkPI, out.size(), glid, item);

            process_block(results, results_num, a_data, v_data, v_size, op,
                          red);

            auto *const out_ptr = out.begin();
            auto *const out_end = out.end();
            results.store(&out_ptr[glid],
                          [out_end](auto &&ptr) { return ptr < out_end; });
        });
}

void validate(const usm_ndarray &a,
              const usm_ndarray &v,
              const usm_ndarray &out,
              const size_t l_pad,
              const size_t r_pad);
} // namespace sliding_window1d
} // namespace statistics

1	//*****************************************************************************
2	// Copyright (c) 2024, Intel Corporation
3	// All rights reserved.
4	//
5	// Redistribution and use in source and binary forms, with or without
6	// modification, are permitted provided that the following conditions are met:
7	// - Redistributions of source code must retain the above copyright notice,
8	// this list of conditions and the following disclaimer.
9	// - Redistributions in binary form must reproduce the above copyright notice,
10	// this list of conditions and the following disclaimer in the documentation
11	// and/or other materials provided with the distribution.
12	//
13	// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14	// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15	// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16	// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17	// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18	// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19	// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20	// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21	// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22	// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23	// THE POSSIBILITY OF SUCH DAMAGE.
24	//*****************************************************************************
25
26	#pragma once
27
28	#include "utils/math_utils.hpp"
29	#include <complex>
30	#include <sycl/sycl.hpp>
31	#include <tuple>
32	#include <type_traits>
33
34	#include <stdio.h>
35
36	#include "common.hpp"
37
38	namespace dpctl
39	{
40	namespace tensor
41	{
42	class usm_ndarray;
43	}
44	} // namespace dpctl
45
46	using dpctl::tensor::usm_ndarray;
47
48	namespace statistics
49	{
50	using common::Align;
51	using common::CeilDiv;
52
53	namespace sliding_window1d
54	{
55
56	template <typename T, uint32_t Size>
57	class _RegistryDataStorage
58	{
59	public:
60	using ncT = typename std::remove_const_t<T>;
61	using SizeT = decltype(Size);
62	static constexpr SizeT _size = Size;
63
64	_RegistryDataStorage(const sycl::nd_item<1> &item)
NEW 65	: sbgroup(item.get_sub_group())	×
NEW 66	{	×
NEW 67	}	×
68
69	template <typename yT>
70	T &operator[](const yT &idx)
NEW 71	{	×
NEW 72	static_assert(std::is_integral_v<yT>,	×
NEW 73	"idx must be of an integral type");	×
NEW 74	return data[idx];	×
NEW 75	}	×
76
77	template <typename yT>
78	const T &operator[](const yT &idx) const
79	{
80	static_assert(std::is_integral_v<yT>,
81	"idx must be of an integral type");
82	return data[idx];
83	}
84
85	T &value()
86	{
87	static_assert(Size == 1,
88	"Size is not equal to 1. Use value(idx) instead");
89	return data[0];
90	}
91
92	const T &value() const
93	{
94	static_assert(Size == 1,
95	"Size is not equal to 1. Use value(idx) instead");
96	return data[0];
97	}
98
99	template <typename yT, typename xT>
100	T broadcast(const yT &y, const xT &x) const
NEW 101	{	×
NEW 102	static_assert(std::is_integral_v<std::remove_reference_t<yT>>,	×
NEW 103	"y must be of an integral type");	×
NEW 104	static_assert(std::is_integral_v<std::remove_reference_t<xT>>,	×
NEW 105	"x must be of an integral type");	×
106
NEW 107	return sycl::select_from_group(sbgroup, data[y], x);	×
NEW 108	}	×
109
110	template <typename iT>
111	T broadcast(const iT &idx) const
NEW 112	{	×
NEW 113	if constexpr (Size == 1) {	×
NEW 114	return broadcast(0, idx);	×
115	}
116	else {
117	return broadcast(idx / size_x(), idx % size_x());
118	}
NEW 119	}	×
120
121	template <typename yT, typename xT>
122	T shift_left(const yT &y, const xT &x) const
NEW 123	{	×
NEW 124	static_assert(std::is_integral_v<yT>, "y must be of an integral type");	×
NEW 125	static_assert(std::is_integral_v<xT>, "x must be of an integral type");	×
126
NEW 127	return sycl::shift_group_left(sbgroup, data[y], x);	×
NEW 128	}	×
129
130	template <typename yT, typename xT>
131	T shift_right(const yT &y, const xT &x) const
NEW 132	{	×
NEW 133	static_assert(std::is_integral_v<yT>, "y must be of an integral type");	×
NEW 134	static_assert(std::is_integral_v<xT>, "x must be of an integral type");	×
135
NEW 136	return sycl::shift_group_right(sbgroup, data[y], x);	×
NEW 137	}	×
138
139	constexpr SizeT size_y() const
NEW 140	{	×
NEW 141	return _size;	×
NEW 142	}	×
143
144	SizeT size_x() const
NEW 145	{	×
NEW 146	return sbgroup.get_max_local_range()[0];	×
NEW 147	}	×
148
149	SizeT total_size() const
NEW 150	{	×
NEW 151	return size_x() * size_y();	×
NEW 152	}	×
153
154	ncT *ptr()
155	{
156	return data;
157	}
158
159	SizeT x() const
NEW 160	{	×
NEW 161	return sbgroup.get_local_linear_id();	×
NEW 162	}	×
163
164	protected:
165	const sycl::sub_group sbgroup;
166	ncT data[Size];
167	};
168
169	template <typename T, uint32_t Size = 1>
170	struct RegistryData : public _RegistryDataStorage<T, Size>
171	{
172	using SizeT = typename _RegistryDataStorage<T, Size>::SizeT;
173
174	using _RegistryDataStorage<T, Size>::_RegistryDataStorage;
175
176	template <typename LaneIdT,
177	typename Condition,
178	typename = std::enable_if_t<
179	std::is_invocable_r_v<bool, Condition, SizeT>>>
180	void fill_lane(const LaneIdT &lane_id, const T &value, Condition &&mask)
NEW 181	{	×
NEW 182	static_assert(std::is_integral_v<LaneIdT>,	×
NEW 183	"lane_id must be of an integral type");	×
NEW 184	if (mask(this->x())) {	×
NEW 185	this->data[lane_id] = value;	×
NEW 186	}	×
NEW 187	}	×
188
189	template <typename LaneIdT>
190	void fill_lane(const LaneIdT &lane_id, const T &value, const bool &mask)
NEW 191	{	×
NEW 192	fill_lane(lane_id, value, [mask](auto &&) { return mask; });	×
NEW 193	}	×
194
195	template <typename LaneIdT>
196	void fill_lane(const LaneIdT &lane_id, const T &value)
197	{
198	fill_lane(lane_id, value, true);
199	}
200
201	template <typename Condition,
202	typename = std::enable_if_t<
203	std::is_invocable_r_v<bool, Condition, SizeT, SizeT>>>
204	void fill(const T &value, Condition &&mask)
NEW 205	{	×
NEW 206	for (SizeT i = 0; i < Size; ++i) {	×
NEW 207	fill_lane(i, value, mask(i, this->x()));	×
NEW 208	}	×
NEW 209	}	×
210
211	void fill(const T &value)
NEW 212	{	×
NEW 213	fill(value, [](auto &&, auto &&) { return true; });	×
NEW 214	}	×
215
216	template <typename LaneIdT,
217	typename Condition,
218	typename = std::enable_if_t<
219	std::is_invocable_r_v<bool, Condition, const T *const>>>
220	T *load_lane(const LaneIdT &lane_id,
221	const T *const data,
222	Condition &&mask,
223	const T &default_v)
NEW 224	{	×
NEW 225	static_assert(std::is_integral_v<LaneIdT>,	×
NEW 226	"lane_id must be of an integral type");	×
NEW 227	this->data[lane_id] = mask(data) ? data[0] : default_v;	×
228
NEW 229	return data + this->size_x();	×
NEW 230	}	×
231
232	template <typename LaneIdT>
233	T *load_lane(const LaneIdT &laned_id,
234	const T *const data,
235	const bool &mask,
236	const T &default_v)
237	{
238	return load_lane(
239	laned_id, data, [mask](auto &&) { return mask; }, default_v);
240	}
241
242	template <typename LaneIdT>
243	T load_lane(const LaneIdT &laned_id, const T const data)
244	{
245	constexpr T default_v = 0;
246	return load_lane(laned_id, data, true, default_v);
247	}
248
249	template <typename yStrideT,
250	typename Condition,
251	typename = std::enable_if_t<
252	std::is_invocable_r_v<bool, Condition, const T *const>>>
253	T load(const T const data,
254	const yStrideT &y_stride,
255	Condition &&mask,
256	const T &default_v)
NEW 257	{	×
NEW 258	auto *it = data;	×
NEW 259	for (SizeT i = 0; i < Size; ++i) {	×
NEW 260	load_lane(i, it, mask, default_v);	×
NEW 261	it += y_stride;	×
NEW 262	}	×
263
NEW 264	return it;	×
NEW 265	}	×
266
267	template <typename yStrideT>
268	T load(const T const data,
269	const yStrideT &y_stride,
270	const bool &mask,
271	const T &default_v)
272	{
273	return load(
274	data, y_stride, [mask](auto &&) { return mask; }, default_v);
275	}
276
277	template <typename Condition,
278	typename = std::enable_if_t<
279	std::is_invocable_r_v<bool, Condition, const T *const>>>
280	T load(const T const data, Condition &&mask, const T &default_v)
NEW 281	{	×
NEW 282	return load(data, this->size_x(), mask, default_v);	×
NEW 283	}	×
284
285	T load(const T const data, const bool &mask, const T &default_v)
NEW 286	{	×
NEW 287	return load(	×
NEW 288	data, [mask](auto &&) { return mask; }, default_v);	×
NEW 289	}	×
290
291	T load(const T const data)
292	{
293	constexpr T default_v = 0;
294	return load(data, true, default_v);
295	}
296
297	template <typename LaneIdT,
298	typename Condition,
299	typename = std::enable_if_t<
300	std::is_invocable_r_v<bool, Condition, const T *const>>>
301	T store_lane(const LaneIdT &lane_id, T const data, Condition &&mask)
NEW 302	{	×
NEW 303	static_assert(std::is_integral_v<LaneIdT>,	×
NEW 304	"lane_id must be of an integral type");	×
305
NEW 306	if (mask(data)) {	×
NEW 307	data[0] = this->data[lane_id];	×
NEW 308	}	×
309
NEW 310	return data + this->size_x();	×
NEW 311	}	×
312
313	template <typename LaneIdT>
314	T store_lane(const LaneIdT &lane_id, T const data, const bool &mask)
315	{
316	return store_lane(lane_id, data, [mask](auto &&) { return mask; });
317	}
318
319	template <typename LaneIdT>
320	T store_lane(const LaneIdT &lane_id, T const data)
321	{
322	return store_lane(lane_id, data, true);
323	}
324
325	template <typename yStrideT,
326	typename Condition,
327	typename = std::enable_if_t<
328	std::is_invocable_r_v<bool, Condition, const T *const>>>
329	T store(T const data, const yStrideT &y_stride, Condition &&condition)
NEW 330	{	×
NEW 331	auto *it = data;	×
NEW 332	for (SizeT i = 0; i < Size; ++i) {	×
NEW 333	store_lane(i, it, condition);	×
NEW 334	it += y_stride;	×
NEW 335	}	×
336
NEW 337	return it;	×
NEW 338	}	×
339
340	template <typename yStrideT>
341	T store(T const data, const yStrideT &y_stride, const bool &mask)
342	{
343	return store(data, y_stride, [mask](auto &&) { return mask; });
344	}
345
346	template <typename Condition,
347	typename = std::enable_if_t<
348	std::is_invocable_r_v<bool, Condition, const T *const>>>
349	T store(T const data, Condition &&condition)
NEW 350	{	×
NEW 351	return store(data, this->size_x(), condition);	×
NEW 352	}	×
353
354	T store(T const data, const bool &mask)
355	{
356	return store(data, [mask](auto &&) { return mask; });
357	}
358
359	T store(T const data)
360	{
361	return store(data, true);
362	}
363	};
364
365	template <typename T, uint32_t Size>
366	struct RegistryWindow : public RegistryData<T, Size>
367	{
368	using SizeT = typename RegistryData<T, Size>::SizeT;
369
370	using RegistryData<T, Size>::RegistryData;
371
372	template <typename shT>
373	void advance_left(const shT &shift, const T &fill_value)
NEW 374	{	×
NEW 375	static_assert(std::is_integral_v<shT>,	×
NEW 376	"shift must be of an integral type");	×
377
NEW 378	uint32_t shift_r = this->size_x() - shift;	×
NEW 379	for (SizeT i = 0; i < Size; ++i) {	×
NEW 380	this->data[i] = this->shift_left(i, shift);	×
NEW 381	auto border =	×
NEW 382	i < Size - 1 ? this->shift_right(i + 1, shift_r) : fill_value;	×
NEW 383	if (this->x() >= shift_r) {	×
NEW 384	this->data[i] = border;	×
NEW 385	}	×
NEW 386	}	×
NEW 387	}	×
388
389	void advance_left(const T &fill_value)
NEW 390	{	×
NEW 391	advance_left(1, fill_value);	×
NEW 392	}	×
393
394	void advance_left()
NEW 395	{	×
NEW 396	constexpr T fill_value = 0;	×
NEW 397	advance_left(fill_value);	×
NEW 398	}	×
399	};
400
401	template <typename T, typename SizeT = size_t>
402	class Span
403	{
404	public:
405	using value_type = T;
406	using size_type = SizeT;
407
408	Span(T *const data, const SizeT size) : data_(data), size_(size) {}	1,020✔
409
410	T *begin() const
NEW 411	{	×
NEW 412	return data();	×
NEW 413	}	×
414
415	T *end() const
NEW 416	{	×
NEW 417	return data() + size();	×
NEW 418	}	×
419
420	SizeT size() const
NEW 421	{	×
NEW 422	return size_;	×
NEW 423	}	×
424
425	T *data() const
NEW 426	{	×
NEW 427	return data_;	×
NEW 428	}	×
429
430	protected:
431	T *const data_;
432	const SizeT size_;
433	};
434
435	template <typename T, typename SizeT = size_t>
436	Span<T, SizeT> make_span(T *const data, const SizeT size)
437	{	680✔
438	return Span<T, SizeT>(data, size);	680✔
439	}	680✔
440
441	template <typename T, typename SizeT = size_t>
442	class PaddedSpan : public Span<T, SizeT>
443	{
444	public:
445	using value_type = T;
446	using size_type = SizeT;
447
448	PaddedSpan(T *const data, const SizeT size, const SizeT pad)
449	: Span<T>(data, size), pad_(pad)	340✔
450	{	340✔
451	}	340✔
452
453	T *padded_begin() const
NEW 454	{	×
NEW 455	return this->begin() - pad();	×
NEW 456	}	×
457
458	SizeT pad() const
NEW 459	{	×
NEW 460	return pad_;	×
NEW 461	}	×
462
463	protected:
464	const SizeT pad_;
465	};
466
467	template <typename T, typename SizeT = size_t>
468	PaddedSpan<T, SizeT>
469	make_padded_span(T *const data, const SizeT size, const SizeT offset)
470	{	340✔
471	return PaddedSpan<T, SizeT>(data, size, offset);	340✔
472	}	340✔
473
474	template <typename Results,
475	typename AData,
476	typename VData,
477	typename Op,
478	typename Red>
479	void process_block(Results &results,
480	uint32_t r_size,
481	AData &a_data,
482	VData &v_data,
483	uint32_t block_size,
484	Op op,
485	Red red)
NEW 486	{	×
NEW 487	for (uint32_t i = 0; i < block_size; ++i) {	×
NEW 488	auto v_val = v_data.broadcast(i);	×
NEW 489	for (uint32_t r = 0; r < r_size; ++r) {	×
NEW 490	results[r] = red(results[r], op(a_data[r], v_val));	×
NEW 491	}	×
NEW 492	a_data.advance_left();	×
NEW 493	}	×
NEW 494	}	×
495
496	template <typename SizeT>
497	SizeT get_global_linear_id(const uint32_t wpi, const sycl::nd_item<1> &item)
NEW 498	{	×
NEW 499	auto sbgroup = item.get_sub_group();	×
NEW 500	const auto sg_loc_id = sbgroup.get_local_linear_id();	×
501
NEW 502	const SizeT sg_base_id = wpi * (item.get_global_linear_id() - sg_loc_id);	×
NEW 503	const SizeT id = sg_base_id + sg_loc_id;	×
504
NEW 505	return id;	×
NEW 506	}	×
507
508	template <typename SizeT>
509	uint32_t get_results_num(const uint32_t wpi,
510	const SizeT size,
511	const SizeT global_id,
512	const sycl::nd_item<1> &item)
NEW 513	{	×
NEW 514	auto sbgroup = item.get_sub_group();	×
515
NEW 516	const auto sbg_size = sbgroup.get_max_local_range()[0];	×
NEW 517	const auto size_ = sycl::sub_sat(size, global_id);	×
NEW 518	return std::min(SizeT(wpi), CeilDiv(size_, sbg_size));	×
NEW 519	}	×
520
521	template <uint32_t WorkPI,
522	typename T,
523	typename SizeT,
524	typename Op,
525	typename Red>
526	class sliding_window1d_kernel;
527
528	template <uint32_t WorkPI,
529	typename T,
530	typename SizeT,
531	typename Op,
532	typename Red>
533	void submit_sliding_window1d(const PaddedSpan<const T, SizeT> &a,
534	const Span<const T, SizeT> &v,
535	const Op &op,
536	const Red &red,
537	Span<T, SizeT> &out,
538	sycl::nd_range<1> nd_range,
539	sycl::handler &cgh)
540	{	76✔
541	cgh.parallel_for<sliding_window1d_kernel<WorkPI, T, SizeT, Op, Red>>(	76✔
542	nd_range, [=](sycl::nd_item<1> item) {	76✔
NEW 543	auto glid = get_global_linear_id<SizeT>(WorkPI, item);	×
544
NEW 545	auto results = RegistryData<T, WorkPI>(item);	×
NEW 546	results.fill(0);	×
547
NEW 548	auto results_num = get_results_num(WorkPI, out.size(), glid, item);	×
549
NEW 550	const auto *a_begin = a.begin();	×
NEW 551	const auto *a_end = a.end();	×
552
NEW 553	auto sbgroup = item.get_sub_group();	×
554
NEW 555	const auto chunks_count =	×
NEW 556	CeilDiv(v.size(), sbgroup.get_max_local_range()[0]);	×
557
NEW 558	const auto *a_ptr = &a.padded_begin()[glid];	×
559
NEW 560	auto _a_load_cond = [a_begin, a_end](auto &&ptr) {	×
NEW 561	return ptr >= a_begin && ptr < a_end;	×
NEW 562	};	×
563
NEW 564	auto a_data = RegistryWindow<const T, WorkPI + 1>(item);	×
NEW 565	a_ptr = a_data.load(a_ptr, _a_load_cond, 0);	×
566
NEW 567	const auto *v_ptr = &v.begin()[sbgroup.get_local_linear_id()];	×
NEW 568	auto v_size = v.size();	×
569
NEW 570	for (uint32_t b = 0; b < chunks_count; ++b) {	×
NEW 571	auto v_data = RegistryData<const T>(item);	×
NEW 572	v_ptr = v_data.load(v_ptr, v_data.x() < v_size, 0);	×
573
NEW 574	uint32_t chunk_size_ =	×
NEW 575	std::min(v_size, SizeT(v_data.total_size()));	×
NEW 576	process_block(results, results_num, a_data, v_data, chunk_size_,	×
NEW 577	op, red);	×
578
NEW 579	if (b != chunks_count - 1) {	×
NEW 580	a_ptr = a_data.load_lane(a_data.size_y() - 1, a_ptr,	×
NEW 581	_a_load_cond, 0);	×
NEW 582	v_size -= v_data.total_size();	×
NEW 583	}	×
NEW 584	}	×
585
NEW 586	auto *const out_ptr = out.begin();	×
NEW 587	auto *const out_end = out.end();	×
NEW 588	results.store(&out_ptr[glid],	×
NEW 589	[out_end](auto &&ptr) { return ptr < out_end; });	×
NEW 590	});	×
591	}	76✔
592
593	template <uint32_t WorkPI,
594	typename T,
595	typename SizeT,
596	typename Op,
597	typename Red>
598	class sliding_window1d_small_kernel;
599
600	template <uint32_t WorkPI,
601	typename T,
602	typename SizeT,
603	typename Op,
604	typename Red>
605	void submit_sliding_window1d_small_kernel(const PaddedSpan<const T, SizeT> &a,
606	const Span<const T, SizeT> &v,
607	const Op &op,
608	const Red &red,
609	Span<T, SizeT> &out,
610	sycl::nd_range<1> nd_range,
611	sycl::handler &cgh)
612	{	264✔
613	cgh.parallel_for<sliding_window1d_small_kernel<WorkPI, T, SizeT, Op, Red>>(	264✔
614	nd_range, [=](sycl::nd_item<1> item) {	264✔
NEW 615	auto glid = get_global_linear_id<SizeT>(WorkPI, item);	×
616
NEW 617	auto results = RegistryData<T, WorkPI>(item);	×
NEW 618	results.fill(0);	×
619
NEW 620	auto sbgroup = item.get_sub_group();	×
NEW 621	auto sg_size = sbgroup.get_max_local_range()[0];	×
622
NEW 623	const uint32_t to_read = WorkPI * sg_size + v.size();	×
NEW 624	const auto *a_begin = a.begin();	×
625
NEW 626	const auto *a_ptr = &a.padded_begin()[glid];	×
NEW 627	const auto *a_end = std::min(a_ptr + to_read, a.end());	×
628
NEW 629	auto _a_load_cond = [a_begin, a_end](auto &&ptr) {	×
NEW 630	return ptr >= a_begin && ptr < a_end;	×
NEW 631	};	×
632
NEW 633	auto a_data = RegistryWindow<const T, WorkPI + 1>(item);	×
NEW 634	a_data.load(a_ptr, _a_load_cond, 0);	×
635
NEW 636	const auto *v_ptr = &v.begin()[sbgroup.get_local_linear_id()];	×
NEW 637	auto v_size = v.size();	×
638
NEW 639	auto v_data = RegistryData<const T>(item);	×
NEW 640	v_ptr = v_data.load(v_ptr, v_data.x() < v_size, 0);	×
641
NEW 642	auto results_num = get_results_num(WorkPI, out.size(), glid, item);	×
643
NEW 644	process_block(results, results_num, a_data, v_data, v_size, op,	×
NEW 645	red);	×
646
NEW 647	auto *const out_ptr = out.begin();	×
NEW 648	auto *const out_end = out.end();	×
NEW 649	results.store(&out_ptr[glid],	×
NEW 650	[out_end](auto &&ptr) { return ptr < out_end; });	×
NEW 651	});	×
652	}	264✔
653
654	void validate(const usm_ndarray &a,
655	const usm_ndarray &v,
656	const usm_ndarray &out,
657	const size_t l_pad,
658	const size_t r_pad);
659	} // namespace sliding_window1d
660	} // namespace statistics

IntelPython / dpnp / 12076405164

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous