@@ -50,9 +50,9 @@ struct EdgesDd
5050 using boundsT = std::tuple<EdgesT *, EdgesT *>;
5151
5252 EdgesDd (const EdgesT *global_edges,
53- size_t total_size,
53+ const size_t total_size,
5454 const EdgesCountT *global_edges_sizes,
55- size_t dims,
55+ const size_t dims,
5656 sycl::handler &cgh)
5757 : edges(global_edges, sycl::range<1 >(total_size), cgh),
5858 edges_size (global_edges_sizes, sycl::range<1 >(dims), cgh)
@@ -72,13 +72,13 @@ struct EdgesDd
7272 sycl::group_barrier (group, sycl::memory_scope::work_group);
7373 }
7474
75- auto *edges_ptr = edges.get_ptr ();
75+ const auto *edges_ptr = edges.get_ptr ();
7676 auto *min_ptr = &min[0 ];
7777 auto *max_ptr = &max[0 ];
7878
7979 if (group.leader ()) {
80- for (size_t i = 0 ; i < edges_size.size (); ++i) {
81- auto size = edges_size[i];
80+ for (uint32_t i = 0 ; i < edges_size.size (); ++i) {
81+ const auto size = edges_size[i];
8282 min_ptr[i] = edges_ptr[0 ];
8383 max_ptr[i] = edges_ptr[size - 1 ];
8484 edges_ptr += size;
@@ -91,11 +91,11 @@ struct EdgesDd
9191 return {&min[0 ], &max[0 ]};
9292 }
9393
94- size_t get_bin_for_dim (const EdgesT &val,
95- const EdgesT *edges_data,
96- uint32_t edges_count) const
94+ auto get_bin_for_dim (const EdgesT &val,
95+ const EdgesT *edges_data,
96+ const uint32_t edges_count) const
9797 {
98- uint32_t bins_count = edges_count - 1 ;
98+ const uint32_t bins_count = edges_count - 1 ;
9999
100100 uint32_t bin = std::upper_bound (edges_data, edges_data + edges_count,
101101 val, Less<EdgesT>{}) -
@@ -106,18 +106,19 @@ struct EdgesDd
106106 }
107107
108108 template <int _Dims, typename dT>
109- size_t get_bin (const sycl::nd_item<_Dims> &,
110- const dT *val,
111- const boundsT &) const
109+ auto get_bin (const sycl::nd_item<_Dims> &,
110+ const dT *val,
111+ const boundsT &) const
112112 {
113- int32_t resulting_bin = 0 ;
114- auto *edges_ptr = &edges[0 ];
115- int32_t dims = edges_size.size ();
113+ uint32_t resulting_bin = 0 ;
114+ const auto *edges_ptr = &edges[0 ];
115+ const uint32_t dims = edges_size.size ();
116116
117- for (int i = 0 ; i < dims; ++i) {
118- int32_t curr_edges_count = edges_size[i];
117+ for (uint32_t i = 0 ; i < dims; ++i) {
118+ const uint32_t curr_edges_count = edges_size[i];
119119
120- auto bin_id = get_bin_for_dim (val[i], edges_ptr, curr_edges_count);
120+ const auto bin_id =
121+ get_bin_for_dim (val[i], edges_ptr, curr_edges_count);
121122 resulting_bin = resulting_bin * (curr_edges_count - 1 ) + bin_id;
122123 edges_ptr += curr_edges_count;
123124 }
@@ -137,12 +138,12 @@ struct EdgesDd
137138 template <typename dT>
138139 bool in_bounds (const dT *val, const boundsT &bounds) const
139140 {
140- EdgesT *min = std::get<0 >(bounds);
141- EdgesT *max = std::get<1 >(bounds);
142- int32_t dims = edges_size.size ();
141+ const EdgesT *min = std::get<0 >(bounds);
142+ const EdgesT *max = std::get<1 >(bounds);
143+ const uint32_t dims = edges_size.size ();
143144
144145 auto in_bounds = true ;
145- for (int i = 0 ; i < dims; ++i) {
146+ for (uint32_t i = 0 ; i < dims; ++i) {
146147 in_bounds &= in_bounds_for_dim (val[i], min[i], max[i]);
147148 }
148149
@@ -198,8 +199,8 @@ struct HistogramddF
198199 return exec_q.submit ([&](sycl::handler &cgh) {
199200 cgh.depends_on (depends);
200201
201- auto dispatch_edges = [&](uint32_t local_mem, const auto &weights ,
202- auto &hist) {
202+ auto dispatch_edges = [&](const uint32_t local_mem ,
203+ const auto &weights, auto &hist) {
203204 if (device.is_gpu () && (local_mem >= total_edges)) {
204205 auto edges = CachedEdgesDd<BinsT, EdgesCountT>(
205206 bins_edges, total_edges, bins_edges_count, dims, cgh);
@@ -226,7 +227,8 @@ struct HistogramddF
226227
227228 auto hist = HistWithLocalCopies<HistType>(
228229 out, bins_count, local_hist_count, cgh);
229- uint32_t free_local_mem = local_mem_size - hist.size ();
230+ const uint32_t free_local_mem =
231+ local_mem_size - hist.size ();
230232
231233 dispatch_edges (free_local_mem, weights, hist);
232234 }
0 commit comments