@@ -74,18 +74,18 @@ struct CachedData
7474 template <int _Dims>
7575 void init (const sycl::nd_item<_Dims> &item) const
7676 {
77- uint32_t llid = item.get_local_linear_id ();
77+ std:: uint32_t llid = item.get_local_linear_id ();
7878 auto local_ptr = &local_data[0 ];
79- uint32_t size = local_data.size ();
79+ std:: uint32_t size = local_data.size ();
8080 auto group = item.get_group ();
81- uint32_t local_size = group.get_local_linear_range ();
81+ std:: uint32_t local_size = group.get_local_linear_range ();
8282
83- for (uint32_t i = llid; i < size; i += local_size) {
83+ for (std:: uint32_t i = llid; i < size; i += local_size) {
8484 local_ptr[i] = global_data[i];
8585 }
8686 }
8787
88- size_t size () const
88+ std:: size_t size () const
8989 {
9090 return local_data.size ();
9191 }
@@ -96,7 +96,7 @@ struct CachedData
9696 }
9797
9898 template <typename = std::enable_if_t <Dims == 1 >>
99- T &operator [](const size_t id) const
99+ T &operator [](const std:: size_t id) const
100100 {
101101 return local_data[id];
102102 }
@@ -131,7 +131,7 @@ struct UncachedData
131131 {
132132 }
133133
134- size_t size () const
134+ std:: size_t size () const
135135 {
136136 return _shape.size ();
137137 }
@@ -142,7 +142,7 @@ struct UncachedData
142142 }
143143
144144 template <typename = std::enable_if_t <Dims == 1 >>
145- T &operator [](const size_t id) const
145+ T &operator [](const std:: size_t id) const
146146 {
147147 return global_data[id];
148148 }
@@ -159,15 +159,15 @@ struct HistLocalType
159159};
160160
161161template <>
162- struct HistLocalType <uint64_t >
162+ struct HistLocalType <std:: uint64_t >
163163{
164- using type = uint32_t ;
164+ using type = std:: uint32_t ;
165165};
166166
167167template <>
168- struct HistLocalType <int64_t >
168+ struct HistLocalType <std:: int64_t >
169169{
170- using type = int32_t ;
170+ using type = std:: int32_t ;
171171};
172172
173173template <typename T, typename localT = typename HistLocalType<T>::type>
@@ -179,8 +179,8 @@ struct HistWithLocalCopies
179179 using LocalHist = sycl::local_accessor<localT, 2 >;
180180
181181 HistWithLocalCopies (T *global_data,
182- size_t bins_count,
183- int32_t copies_count,
182+ std:: size_t bins_count,
183+ std:: int32_t copies_count,
184184 sycl::handler &cgh)
185185 {
186186 local_hist = LocalHist (sycl::range<2 >(copies_count, bins_count), cgh);
@@ -190,23 +190,25 @@ struct HistWithLocalCopies
190190 template <int _Dims>
191191 void init (const sycl::nd_item<_Dims> &item, localT val = 0 ) const
192192 {
193- uint32_t llid = item.get_local_linear_id ();
193+ std:: uint32_t llid = item.get_local_linear_id ();
194194 auto *local_ptr = &local_hist[0 ][0 ];
195- uint32_t size = local_hist.size ();
195+ std:: uint32_t size = local_hist.size ();
196196 auto group = item.get_group ();
197- uint32_t local_size = group.get_local_linear_range ();
197+ std:: uint32_t local_size = group.get_local_linear_range ();
198198
199- for (uint32_t i = llid; i < size; i += local_size) {
199+ for (std:: uint32_t i = llid; i < size; i += local_size) {
200200 local_ptr[i] = val;
201201 }
202202 }
203203
204204 template <int _Dims>
205- void add (const sycl::nd_item<_Dims> &item, int32_t bin, localT value) const
205+ void add (const sycl::nd_item<_Dims> &item,
206+ std::int32_t bin,
207+ localT value) const
206208 {
207- int32_t llid = item.get_local_linear_id ();
208- int32_t local_hist_count = local_hist.get_range ().get (0 );
209- int32_t local_copy_id =
209+ std:: int32_t llid = item.get_local_linear_id ();
210+ std:: int32_t local_hist_count = local_hist.get_range ().get (0 );
211+ std:: int32_t local_copy_id =
210212 local_hist_count == 1 ? 0 : llid % local_hist_count;
211213
212214 AtomicOp<localT, sycl::memory_order::relaxed,
@@ -218,15 +220,15 @@ struct HistWithLocalCopies
218220 template <int _Dims>
219221 void finalize (const sycl::nd_item<_Dims> &item) const
220222 {
221- uint32_t llid = item.get_local_linear_id ();
222- uint32_t bins_count = local_hist.get_range ().get (1 );
223- uint32_t local_hist_count = local_hist.get_range ().get (0 );
223+ std:: uint32_t llid = item.get_local_linear_id ();
224+ std:: uint32_t bins_count = local_hist.get_range ().get (1 );
225+ std:: uint32_t local_hist_count = local_hist.get_range ().get (0 );
224226 auto group = item.get_group ();
225- uint32_t local_size = group.get_local_linear_range ();
227+ std:: uint32_t local_size = group.get_local_linear_range ();
226228
227- for (uint32_t i = llid; i < bins_count; i += local_size) {
229+ for (std:: uint32_t i = llid; i < bins_count; i += local_size) {
228230 auto value = local_hist[0 ][i];
229- for (uint32_t lhc = 1 ; lhc < local_hist_count; ++lhc) {
231+ for (std:: uint32_t lhc = 1 ; lhc < local_hist_count; ++lhc) {
230232 value += local_hist[lhc][i];
231233 }
232234 if (value != T (0 )) {
@@ -237,7 +239,7 @@ struct HistWithLocalCopies
237239 }
238240 }
239241
240- uint32_t size () const
242+ std:: uint32_t size () const
241243 {
242244 return local_hist.size ();
243245 }
@@ -264,7 +266,7 @@ struct HistGlobalMemory
264266 }
265267
266268 template <int _Dims>
267- void add (const sycl::nd_item<_Dims> &, int32_t bin, T value) const
269+ void add (const sycl::nd_item<_Dims> &, std:: int32_t bin, T value) const
268270 {
269271 AtomicOp<T, sycl::memory_order::relaxed,
270272 sycl::memory_scope::device>::add (global_hist[bin], value);
@@ -279,10 +281,10 @@ struct HistGlobalMemory
279281 T *global_hist = nullptr ;
280282};
281283
282- template <typename T = uint32_t >
284+ template <typename T = std:: uint32_t >
283285struct NoWeights
284286{
285- constexpr T get (size_t ) const
287+ constexpr T get (std:: size_t ) const
286288 {
287289 return 1 ;
288290 }
@@ -296,7 +298,7 @@ struct Weights
296298 data = weights;
297299 }
298300
299- T get (size_t id) const
301+ T get (std:: size_t id) const
300302 {
301303 return data[id];
302304 }
@@ -314,9 +316,9 @@ bool check_in_bounds(const dT &val, const dT &min, const dT &max)
314316
315317template <typename T, typename HistImpl, typename Edges, typename Weights>
316318void submit_histogram (const T *in,
317- const size_t size,
318- const size_t dims,
319- const uint32_t WorkPI,
319+ const std:: size_t size,
320+ const std:: size_t dims,
321+ const std:: uint32_t WorkPI,
320322 const HistImpl &hist,
321323 const Edges &edges,
322324 const Weights &weights,
@@ -336,8 +338,8 @@ void validate(const usm_ndarray &sample,
336338 const std::optional<const dpctl::tensor::usm_ndarray> &weights,
337339 const usm_ndarray &histogram);
338340
339- uint32_t get_local_hist_copies_count (uint32_t loc_mem_size_in_items,
340- uint32_t local_size,
341- uint32_t hist_size_in_items);
341+ std:: uint32_t get_local_hist_copies_count (std:: uint32_t loc_mem_size_in_items,
342+ std:: uint32_t local_size,
343+ std:: uint32_t hist_size_in_items);
342344
343345} // namespace statistics::histogram
0 commit comments