Skip to content

Commit c0ba767

Browse files
committed
Use dpctl::tensor scope for ssize_t
1 parent a7c817f commit c0ba767

File tree

1 file changed

+28
-27
lines changed

1 file changed

+28
-27
lines changed

dpnp/backend/extensions/indexing/choose_kernel.hpp

Lines changed: 28 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -43,31 +43,32 @@ namespace dpnp::extensions::indexing::strides_detail
4343
struct NthStrideOffsetUnpacked
4444
{
4545
NthStrideOffsetUnpacked(int common_nd,
46-
ssize_t const *_offsets,
47-
ssize_t const *_shape,
48-
ssize_t const *_strides)
46+
dpctl::tensor::ssize_t const *_offsets,
47+
dpctl::tensor::ssize_t const *_shape,
48+
dpctl::tensor::ssize_t const *_strides)
4949
: _ind(common_nd), nd(common_nd), offsets(_offsets), shape(_shape),
5050
strides(_strides)
5151
{
5252
}
5353

5454
template <typename nT>
55-
size_t operator()(ssize_t gid, nT n) const
55+
size_t operator()(dpctl::tensor::ssize_t gid, nT n) const
5656
{
57-
ssize_t relative_offset(0);
58-
_ind.get_displacement<const ssize_t *, const ssize_t *>(
57+
dpctl::tensor::ssize_t relative_offset(0);
58+
_ind.get_displacement<const dpctl::tensor::ssize_t *,
59+
const dpctl::tensor::ssize_t *>(
5960
gid, shape, strides + (n * nd), relative_offset);
6061

6162
return relative_offset + offsets[n];
6263
}
6364

6465
private:
65-
dpctl::tensor::strides::CIndexer_vector<ssize_t> _ind;
66+
dpctl::tensor::strides::CIndexer_vector<dpctl::tensor::ssize_t> _ind;
6667

6768
int nd;
68-
ssize_t const *offsets;
69-
ssize_t const *shape;
70-
ssize_t const *strides;
69+
dpctl::tensor::ssize_t const *offsets;
70+
dpctl::tensor::ssize_t const *shape;
71+
dpctl::tensor::ssize_t const *strides;
7172
};
7273

7374
} // namespace dpnp::extensions::indexing::strides_detail
@@ -93,15 +94,15 @@ class ChooseFunctor
9394
const IndT *ind = nullptr;
9495
T *dst = nullptr;
9596
char **chcs = nullptr;
96-
ssize_t n_chcs;
97+
dpctl::tensor::ssize_t n_chcs;
9798
const IndOutIndexerT ind_out_indexer;
9899
const ChoicesIndexerT chcs_indexer;
99100

100101
public:
101102
ChooseFunctor(const IndT *ind_,
102103
T *dst_,
103104
char **chcs_,
104-
ssize_t n_chcs_,
105+
dpctl::tensor::ssize_t n_chcs_,
105106
const IndOutIndexerT &ind_out_indexer_,
106107
const ChoicesIndexerT &chcs_indexer_)
107108
: ind(ind_), dst(dst_), chcs(chcs_), n_chcs(n_chcs_),
@@ -113,17 +114,17 @@ class ChooseFunctor
113114
{
114115
const ProjectorT proj{};
115116

116-
ssize_t i = id[0];
117+
dpctl::tensor::ssize_t i = id[0];
117118

118119
auto ind_dst_offsets = ind_out_indexer(i);
119-
ssize_t ind_offset = ind_dst_offsets.get_first_offset();
120-
ssize_t dst_offset = ind_dst_offsets.get_second_offset();
120+
dpctl::tensor::ssize_t ind_offset = ind_dst_offsets.get_first_offset();
121+
dpctl::tensor::ssize_t dst_offset = ind_dst_offsets.get_second_offset();
121122

122123
IndT chc_idx = ind[ind_offset];
123124
// proj produces an index in the range of n_chcs
124-
ssize_t projected_idx = proj(n_chcs, chc_idx);
125+
dpctl::tensor::ssize_t projected_idx = proj(n_chcs, chc_idx);
125126

126-
ssize_t chc_offset = chcs_indexer(i, projected_idx);
127+
dpctl::tensor::ssize_t chc_offset = chcs_indexer(i, projected_idx);
127128

128129
T *chc = reinterpret_cast<T *>(chcs[projected_idx]);
129130

@@ -133,29 +134,29 @@ class ChooseFunctor
133134

134135
typedef sycl::event (*choose_fn_ptr_t)(sycl::queue &,
135136
size_t,
136-
ssize_t,
137+
dpctl::tensor::ssize_t,
137138
int,
138-
const ssize_t *,
139+
const dpctl::tensor::ssize_t *,
139140
const char *,
140141
char *,
141142
char **,
142-
ssize_t,
143-
ssize_t,
144-
const ssize_t *,
143+
dpctl::tensor::ssize_t,
144+
dpctl::tensor::ssize_t,
145+
const dpctl::tensor::ssize_t *,
145146
const std::vector<sycl::event> &);
146147

147148
template <typename ProjectorT, typename indTy, typename Ty>
148149
sycl::event choose_impl(sycl::queue &q,
149150
size_t nelems,
150-
ssize_t n_chcs,
151+
dpctl::tensor::ssize_t n_chcs,
151152
int nd,
152-
const ssize_t *shape_and_strides,
153+
const dpctl::tensor::ssize_t *shape_and_strides,
153154
const char *ind_cp,
154155
char *dst_cp,
155156
char **chcs_cp,
156-
ssize_t ind_offset,
157-
ssize_t dst_offset,
158-
const ssize_t *chc_offsets,
157+
dpctl::tensor::ssize_t ind_offset,
158+
dpctl::tensor::ssize_t dst_offset,
159+
const dpctl::tensor::ssize_t *chc_offsets,
159160
const std::vector<sycl::event> &depends)
160161
{
161162
dpctl::tensor::type_utils::validate_type_for_device<Ty>(q);

0 commit comments

Comments
 (0)