11385101423

Committed 17 Oct 2024 12:38PM UTC coverage: 95.216% (-0.09%) from 95.303%

Build # 11385101423

Build Type

push

github

Committed by

fknorr

Commit Message

Add type trait to match ReductionsAndKernel parameter lists

Run Details

3050 of 3442 branches covered (88.61%)

Branch coverage included in aggregate %.

5 of 5 new or added lines in 1 file covered. (100.0%)

24 existing lines in 8 files now uncovered.

6821 of 6925 relevant lines covered (98.5%)

1449510.68 hits per line

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

86.0

/include/item.h

#pragma once

#include "ranges.h"

namespace celerity {

template <int Dims>
class item;
template <int Dims>
class group;
template <int Dims>
class nd_item;

namespace detail {

        template <int Dims>
        inline item<Dims> make_item(id<Dims> absolute_global_id, id<Dims> global_offset, range<Dims> global_range) {
                return item<Dims>{absolute_global_id, global_offset, global_range};
        }

        template <int Dims>
        inline group<Dims> make_group(const sycl::group<std::max(1, Dims)>& sycl_group, const id<Dims>& group_id, const range<Dims>& group_range) {
                return group<Dims>{sycl_group, group_id, group_range};
        }

        template <int Dims>
        nd_item<Dims> make_nd_item(const sycl::nd_item<std::max(1, Dims)>& sycl_item, const range<Dims>& global_range, const id<Dims>& global_offset,
            const id<Dims>& chunk_offset, const range<Dims>& group_range, const id<Dims>& group_offset) {
                return nd_item<Dims>{sycl_item, global_range, global_offset, chunk_offset, group_range, group_offset};
        }

        template <int Dims>
        inline sycl::nd_item<std::max(1, Dims)>& get_sycl_item(nd_item<Dims>& nd_item) {
                return nd_item.m_sycl_item;
        }

        template <int Dims>
        inline const sycl::nd_item<std::max(1, Dims)>& get_sycl_item(const nd_item<Dims>& nd_item) {
                return nd_item.m_sycl_item;
        }

        template <int Dims>
        inline sycl::group<std::max(1, Dims)>& get_sycl_group(group<Dims>& g) {
                return g.m_sycl_group;
        }

        template <int Dims>
        inline const sycl::group<std::max(1, Dims)>& get_sycl_group(const group<Dims>& g) {
                return g.m_sycl_group;
        }

} // namespace detail

// We replace sycl::item with celerity::item to correctly expose the cluster global size instead of the chunk size to the user.
template <int Dims = 1>
class item {
  public:
        item() = delete;

        friend bool operator==(const item& lhs, const item& rhs) {
                return lhs.m_absolute_global_id == rhs.m_absolute_global_id && lhs.m_global_offset == rhs.m_global_offset && lhs.m_global_range == rhs.m_global_range;
        }

        friend bool operator!=(const item& lhs, const item& rhs) { return !(lhs == rhs); }

        id<Dims> get_id() const { return m_absolute_global_id; }

        size_t get_id(int dimension) const { return m_absolute_global_id[dimension]; }

        operator id<Dims>() const { return m_absolute_global_id; } // NOLINT(google-explicit-constructor)

        size_t operator[](int dimension) const { return m_absolute_global_id[dimension]; }

        range<Dims> get_range() const { return m_global_range; }

        size_t get_range(int dimension) const { return m_global_range[dimension]; }

        size_t get_linear_id() const { return detail::get_linear_index(m_global_range, m_absolute_global_id - m_global_offset); }

        id<Dims> get_offset() const { return m_global_offset; }

  private:
        template <int D>
        friend item<D> celerity::detail::make_item(id<D>, id<D>, range<D>);

        id<Dims> m_absolute_global_id;
        id<Dims> m_global_offset;
        range<Dims> m_global_range;

        explicit item(id<Dims> absolute_global_id, id<Dims> global_offset, range<Dims> global_range)
            : m_absolute_global_id(absolute_global_id), m_global_offset(global_offset), m_global_range(global_range) {}
};


template <int Dims = 1>
class group {
  public:
        using id_type = id<Dims>;
        using range_type = range<Dims>;
        using linear_id_type = size_t;
        static constexpr int dimensions = Dims;
        static constexpr memory_scope fence_scope = memory_scope_work_group;

        id<Dims> get_group_id() const { return m_group_id; }

        size_t get_group_id(int dimension) const { return m_group_id[dimension]; }

        id<Dims> get_local_id() const { return m_sycl_group.get_local_id(); }

        size_t get_local_id(int dimension) const { return m_sycl_group.get_local_id(dimension); }

        range<Dims> get_local_range() const { return m_sycl_group.get_local_range(); }

        size_t get_local_range(int dimension) const { return m_sycl_group.get_local_range(dimension); }

        range<Dims> get_group_range() const { return m_group_range; }

        size_t get_group_range(int dimension) const { return m_group_range[dimension]; }

        range<Dims> get_max_local_range() const { return m_sycl_group.get_max_local_range(); }

        size_t operator[](int dimension) const { return m_group_id[dimension]; }

        size_t get_group_linear_id() const { return detail::get_linear_index(m_group_range, m_group_id); }

        size_t get_local_linear_id() const { return m_sycl_group.get_local_linear_id(); }

        size_t get_group_linear_range() const { return m_group_range.size(); }

        size_t get_local_linear_range() const { return m_sycl_group.get_local_range().size(); }

        bool leader() const { return m_sycl_group.get_local_id() == id<Dims>{}; }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements) const {
                return m_sycl_group.async_work_group_copy(dest, src, num_elements);
        }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements) const {
                return m_sycl_group.async_work_group_copy(dest, src, num_elements);
        }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements, size_t src_stride) const {
                return m_sycl_group.async_work_group_copy(dest, src, num_elements, src_stride);
        }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements, size_t dest_stride) const {
                return m_sycl_group.async_work_group_copy(dest, src, num_elements, dest_stride);
        }

        template <typename... DeviceEvents>
        void wait_for(DeviceEvents... events) const {
                m_sycl_group.wait_for(events...);
        }

  private:
        constexpr static int sycl_dims = std::max(1, Dims);

        sycl::group<sycl_dims> m_sycl_group;
        id<Dims> m_group_id;
        range<Dims> m_group_range;

        template <int D>
        friend group<D> celerity::detail::make_group(const sycl::group<std::max(1, D)>& sycl_group, const id<D>& group_id, const range<D>& group_range);

        template <int D>
        friend sycl::group<std::max(1, D)>& celerity::detail::get_sycl_group(group<D>&);

        template <int D>
        friend const sycl::group<std::max(1, D)>& celerity::detail::get_sycl_group(const group<D>&);

        explicit group(const sycl::group<sycl_dims>& sycl_group, const id<Dims>& group_id, const range<Dims>& group_range)
            : m_sycl_group(sycl_group), m_group_id(group_id), m_group_range(group_range) {}
};


// We replace sycl::nd_item with celerity::nd_item to correctly expose the cluster global size instead of the chunk size to the user.
template <int Dims = 1>
class nd_item {
  public:
        nd_item() = delete;

        id<Dims> get_global_id() const { return m_global_id; }

        size_t get_global_id(const int dimension) const { return m_global_id[dimension]; }

        size_t get_global_linear_id() const { return detail::get_linear_index(m_global_range, m_global_id); }

        id<Dims> get_local_id() const { return m_sycl_item.get_local_id(); }

        size_t get_local_id(int dimension) const { return m_sycl_item.get_local_id(dimension); }

        size_t get_local_linear_id() const { return m_sycl_item.get_local_linear_id(); }

        group<Dims> get_group() const { return detail::make_group<Dims>(m_sycl_item.get_group(), m_group_id, m_group_range); }

        size_t get_group(const int dimension) const { return m_group_id[dimension]; }

        size_t get_group_linear_id() const { return detail::get_linear_index(m_group_range, m_group_id); }

        range<Dims> get_group_range() const { return m_group_range; }

        size_t get_group_range(const int dimension) const { return m_group_range[dimension]; }

        sycl::sub_group get_sub_group() const { return m_sycl_item.get_sub_group(); }

        range<Dims> get_global_range() const { return m_global_range; }

        size_t get_global_range(const int dimension) const { return m_global_range[dimension]; }

        range<Dims> get_local_range() const { return m_sycl_item.get_local_range(); }

        size_t get_local_range(const int dimension) const { return m_sycl_item.get_local_range(dimension); }

        id<Dims> get_offset() const { return m_global_offset; }

        celerity::nd_range<Dims> get_nd_range() const { return celerity::nd_range<Dims>(get_global_range(), get_local_range(), get_offset()); }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements) const {
                return m_sycl_item.async_work_group_copy(dest, src, num_elements);
        }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements) const {
                return m_sycl_item.async_work_group_copy(dest, src, num_elements);
        }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements, size_t src_stride) const {
                return m_sycl_item.async_work_group_copy(dest, src, num_elements, src_stride);
        }

        template <typename T>
        sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements, size_t dest_stride) const {
                return m_sycl_item.async_work_group_copy(dest, src, num_elements, dest_stride);
        }

        template <typename... DeviceEvents>
        void wait_for(DeviceEvents... events) const {
                m_sycl_item.wait_for(events...);
        }

  private:
        constexpr static int sycl_dims = std::max(1, Dims);

        sycl::nd_item<sycl_dims> m_sycl_item;
        id<Dims> m_global_id;
        id<Dims> m_global_offset;
        range<Dims> m_global_range;
        id<Dims> m_group_id;
        range<Dims> m_group_range;

        template <int D>
        friend nd_item<D> celerity::detail::make_nd_item(
            const sycl::nd_item<std::max(1, D)>&, const range<D>&, const id<D>&, const id<D>&, const range<D>&, const id<D>&);

        template <int D>
        friend sycl::nd_item<std::max(1, D)>& celerity::detail::get_sycl_item(group<D>& nd_item);

        template <int D>
        friend const sycl::nd_item<std::max(1, D)>& celerity::detail::get_sycl_item(const group<D>& nd_item);

        explicit nd_item(const sycl::nd_item<std::max(1, Dims)>& sycl_item, const range<Dims>& global_range, const id<Dims>& global_offset,
            const id<Dims>& chunk_offset, const range<Dims>& group_range, const id<Dims>& group_offset)
            : m_sycl_item(sycl_item), m_global_id(chunk_offset + detail::id_cast<Dims>(celerity::id(sycl_item.get_global_id()))), m_global_offset(global_offset),
              m_global_range(global_range), m_group_id(group_offset + detail::id_cast<Dims>(celerity::id(sycl_item.get_group().get_group_id()))),
              m_group_range(group_range) {}
};


using sycl::group_barrier;

template <int Dims>
void group_barrier(const group<Dims>& g, memory_scope scope = memory_scope_work_group) {
        sycl::group_barrier(detail::get_sycl_group(g), static_cast<sycl::memory_scope>(scope)); // identical representation
}

using sycl::group_broadcast;

template <int Dims, typename T>
inline T group_broadcast(const group<Dims>& g, T x) {
        return sycl::group_broadcast(detail::get_sycl_group(g), x);
}

template <int Dims, typename T>
inline T group_broadcast(const group<Dims>& g, T x, size_t local_linear_id) {
        return sycl::group_broadcast(detail::get_sycl_group(g), x, local_linear_id);
}

template <int Dims, typename T>
inline T group_broadcast(const group<Dims>& g, T x, const id<Dims>& local_id) {
        return sycl::group_broadcast(detail::get_sycl_group(g), x, sycl::id<Dims>(local_id));
};


using sycl::joint_any_of;

template <int Dims, typename Ptr, typename Predicate>
bool joint_any_of(const group<Dims>& g, Ptr first, Ptr last, Predicate pred) {
        return sycl::joint_any_of(detail::get_sycl_group(g), first, last, pred);
}


using sycl::any_of_group;

template <int Dims, typename T, typename Predicate>
bool any_of_group(const group<Dims>& g, T x, Predicate pred) {
        return sycl::any_of_group(detail::get_sycl_group(g), x, pred);
}

template <int Dims>
bool any_of_group(const group<Dims>& g, bool pred) {
        return sycl::any_of_group(detail::get_sycl_group(g), pred);
}


using sycl::joint_all_of;

template <int Dims, typename Ptr, typename Predicate>
bool joint_all_of(const group<Dims>& g, Ptr first, Ptr last, Predicate pred) {
        return sycl::joint_all_of(detail::get_sycl_group(g), first, last, pred);
}


using sycl::all_of_group;

template <int Dims, typename T, typename Predicate>
bool all_of_group(const group<Dims>& g, T x, Predicate pred) {
        return sycl::all_of_group(detail::get_sycl_group(g), x, pred);
}

template <int Dims>
bool all_of_group(const group<Dims>& g, bool pred) {
        return sycl::all_of_group(detail::get_sycl_group(g), pred);
}


using sycl::joint_none_of;

template <int Dims, typename Ptr, typename Predicate>
bool joint_none_of(const group<Dims>& g, Ptr first, Ptr last, Predicate pred) {
        return sycl::joint_none_of(detail::get_sycl_group(g), first, last, pred);
}


using sycl::none_of_group;

template <int Dims, typename T, typename Predicate>
bool none_of_group(const group<Dims>& g, T x, Predicate pred) {
        return sycl::none_of_group(detail::get_sycl_group(g), x, pred);
}

template <int Dims>
bool none_of_group(const group<Dims>& g, bool pred) {
        return sycl::none_of_group(detail::get_sycl_group(g), pred);
}


using sycl::permute_group_by_xor;
using sycl::shift_group_left;
using sycl::shift_group_right;

template <int Dims, typename T>
T shift_group_left(const group<Dims>& g, T x, size_t delta = 1) {
        return sycl::shift_group_left(detail::get_sycl_group(g), x, delta);
}

template <int Dims, typename T>
T shift_group_right(const group<Dims>& g, T x, size_t delta = 1) {
        return sycl::shift_group_right(detail::get_sycl_group(g), x, delta);
}

template <int Dims, typename T>
T permute_group_by_xor(const group<Dims>& g, T x, size_t mask) {
        return sycl::permute_group_by_xor(detail::get_sycl_group(g), x, mask);
}


using sycl::select_from_group;

template <int Dims, typename T>
T select_from_group(const group<Dims>& g, T x, size_t remote_local_id) {
        return sycl::select_from_group(detail::get_sycl_group(g), x, sycl::id<Dims>(remote_local_id));
}


using sycl::joint_reduce;

template <int Dims, typename Ptr, typename BinaryOperation>
typename std::iterator_traits<Ptr>::value_type joint_reduce(const group<Dims>& g, Ptr first, Ptr last, BinaryOperation binary_op) {
        return sycl::joint_reduce(detail::get_sycl_group(g), first, last, binary_op);
}

template <int Dims, typename Ptr, typename T, typename BinaryOperation>
T joint_reduce(const group<Dims>& g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
        return sycl::joint_reduce(detail::get_sycl_group(g), first, last, init, binary_op);
}


using sycl::reduce_over_group;

template <int Dims, typename T, typename BinaryOperation>
T reduce_over_group(const group<Dims>& g, T x, BinaryOperation binary_op) {
        return sycl::reduce_over_group(detail::get_sycl_group(g), x, binary_op);
}

template <int Dims, typename V, typename T, typename BinaryOperation>
T reduce_over_group(const group<Dims>& g, V x, T init, BinaryOperation binary_op) {
        return sycl::reduce_over_group(detail::get_sycl_group(g), x, init, binary_op);
}


using sycl::joint_exclusive_scan;

template <int Dims, typename InPtr, typename OutPtr, typename BinaryOperation>
OutPtr joint_exclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) {
        return sycl::joint_exclusive_scan(detail::get_sycl_group(g), first, last, result, binary_op);
}

template <int Dims, typename InPtr, typename OutPtr, typename T, typename BinaryOperation>
T joint_exclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) {
        return sycl::joint_exclusive_scan(detail::get_sycl_group(g), first, last, result, init, binary_op);
}


using sycl::exclusive_scan_over_group;

template <int Dims, typename T, typename BinaryOperation>
T exclusive_scan_over_group(const group<Dims>& g, T x, BinaryOperation binary_op) {
        return sycl::exclusive_scan_over_group(detail::get_sycl_group(g), x, binary_op);
}

template <int Dims, typename V, typename T, typename BinaryOperation>
T exclusive_scan_over_group(const group<Dims>& g, V x, T init, BinaryOperation binary_op) {
        return sycl::exclusive_scan_over_group(detail::get_sycl_group(g), x, init, binary_op);
}


using sycl::joint_inclusive_scan;

template <int Dims, typename InPtr, typename OutPtr, typename BinaryOperation>
OutPtr joint_inclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) {
        return sycl::joint_inclusive_scan(detail::get_sycl_group(g), first, last, result, binary_op);
}

template <int Dims, typename InPtr, typename OutPtr, typename T, typename BinaryOperation>
T joint_inclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) {
        return sycl::joint_inclusive_scan(detail::get_sycl_group(g), first, last, result, binary_op, init);
}

template <int Dims, typename T, typename BinaryOperation>
T inclusive_scan_over_group(const group<Dims>& g, T x, BinaryOperation binary_op) {
        return sycl::inclusive_scan_over_group(detail::get_sycl_group(g), x, binary_op);
}

using sycl::inclusive_scan_over_group;

template <int Dims, typename V, typename T, typename BinaryOperation>
T inclusive_scan_over_group(const group<Dims>& g, V x, BinaryOperation binary_op, T init) {
        return sycl::inclusive_scan_over_group(detail::get_sycl_group(g), x, binary_op, init);
}

} // namespace celerity

1	#pragma once
2
3	#include "ranges.h"
4
5	namespace celerity {
6
7	template <int Dims>
8	class item;
9	template <int Dims>
10	class group;
11	template <int Dims>
12	class nd_item;
13
14	namespace detail {
15
16	template <int Dims>
17	inline item<Dims> make_item(id<Dims> absolute_global_id, id<Dims> global_offset, range<Dims> global_range) {	41,970,009✔
18	return item<Dims>{absolute_global_id, global_offset, global_range};	41,970,009✔
19	}
20
21	template <int Dims>
22	inline group<Dims> make_group(const sycl::group<std::max(1, Dims)>& sycl_group, const id<Dims>& group_id, const range<Dims>& group_range) {	25,008✔
23	return group<Dims>{sycl_group, group_id, group_range};	25,008✔
24	}
25
26	template <int Dims>
27	nd_item<Dims> make_nd_item(const sycl::nd_item<std::max(1, Dims)>& sycl_item, const range<Dims>& global_range, const id<Dims>& global_offset,	25,330✔
28	const id<Dims>& chunk_offset, const range<Dims>& group_range, const id<Dims>& group_offset) {
29	return nd_item<Dims>{sycl_item, global_range, global_offset, chunk_offset, group_range, group_offset};	25,330✔
30	}
31
32	template <int Dims>
33	inline sycl::nd_item<std::max(1, Dims)>& get_sycl_item(nd_item<Dims>& nd_item) {
34	return nd_item.m_sycl_item;
35	}
36
37	template <int Dims>
38	inline const sycl::nd_item<std::max(1, Dims)>& get_sycl_item(const nd_item<Dims>& nd_item) {
39	return nd_item.m_sycl_item;
40	}
41
42	template <int Dims>
43	inline sycl::group<std::max(1, Dims)>& get_sycl_group(group<Dims>& g) {
44	return g.m_sycl_group;
45	}
46
47	template <int Dims>
48	inline const sycl::group<std::max(1, Dims)>& get_sycl_group(const group<Dims>& g) {	96✔
49	return g.m_sycl_group;	96✔
50	}
51
52	} // namespace detail
53
54	// We replace sycl::item with celerity::item to correctly expose the cluster global size instead of the chunk size to the user.
55	template <int Dims = 1>
56	class item {
57	public:
58	item() = delete;
59
60	friend bool operator==(const item& lhs, const item& rhs) {
61	return lhs.m_absolute_global_id == rhs.m_absolute_global_id && lhs.m_global_offset == rhs.m_global_offset && lhs.m_global_range == rhs.m_global_range;
62	}
63
64	friend bool operator!=(const item& lhs, const item& rhs) { return !(lhs == rhs); }
65
66	id<Dims> get_id() const { return m_absolute_global_id; }	25,165,833✔
67
68	size_t get_id(int dimension) const { return m_absolute_global_id[dimension]; }	4,005✔
69
70	operator id<Dims>() const { return m_absolute_global_id; } // NOLINT(google-explicit-constructor)	41,947,572✔
71
72	size_t operator[](int dimension) const { return m_absolute_global_id[dimension]; }	16,777,438✔
73
74	range<Dims> get_range() const { return m_global_range; }	18✔
75
76	size_t get_range(int dimension) const { return m_global_range[dimension]; }
77
78	size_t get_linear_id() const { return detail::get_linear_index(m_global_range, m_absolute_global_id - m_global_offset); }	2,578✔
79
80	id<Dims> get_offset() const { return m_global_offset; }	27✔
81
82	private:
83	template <int D>
84	friend item<D> celerity::detail::make_item(id<D>, id<D>, range<D>);
85
86	id<Dims> m_absolute_global_id;
87	id<Dims> m_global_offset;
88	range<Dims> m_global_range;
89
90	explicit item(id<Dims> absolute_global_id, id<Dims> global_offset, range<Dims> global_range)	41,970,009✔
91	: m_absolute_global_id(absolute_global_id), m_global_offset(global_offset), m_global_range(global_range) {}	41,970,009✔
92	};
93
94
95	template <int Dims = 1>
96	class group {
97	public:
98	using id_type = id<Dims>;
99	using range_type = range<Dims>;
100	using linear_id_type = size_t;
101	static constexpr int dimensions = Dims;
102	static constexpr memory_scope fence_scope = memory_scope_work_group;
103
104	id<Dims> get_group_id() const { return m_group_id; }	24,912✔
105
106	size_t get_group_id(int dimension) const { return m_group_id[dimension]; }
107
108	id<Dims> get_local_id() const { return m_sycl_group.get_local_id(); }	48,672✔
109
110	size_t get_local_id(int dimension) const { return m_sycl_group.get_local_id(dimension); }
111
112	range<Dims> get_local_range() const { return m_sycl_group.get_local_range(); }	48,672✔
113
114	size_t get_local_range(int dimension) const { return m_sycl_group.get_local_range(dimension); }
115
116	range<Dims> get_group_range() const { return m_group_range; }	24,912✔
117
118	size_t get_group_range(int dimension) const { return m_group_range[dimension]; }
119
120	range<Dims> get_max_local_range() const { return m_sycl_group.get_max_local_range(); }
121
122	size_t operator[](int dimension) const { return m_group_id[dimension]; }
123
124	size_t get_group_linear_id() const { return detail::get_linear_index(m_group_range, m_group_id); }	24,912✔
125
126	size_t get_local_linear_id() const { return m_sycl_group.get_local_linear_id(); }	24,912✔
127
128	size_t get_group_linear_range() const { return m_group_range.size(); }
129
130	size_t get_local_linear_range() const { return m_sycl_group.get_local_range().size(); }
131
132	bool leader() const { return m_sycl_group.get_local_id() == id<Dims>{}; }
133
134	template <typename T>
135	sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements) const {
136	return m_sycl_group.async_work_group_copy(dest, src, num_elements);
137	}
138
139	template <typename T>
140	sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements) const {
141	return m_sycl_group.async_work_group_copy(dest, src, num_elements);
142	}
143
144	template <typename T>
145	sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements, size_t src_stride) const {
146	return m_sycl_group.async_work_group_copy(dest, src, num_elements, src_stride);
147	}
148
149	template <typename T>
150	sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements, size_t dest_stride) const {
151	return m_sycl_group.async_work_group_copy(dest, src, num_elements, dest_stride);
152	}
153
154	template <typename... DeviceEvents>
155	void wait_for(DeviceEvents... events) const {
156	m_sycl_group.wait_for(events...);
157	}
158
159	private:
160	constexpr static int sycl_dims = std::max(1, Dims);
161
162	sycl::group<sycl_dims> m_sycl_group;
163	id<Dims> m_group_id;
164	range<Dims> m_group_range;
165
166	template <int D>
167	friend group<D> celerity::detail::make_group(const sycl::group<std::max(1, D)>& sycl_group, const id<D>& group_id, const range<D>& group_range);
168
169	template <int D>
170	friend sycl::group<std::max(1, D)>& celerity::detail::get_sycl_group(group<D>&);
171
172	template <int D>
173	friend const sycl::group<std::max(1, D)>& celerity::detail::get_sycl_group(const group<D>&);
174
175	explicit group(const sycl::group<sycl_dims>& sycl_group, const id<Dims>& group_id, const range<Dims>& group_range)	25,008✔
176	: m_sycl_group(sycl_group), m_group_id(group_id), m_group_range(group_range) {}	25,008✔
177	};
178
179
180	// We replace sycl::nd_item with celerity::nd_item to correctly expose the cluster global size instead of the chunk size to the user.
181	template <int Dims = 1>
182	class nd_item {
183	public:
184	nd_item() = delete;
185
186	id<Dims> get_global_id() const { return m_global_id; }	50,016✔
187
188	size_t get_global_id(const int dimension) const { return m_global_id[dimension]; }
189
190	size_t get_global_linear_id() const { return detail::get_linear_index(m_global_range, m_global_id); }	25,168✔
191
192	id<Dims> get_local_id() const { return m_sycl_item.get_local_id(); }	48,768✔
193
194	size_t get_local_id(int dimension) const { return m_sycl_item.get_local_id(dimension); }	64✔
195
196	size_t get_local_linear_id() const { return m_sycl_item.get_local_linear_id(); }	24,912✔
197
198	group<Dims> get_group() const { return detail::make_group<Dims>(m_sycl_item.get_group(), m_group_id, m_group_range); }	50,016✔
199
200	size_t get_group(const int dimension) const { return m_group_id[dimension]; }	×
201
202	size_t get_group_linear_id() const { return detail::get_linear_index(m_group_range, m_group_id); }	24,912✔
203
204	range<Dims> get_group_range() const { return m_group_range; }	24,912✔
205
206	size_t get_group_range(const int dimension) const { return m_group_range[dimension]; }
207
208	sycl::sub_group get_sub_group() const { return m_sycl_item.get_sub_group(); }
209
210	range<Dims> get_global_range() const { return m_global_range; }	24,912✔
211
212	size_t get_global_range(const int dimension) const { return m_global_range[dimension]; }
213
214	range<Dims> get_local_range() const { return m_sycl_item.get_local_range(); }	48,672✔
215
216	size_t get_local_range(const int dimension) const { return m_sycl_item.get_local_range(dimension); }	64✔
217
218	id<Dims> get_offset() const { return m_global_offset; }
219
220	celerity::nd_range<Dims> get_nd_range() const { return celerity::nd_range<Dims>(get_global_range(), get_local_range(), get_offset()); }
221
222	template <typename T>
223	sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements) const {
224	return m_sycl_item.async_work_group_copy(dest, src, num_elements);
225	}
226
227	template <typename T>
228	sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements) const {
229	return m_sycl_item.async_work_group_copy(dest, src, num_elements);
230	}
231
232	template <typename T>
233	sycl::device_event async_work_group_copy(decorated_local_ptr<T> dest, decorated_global_ptr<T> src, size_t num_elements, size_t src_stride) const {
234	return m_sycl_item.async_work_group_copy(dest, src, num_elements, src_stride);
235	}
236
237	template <typename T>
238	sycl::device_event async_work_group_copy(decorated_global_ptr<T> dest, decorated_local_ptr<T> src, size_t num_elements, size_t dest_stride) const {
239	return m_sycl_item.async_work_group_copy(dest, src, num_elements, dest_stride);
240	}
241
242	template <typename... DeviceEvents>
243	void wait_for(DeviceEvents... events) const {
244	m_sycl_item.wait_for(events...);
245	}
246
247	private:
248	constexpr static int sycl_dims = std::max(1, Dims);
249
250	sycl::nd_item<sycl_dims> m_sycl_item;
251	id<Dims> m_global_id;
252	id<Dims> m_global_offset;
253	range<Dims> m_global_range;
254	id<Dims> m_group_id;
255	range<Dims> m_group_range;
256
257	template <int D>
258	friend nd_item<D> celerity::detail::make_nd_item(
259	const sycl::nd_item<std::max(1, D)>&, const range<D>&, const id<D>&, const id<D>&, const range<D>&, const id<D>&);
260
261	template <int D>
262	friend sycl::nd_item<std::max(1, D)>& celerity::detail::get_sycl_item(group<D>& nd_item);
263
264	template <int D>
265	friend const sycl::nd_item<std::max(1, D)>& celerity::detail::get_sycl_item(const group<D>& nd_item);
266
267	explicit nd_item(const sycl::nd_item<std::max(1, Dims)>& sycl_item, const range<Dims>& global_range, const id<Dims>& global_offset,	25,330✔
268	const id<Dims>& chunk_offset, const range<Dims>& group_range, const id<Dims>& group_offset)
269	: m_sycl_item(sycl_item), m_global_id(chunk_offset + detail::id_cast<Dims>(celerity::id(sycl_item.get_global_id()))), m_global_offset(global_offset),	25,330✔
270	m_global_range(global_range), m_group_id(group_offset + detail::id_cast<Dims>(celerity::id(sycl_item.get_group().get_group_id()))),	25,330✔
271	m_group_range(group_range) {}	25,330✔
272	};
273
274
275	using sycl::group_barrier;
276
277	template <int Dims>
278	void group_barrier(const group<Dims>& g, memory_scope scope = memory_scope_work_group) {	96✔
279	sycl::group_barrier(detail::get_sycl_group(g), static_cast<sycl::memory_scope>(scope)); // identical representation	96✔
280	}	96✔
281
282	using sycl::group_broadcast;
283
284	template <int Dims, typename T>
UNCOV 285	inline T group_broadcast(const group<Dims>& g, T x) {	×
UNCOV 286	return sycl::group_broadcast(detail::get_sycl_group(g), x);	×
287	}
288
289	template <int Dims, typename T>
UNCOV 290	inline T group_broadcast(const group<Dims>& g, T x, size_t local_linear_id) {	×
UNCOV 291	return sycl::group_broadcast(detail::get_sycl_group(g), x, local_linear_id);	×
292	}
293
294	template <int Dims, typename T>
UNCOV 295	inline T group_broadcast(const group<Dims>& g, T x, const id<Dims>& local_id) {	×
UNCOV 296	return sycl::group_broadcast(detail::get_sycl_group(g), x, sycl::id<Dims>(local_id));	×
297	};
298
299
300	using sycl::joint_any_of;
301
302	template <int Dims, typename Ptr, typename Predicate>
303	bool joint_any_of(const group<Dims>& g, Ptr first, Ptr last, Predicate pred) {
304	return sycl::joint_any_of(detail::get_sycl_group(g), first, last, pred);
305	}
306
307
308	using sycl::any_of_group;
309
310	template <int Dims, typename T, typename Predicate>
311	bool any_of_group(const group<Dims>& g, T x, Predicate pred) {
312	return sycl::any_of_group(detail::get_sycl_group(g), x, pred);
313	}
314
315	template <int Dims>
316	bool any_of_group(const group<Dims>& g, bool pred) {
317	return sycl::any_of_group(detail::get_sycl_group(g), pred);
318	}
319
320
321	using sycl::joint_all_of;
322
323	template <int Dims, typename Ptr, typename Predicate>
324	bool joint_all_of(const group<Dims>& g, Ptr first, Ptr last, Predicate pred) {
325	return sycl::joint_all_of(detail::get_sycl_group(g), first, last, pred);
326	}
327
328
329	using sycl::all_of_group;
330
331	template <int Dims, typename T, typename Predicate>
332	bool all_of_group(const group<Dims>& g, T x, Predicate pred) {
333	return sycl::all_of_group(detail::get_sycl_group(g), x, pred);
334	}
335
336	template <int Dims>
337	bool all_of_group(const group<Dims>& g, bool pred) {
338	return sycl::all_of_group(detail::get_sycl_group(g), pred);
339	}
340
341
342	using sycl::joint_none_of;
343
344	template <int Dims, typename Ptr, typename Predicate>
345	bool joint_none_of(const group<Dims>& g, Ptr first, Ptr last, Predicate pred) {
346	return sycl::joint_none_of(detail::get_sycl_group(g), first, last, pred);
347	}
348
349
350	using sycl::none_of_group;
351
352	template <int Dims, typename T, typename Predicate>
353	bool none_of_group(const group<Dims>& g, T x, Predicate pred) {
354	return sycl::none_of_group(detail::get_sycl_group(g), x, pred);
355	}
356
357	template <int Dims>
358	bool none_of_group(const group<Dims>& g, bool pred) {
359	return sycl::none_of_group(detail::get_sycl_group(g), pred);
360	}
361
362
363	using sycl::permute_group_by_xor;
364	using sycl::shift_group_left;
365	using sycl::shift_group_right;
366
367	template <int Dims, typename T>
368	T shift_group_left(const group<Dims>& g, T x, size_t delta = 1) {
369	return sycl::shift_group_left(detail::get_sycl_group(g), x, delta);
370	}
371
372	template <int Dims, typename T>
373	T shift_group_right(const group<Dims>& g, T x, size_t delta = 1) {
374	return sycl::shift_group_right(detail::get_sycl_group(g), x, delta);
375	}
376
377	template <int Dims, typename T>
378	T permute_group_by_xor(const group<Dims>& g, T x, size_t mask) {
379	return sycl::permute_group_by_xor(detail::get_sycl_group(g), x, mask);
380	}
381
382
383	using sycl::select_from_group;
384
385	template <int Dims, typename T>
386	T select_from_group(const group<Dims>& g, T x, size_t remote_local_id) {
387	return sycl::select_from_group(detail::get_sycl_group(g), x, sycl::id<Dims>(remote_local_id));
388	}
389
390
391	using sycl::joint_reduce;
392
393	template <int Dims, typename Ptr, typename BinaryOperation>
394	typename std::iterator_traits<Ptr>::value_type joint_reduce(const group<Dims>& g, Ptr first, Ptr last, BinaryOperation binary_op) {
395	return sycl::joint_reduce(detail::get_sycl_group(g), first, last, binary_op);
396	}
397
398	template <int Dims, typename Ptr, typename T, typename BinaryOperation>
399	T joint_reduce(const group<Dims>& g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
400	return sycl::joint_reduce(detail::get_sycl_group(g), first, last, init, binary_op);
401	}
402
403
404	using sycl::reduce_over_group;
405
406	template <int Dims, typename T, typename BinaryOperation>
407	T reduce_over_group(const group<Dims>& g, T x, BinaryOperation binary_op) {
408	return sycl::reduce_over_group(detail::get_sycl_group(g), x, binary_op);
409	}
410
411	template <int Dims, typename V, typename T, typename BinaryOperation>
412	T reduce_over_group(const group<Dims>& g, V x, T init, BinaryOperation binary_op) {
413	return sycl::reduce_over_group(detail::get_sycl_group(g), x, init, binary_op);
414	}
415
416
417	using sycl::joint_exclusive_scan;
418
419	template <int Dims, typename InPtr, typename OutPtr, typename BinaryOperation>
420	OutPtr joint_exclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) {
421	return sycl::joint_exclusive_scan(detail::get_sycl_group(g), first, last, result, binary_op);
422	}
423
424	template <int Dims, typename InPtr, typename OutPtr, typename T, typename BinaryOperation>
425	T joint_exclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) {
426	return sycl::joint_exclusive_scan(detail::get_sycl_group(g), first, last, result, init, binary_op);
427	}
428
429
430	using sycl::exclusive_scan_over_group;
431
432	template <int Dims, typename T, typename BinaryOperation>
433	T exclusive_scan_over_group(const group<Dims>& g, T x, BinaryOperation binary_op) {
434	return sycl::exclusive_scan_over_group(detail::get_sycl_group(g), x, binary_op);
435	}
436
437	template <int Dims, typename V, typename T, typename BinaryOperation>
438	T exclusive_scan_over_group(const group<Dims>& g, V x, T init, BinaryOperation binary_op) {
439	return sycl::exclusive_scan_over_group(detail::get_sycl_group(g), x, init, binary_op);
440	}
441
442
443	using sycl::joint_inclusive_scan;
444
445	template <int Dims, typename InPtr, typename OutPtr, typename BinaryOperation>
446	OutPtr joint_inclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) {
447	return sycl::joint_inclusive_scan(detail::get_sycl_group(g), first, last, result, binary_op);
448	}
449
450	template <int Dims, typename InPtr, typename OutPtr, typename T, typename BinaryOperation>
451	T joint_inclusive_scan(const group<Dims>& g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) {
452	return sycl::joint_inclusive_scan(detail::get_sycl_group(g), first, last, result, binary_op, init);
453	}
454
455	template <int Dims, typename T, typename BinaryOperation>
456	T inclusive_scan_over_group(const group<Dims>& g, T x, BinaryOperation binary_op) {
457	return sycl::inclusive_scan_over_group(detail::get_sycl_group(g), x, binary_op);
458	}
459
460	using sycl::inclusive_scan_over_group;
461
462	template <int Dims, typename V, typename T, typename BinaryOperation>
463	T inclusive_scan_over_group(const group<Dims>& g, V x, BinaryOperation binary_op, T init) {
464	return sycl::inclusive_scan_over_group(detail::get_sycl_group(g), x, binary_op, init);
465	}
466
467	} // namespace celerity

celerity / celerity-runtime / 11385101423

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