@@ -71,23 +71,26 @@ struct CachedData
7171 template <int _Dims>
7272 void init (const sycl::nd_item<_Dims> &item) const
7373 {
74- uint32_t llid = item.get_local_linear_id ();
74+ std:: uint32_t llid = item.get_local_linear_id ();
7575 auto local_ptr = &local_data[0 ];
76- uint32_t size = local_data.size ();
76+ std:: uint32_t size = local_data.size ();
7777 auto group = item.get_group ();
78- uint32_t local_size = group.get_local_linear_range ();
78+ std:: uint32_t local_size = group.get_local_linear_range ();
7979
80- for (uint32_t i = llid; i < size; i += local_size) {
80+ for (std:: uint32_t i = llid; i < size; i += local_size) {
8181 local_ptr[i] = global_data[i];
8282 }
8383 }
8484
85- size_t size () const { return local_data.size (); }
85+ std::size_t size () const
86+ {
87+ return local_data.size ();
88+ }
8689
8790 T &operator [](const sycl::id<Dims> &id) const { return local_data[id]; }
8891
8992 template <typename = std::enable_if_t <Dims == 1 >>
90- T &operator [](const size_t id) const
93+ T &operator [](const std:: size_t id) const
9194 {
9295 return local_data[id];
9396 }
@@ -119,12 +122,15 @@ struct UncachedData
119122 {
120123 }
121124
122- size_t size () const { return _shape.size (); }
125+ std::size_t size () const
126+ {
127+ return _shape.size ();
128+ }
123129
124130 T &operator [](const sycl::id<Dims> &id) const { return global_data[id]; }
125131
126132 template <typename = std::enable_if_t <Dims == 1 >>
127- T &operator [](const size_t id) const
133+ T &operator [](const std:: size_t id) const
128134 {
129135 return global_data[id];
130136 }
@@ -141,15 +147,15 @@ struct HistLocalType
141147};
142148
143149template <>
144- struct HistLocalType <uint64_t >
150+ struct HistLocalType <std:: uint64_t >
145151{
146- using type = uint32_t ;
152+ using type = std:: uint32_t ;
147153};
148154
149155template <>
150- struct HistLocalType <int64_t >
156+ struct HistLocalType <std:: int64_t >
151157{
152- using type = int32_t ;
158+ using type = std:: int32_t ;
153159};
154160
155161template <typename T, typename localT = typename HistLocalType<T>::type>
@@ -161,8 +167,8 @@ struct HistWithLocalCopies
161167 using LocalHist = sycl::local_accessor<localT, 2 >;
162168
163169 HistWithLocalCopies (T *global_data,
164- size_t bins_count,
165- int32_t copies_count,
170+ std:: size_t bins_count,
171+ std:: int32_t copies_count,
166172 sycl::handler &cgh)
167173 {
168174 local_hist = LocalHist (sycl::range<2 >(copies_count, bins_count), cgh);
@@ -172,23 +178,25 @@ struct HistWithLocalCopies
172178 template <int _Dims>
173179 void init (const sycl::nd_item<_Dims> &item, localT val = 0 ) const
174180 {
175- uint32_t llid = item.get_local_linear_id ();
181+ std:: uint32_t llid = item.get_local_linear_id ();
176182 auto *local_ptr = &local_hist[0 ][0 ];
177- uint32_t size = local_hist.size ();
183+ std:: uint32_t size = local_hist.size ();
178184 auto group = item.get_group ();
179- uint32_t local_size = group.get_local_linear_range ();
185+ std:: uint32_t local_size = group.get_local_linear_range ();
180186
181- for (uint32_t i = llid; i < size; i += local_size) {
187+ for (std:: uint32_t i = llid; i < size; i += local_size) {
182188 local_ptr[i] = val;
183189 }
184190 }
185191
186192 template <int _Dims>
187- void add (const sycl::nd_item<_Dims> &item, int32_t bin, localT value) const
193+ void add (const sycl::nd_item<_Dims> &item,
194+ std::int32_t bin,
195+ localT value) const
188196 {
189- int32_t llid = item.get_local_linear_id ();
190- int32_t local_hist_count = local_hist.get_range ().get (0 );
191- int32_t local_copy_id =
197+ std:: int32_t llid = item.get_local_linear_id ();
198+ std:: int32_t local_hist_count = local_hist.get_range ().get (0 );
199+ std:: int32_t local_copy_id =
192200 local_hist_count == 1 ? 0 : llid % local_hist_count;
193201
194202 AtomicOp<localT, sycl::memory_order::relaxed,
@@ -200,15 +208,15 @@ struct HistWithLocalCopies
200208 template <int _Dims>
201209 void finalize (const sycl::nd_item<_Dims> &item) const
202210 {
203- uint32_t llid = item.get_local_linear_id ();
204- uint32_t bins_count = local_hist.get_range ().get (1 );
205- uint32_t local_hist_count = local_hist.get_range ().get (0 );
211+ std:: uint32_t llid = item.get_local_linear_id ();
212+ std:: uint32_t bins_count = local_hist.get_range ().get (1 );
213+ std:: uint32_t local_hist_count = local_hist.get_range ().get (0 );
206214 auto group = item.get_group ();
207- uint32_t local_size = group.get_local_linear_range ();
215+ std:: uint32_t local_size = group.get_local_linear_range ();
208216
209- for (uint32_t i = llid; i < bins_count; i += local_size) {
217+ for (std:: uint32_t i = llid; i < bins_count; i += local_size) {
210218 auto value = local_hist[0 ][i];
211- for (uint32_t lhc = 1 ; lhc < local_hist_count; ++lhc) {
219+ for (std:: uint32_t lhc = 1 ; lhc < local_hist_count; ++lhc) {
212220 value += local_hist[lhc][i];
213221 }
214222 if (value != T (0 )) {
@@ -219,7 +227,10 @@ struct HistWithLocalCopies
219227 }
220228 }
221229
222- uint32_t size () const { return local_hist.size (); }
230+ std::uint32_t size () const
231+ {
232+ return local_hist.size ();
233+ }
223234
224235private:
225236 LocalHist local_hist;
@@ -240,7 +251,7 @@ struct HistGlobalMemory
240251 }
241252
242253 template <int _Dims>
243- void add (const sycl::nd_item<_Dims> &, int32_t bin, T value) const
254+ void add (const sycl::nd_item<_Dims> &, std:: int32_t bin, T value) const
244255 {
245256 AtomicOp<T, sycl::memory_order::relaxed,
246257 sycl::memory_scope::device>::add (global_hist[bin], value);
@@ -255,18 +266,24 @@ struct HistGlobalMemory
255266 T *global_hist = nullptr ;
256267};
257268
258- template <typename T = uint32_t >
269+ template <typename T = std:: uint32_t >
259270struct NoWeights
260271{
261- constexpr T get (size_t ) const { return 1 ; }
272+ constexpr T get (std::size_t ) const
273+ {
274+ return 1 ;
275+ }
262276};
263277
264278template <typename T>
265279struct Weights
266280{
267281 Weights (T *weights) { data = weights; }
268282
269- T get (size_t id) const { return data[id]; }
283+ T get (std::size_t id) const
284+ {
285+ return data[id];
286+ }
270287
271288private:
272289 T *data = nullptr ;
@@ -281,9 +298,9 @@ bool check_in_bounds(const dT &val, const dT &min, const dT &max)
281298
282299template <typename T, typename HistImpl, typename Edges, typename Weights>
283300void submit_histogram (const T *in,
284- const size_t size,
285- const size_t dims,
286- const uint32_t WorkPI,
301+ const std:: size_t size,
302+ const std:: size_t dims,
303+ const std:: uint32_t WorkPI,
287304 const HistImpl &hist,
288305 const Edges &edges,
289306 const Weights &weights,
@@ -303,8 +320,8 @@ void validate(const usm_ndarray &sample,
303320 const std::optional<const dpctl::tensor::usm_ndarray> &weights,
304321 const usm_ndarray &histogram);
305322
306- uint32_t get_local_hist_copies_count (uint32_t loc_mem_size_in_items,
307- uint32_t local_size,
308- uint32_t hist_size_in_items);
323+ std:: uint32_t get_local_hist_copies_count (std:: uint32_t loc_mem_size_in_items,
324+ std:: uint32_t local_size,
325+ std:: uint32_t hist_size_in_items);
309326
310327} // namespace statistics::histogram
0 commit comments