## On this page

  

 

 # api\_device\_grid\_search

 Last update: 17.07.2025 

<a id="l00001" name="l00001"></a> 1\#pragma once

<a id="l00002" name="l00002"></a> 2 

<a id="l00003" name="l00003"></a> 3\#include &lt;rocky20/core/utils/rocky\_math.hpp&gt;

<a id="l00004" name="l00004"></a> 4 

<a id="l00005" name="l00005"></a> 5\#include &lt;cstddef&gt;

<a id="l00006" name="l00006"></a> 6\#include &lt;vector\_types.h&gt;

<a id="l00007" name="l00007"></a> 7 

<a id="l00008" name="l00008"></a> 8using namespace rocky::vector;

<a id="l00009" name="l00009"></a> 9 

<a id="l00010" name="l00010"></a> 10template &lt;typename NeighborhoodOperator&gt;

<a id="l00011" name="l00011"></a> 11__host__ __device__ size\_t iterate_over_grid_cells(

<a id="l00012" name="l00012"></a> 12 NeighborhoodOperator&amp; neighborhood,

<a id="l00013" name="l00013"></a> 13 const size\_t* grid_starting_index,

<a id="l00014" name="l00014"></a> 14 const size\_t* grid_to_original_map,

<a id="l00015" name="l00015"></a> 15 const longlong3&amp; block_start,

<a id="l00016" name="l00016"></a> 16 const longlong3&amp; block_end,

<a id="l00017" name="l00017"></a> 17 const longlong3&amp; grid_size,

<a id="l00018" name="l00018"></a> 18 const double3&amp; home_position,

<a id="l00019" name="l00019"></a> 19 const size\_t home_entity_id,

<a id="l00020" name="l00020"></a> 20 const double search_radius,

<a id="l00021" name="l00021"></a> 21 const bool not_in_block)

<a id="l00022" name="l00022"></a> 22{

<a id="l00023" name="l00023"></a> 23 const double squared_search_radius = search_radius * search_radius;

<a id="l00024" name="l00024"></a> 24 const size\_t grid_volume = grid_size.x * grid_size.y * grid_size.z;

<a id="l00025" name="l00025"></a> 25 const size\_t near_count = neighborhood.get_near_count();

<a id="l00026" name="l00026"></a> 26 size\_t n_thread_vicinities = 0;

<a id="l00027" name="l00027"></a> 27 for (size\_t x = block_start.x; x &lt; block_end.x; ++x)

<a id="l00028" name="l00028"></a> 28 {

<a id="l00029" name="l00029"></a> 29 const size\_t x_offset = x * grid_size.y * grid_size.z;

<a id="l00030" name="l00030"></a> 30 for (size\_t y = block_start.y; y &lt; block_end.y; ++y)

<a id="l00031" name="l00031"></a> 31 {

<a id="l00032" name="l00032"></a> 32 const size\_t xy_offset = x_offset + (y * grid_size.z);

<a id="l00033" name="l00033"></a> 33 const size\_t first_cell_id = xy_offset + block_start.z;

<a id="l00034" name="l00034"></a> 34 const size\_t last_cell_id = xy_offset + block_end.z;

<a id="l00035" name="l00035"></a> 35 const size\_t index_start = grid_starting_index[first_cell_id];

<a id="l00036" name="l00036"></a> 36 const size\_t index_end =

<a id="l00037" name="l00037"></a> 37 (last_cell_id &lt; grid_volume) ? grid_starting_index[last_cell_id] : near_count;

<a id="l00038" name="l00038"></a> 38 

<a id="l00039" name="l00039"></a> 39 for (size\_t prt_id_grid = index_start; prt_id_grid &lt; index_end; ++prt_id_grid)

<a id="l00040" name="l00040"></a> 40 {

<a id="l00041" name="l00041"></a> 41 const size\_t near_entity_id = grid_to_original_map[prt_id_grid];

<a id="l00042" name="l00042"></a> 42 if (neighborhood.is_valid(home_entity_id, near_entity_id, not_in_block))

<a id="l00043" name="l00043"></a> 43 {

<a id="l00044" name="l00044"></a> 44 const double3 near_position = neighborhood.get_near_centroid(near_entity_id);

<a id="l00045" name="l00045"></a> 45 const double squared_distance = rocky::vector::dot(home_position - near_position);

<a id="l00046" name="l00046"></a> 46 if (squared_distance &lt;= squared_search_radius)

<a id="l00047" name="l00047"></a> 47 {

<a id="l00048" name="l00048"></a> 48 neighborhood.operate(home_entity_id, near_entity_id);

<a id="l00049" name="l00049"></a> 49 n_thread_vicinities++;

<a id="l00050" name="l00050"></a> 50 }

<a id="l00051" name="l00051"></a> 51 }

<a id="l00052" name="l00052"></a> 52 }

<a id="l00053" name="l00053"></a> 53 }

<a id="l00054" name="l00054"></a> 54 }

<a id="l00055" name="l00055"></a> 55 return n_thread_vicinities;

<a id="l00056" name="l00056"></a> 56}

<a id="l00057" name="l00057"></a> 57 

<a id="l00058" name="l00058"></a> 58template &lt;typename NeighborhoodOperator, typename GridStartingIndex&gt;

<a id="l00059" name="l00059"></a> 59__host__ __device__ size\_t iterate_over_grid_cells_2(

<a id="l00060" name="l00060"></a> 60 NeighborhoodOperator&amp; neighborhood,

<a id="l00061" name="l00061"></a> 61 const GridStartingIndex grid_starting_index,

<a id="l00062" name="l00062"></a> 62 const unsigned int* __restrict__ grid_to_original_map,

<a id="l00063" name="l00063"></a> 63 const uint3&amp; block_start,

<a id="l00064" name="l00064"></a> 64 const uint3&amp; block_end,

<a id="l00065" name="l00065"></a> 65 const uint3&amp; grid_size,

<a id="l00066" name="l00066"></a> 66 const unsigned int home_entity_id,

<a id="l00067" name="l00067"></a> 67 const bool not_in_block)

<a id="l00068" name="l00068"></a> 68{

<a id="l00069" name="l00069"></a> 69 const size\_t grid_volume = size_t(grid_size.x) * grid_size.y * grid_size.z;

<a id="l00070" name="l00070"></a> 70 const unsigned int near_count = neighborhood.get_near_count();

<a id="l00071" name="l00071"></a> 71 unsigned int n_thread_vicinities = 0;

<a id="l00072" name="l00072"></a> 72 for (unsigned int x = block_start.x; x &lt; block_end.x; ++x)

<a id="l00073" name="l00073"></a> 73 {

<a id="l00074" name="l00074"></a> 74 const size\_t x_offset = size_t(x) * grid_size.y * grid_size.z;

<a id="l00075" name="l00075"></a> 75 for (unsigned int y = block_start.y; y &lt; block_end.y; ++y)

<a id="l00076" name="l00076"></a> 76 {

<a id="l00077" name="l00077"></a> 77 const size\_t xy_offset = x_offset + (size_t(y) * grid_size.z);

<a id="l00078" name="l00078"></a> 78 const size\_t first_cell_id = xy_offset + block_start.z;

<a id="l00079" name="l00079"></a> 79 const size\_t last_cell_id = xy_offset + block_end.z;

<a id="l00080" name="l00080"></a> 80 unsigned int index_start, index_end;

<a id="l00081" name="l00081"></a> 81 if constexpr (std::is_pointer_v&lt;GridStartingIndex&gt;)

<a id="l00082" name="l00082"></a> 82 {

<a id="l00083" name="l00083"></a> 83 index_start = grid_starting_index[first_cell_id];

<a id="l00084" name="l00084"></a> 84 index_end =

<a id="l00085" name="l00085"></a> 85 (last_cell_id &lt; grid_volume) ? grid_starting_index[last_cell_id] : near_count;

<a id="l00086" name="l00086"></a> 86 }

<a id="l00087" name="l00087"></a> 87 else

<a id="l00088" name="l00088"></a> 88 {

<a id="l00089" name="l00089"></a> 89 grid_starting_index.get_range(

<a id="l00090" name="l00090"></a> 90 first_cell_id, last_cell_id, index_start, index_end, grid_volume);

<a id="l00091" name="l00091"></a> 91 }

<a id="l00092" name="l00092"></a> 92 

<a id="l00093" name="l00093"></a> 93 for (unsigned int prt_id_grid = index_start; prt_id_grid &lt; index_end; ++prt_id_grid)

<a id="l00094" name="l00094"></a> 94 {

<a id="l00095" name="l00095"></a> 95 const unsigned int near_entity_id = grid_to_original_map[prt_id_grid];

<a id="l00096" name="l00096"></a> 96 if (neighborhood.is_valid(home_entity_id, near_entity_id, not_in_block))

<a id="l00097" name="l00097"></a> 97 {

<a id="l00098" name="l00098"></a> 98 neighborhood.operate(home_entity_id, near_entity_id);

<a id="l00099" name="l00099"></a> 99 n_thread_vicinities++;

<a id="l00100" name="l00100"></a> 100 }

<a id="l00101" name="l00101"></a> 101 }

<a id="l00102" name="l00102"></a> 102 }

<a id="l00103" name="l00103"></a> 103 }

<a id="l00104" name="l00104"></a> 104 return n_thread_vicinities;

<a id="l00105" name="l00105"></a> 105}

<a id="l00106" name="l00106"></a> 106 

<a id="l00123" name="l00123"></a> 123template &lt;typename Backend, typename NeighborhoodOperator&gt;

<a id="l00124" name="l00124"></a> 124__host__ __device__ void travel_vicinities_outsider(

<a id="l00125" name="l00125"></a> 125 NeighborhoodOperator&amp; neighborhood,

<a id="l00126" name="l00126"></a> 126 const size\_t* grid_starting_index,

<a id="l00127" name="l00127"></a> 127 const size\_t* grid_to_original_map,

<a id="l00128" name="l00128"></a> 128 const double cell_size,

<a id="l00129" name="l00129"></a> 129 const double search_radius,

<a id="l00130" name="l00130"></a> 130 const longlong3 grid_size,

<a id="l00131" name="l00131"></a> 131 double3 grid_start,

<a id="l00132" name="l00132"></a> 132 size\_t* n_vicinities)

<a id="l00133" name="l00133"></a> 133{

<a id="l00134" name="l00134"></a> 134 auto clamp_to_grid = [&amp;grid_size](const longlong3&amp; position) -&gt; longlong3 {

<a id="l00135" name="l00135"></a> 135 constexpr longlong3 origin { 0, 0, 0 };

<a id="l00136" name="l00136"></a> 136 return clamp_vector(position, origin, grid_size);

<a id="l00137" name="l00137"></a> 137 };

<a id="l00138" name="l00138"></a> 138 

<a id="l00139" name="l00139"></a> 139 // reduce grid start by an epsilon to avoid summing it on every position calculation

<a id="l00140" name="l00140"></a> 140 grid_start -= (cell_size * 0.01);

<a id="l00141" name="l00141"></a> 141 const double cell_size_inverse = 1.0 / cell_size;

<a id="l00142" name="l00142"></a> 142 size\_t n_thread_vicinities = 0;

<a id="l00143" name="l00143"></a> 143 int start, step, stop;

<a id="l00144" name="l00144"></a> 144 

<a id="l00145" name="l00145"></a> 145 Backend::ThreadIteration(start, step, stop, neighborhood.get_home_count());

<a id="l00146" name="l00146"></a> 146 for (size\_t home_entity_id = start; home_entity_id &lt; stop; home_entity_id += step)

<a id="l00147" name="l00147"></a> 147 {

<a id="l00148" name="l00148"></a> 148 const double3 home_position = neighborhood.get_home_centroid(home_entity_id);

<a id="l00149" name="l00149"></a> 149 const auto home_grid_position { rocky::vector::cast&lt;longlong3&gt;(

<a id="l00150" name="l00150"></a> 150 (home_position - grid_start) * cell_size_inverse) };

<a id="l00151" name="l00151"></a> 151 

<a id="l00152" name="l00152"></a> 152 const longlong3 block_start = clamp_to_grid(home_grid_position - 2);

<a id="l00153" name="l00153"></a> 153 const longlong3 block_end = clamp_to_grid(home_grid_position + 3);

<a id="l00154" name="l00154"></a> 154 

<a id="l00155" name="l00155"></a> 155 n_thread_vicinities += iterate_over_grid_cells(

<a id="l00156" name="l00156"></a> 156 neighborhood,

<a id="l00157" name="l00157"></a> 157 grid_starting_index,

<a id="l00158" name="l00158"></a> 158 grid_to_original_map,

<a id="l00159" name="l00159"></a> 159 block_start,

<a id="l00160" name="l00160"></a> 160 block_end,

<a id="l00161" name="l00161"></a> 161 grid_size,

<a id="l00162" name="l00162"></a> 162 home_position,

<a id="l00163" name="l00163"></a> 163 home_entity_id,

<a id="l00164" name="l00164"></a> 164 search_radius,

<a id="l00165" name="l00165"></a> 165 false);

<a id="l00166" name="l00166"></a> 166 }

<a id="l00167" name="l00167"></a> 167 if (n_vicinities != nullptr &amp;&amp; n_thread_vicinities &gt; 0)

<a id="l00168" name="l00168"></a> 168 {

<a id="l00169" name="l00169"></a> 169 Backend::AtomicAdd(n_vicinities, n_thread_vicinities);

<a id="l00170" name="l00170"></a> 170 }

<a id="l00171" name="l00171"></a> 171}

<a id="l00172" name="l00172"></a> 172 

<a id="l00189" name="l00189"></a> 189template &lt;typename Backend, typename NeighborhoodOperator&gt;

<a id="l00190" name="l00190"></a> 190__host__ __device__ void travel_vicinities_insider(

<a id="l00191" name="l00191"></a> 191 NeighborhoodOperator&amp; neighborhood,

<a id="l00192" name="l00192"></a> 192 const size\_t* grid_starting_index,

<a id="l00193" name="l00193"></a> 193 const size\_t* grid_to_original_map,

<a id="l00194" name="l00194"></a> 194 const double cell_size,

<a id="l00195" name="l00195"></a> 195 const double search_radius,

<a id="l00196" name="l00196"></a> 196 const longlong3 grid_size,

<a id="l00197" name="l00197"></a> 197 double3 grid_start,

<a id="l00198" name="l00198"></a> 198 size\_t* n_vicinities)

<a id="l00199" name="l00199"></a> 199{

<a id="l00200" name="l00200"></a> 200 // reduce grid start by an epsilon to avoid summing it on every position calculation

<a id="l00201" name="l00201"></a> 201 grid_start -= (cell_size * 0.01);

<a id="l00202" name="l00202"></a> 202 const double cell_size_inverse = 1.0 / cell_size;

<a id="l00203" name="l00203"></a> 203 size\_t n_thread_vicinities = 0;

<a id="l00204" name="l00204"></a> 204 int start, step, stop;

<a id="l00205" name="l00205"></a> 205 

<a id="l00206" name="l00206"></a> 206 double3 home_position;

<a id="l00207" name="l00207"></a> 207 size\_t home_entity_id;

<a id="l00208" name="l00208"></a> 208 auto dispatch_on_block = [&amp;](const longlong3&amp; _block_start,

<a id="l00209" name="l00209"></a> 209 const longlong3&amp; _block_end,

<a id="l00210" name="l00210"></a> 210 const bool different_cells = true) -&gt; void {

<a id="l00211" name="l00211"></a> 211 n_thread_vicinities += iterate_over_grid_cells(

<a id="l00212" name="l00212"></a> 212 neighborhood,

<a id="l00213" name="l00213"></a> 213 grid_starting_index,

<a id="l00214" name="l00214"></a> 214 grid_to_original_map,

<a id="l00215" name="l00215"></a> 215 _block_start,

<a id="l00216" name="l00216"></a> 216 _block_end,

<a id="l00217" name="l00217"></a> 217 grid_size,

<a id="l00218" name="l00218"></a> 218 home_position,

<a id="l00219" name="l00219"></a> 219 home_entity_id,

<a id="l00220" name="l00220"></a> 220 search_radius,

<a id="l00221" name="l00221"></a> 221 different_cells);

<a id="l00222" name="l00222"></a> 222 };

<a id="l00223" name="l00223"></a> 223 

<a id="l00224" name="l00224"></a> 224 Backend::ThreadIteration(start, step, stop, neighborhood.get_home_count());

<a id="l00225" name="l00225"></a> 225 for (size\_t prt_id_grid = start; prt_id_grid &lt; stop; prt_id_grid += step)

<a id="l00226" name="l00226"></a> 226 {

<a id="l00227" name="l00227"></a> 227 home_entity_id = grid_to_original_map[prt_id_grid];

<a id="l00228" name="l00228"></a> 228 if (home_entity_id == INVALID_ENTITY_ID)

<a id="l00229" name="l00229"></a> 229 {

<a id="l00230" name="l00230"></a> 230 continue;

<a id="l00231" name="l00231"></a> 231 }

<a id="l00232" name="l00232"></a> 232 

<a id="l00233" name="l00233"></a> 233 home_position = neighborhood.get_home_centroid(home_entity_id);

<a id="l00234" name="l00234"></a> 234 const auto home_coordinate { rocky::vector::cast&lt;longlong3&gt;(

<a id="l00235" name="l00235"></a> 235 (home_position - grid_start) * cell_size_inverse) };

<a id="l00236" name="l00236"></a> 236 

<a id="l00237" name="l00237"></a> 237 const longlong3 block_start { max_vector({ 0ll, 0ll, 0ll }, home_coordinate - 2) };

<a id="l00238" name="l00238"></a> 238 const longlong3 block_diag { min_vector(grid_size, home_coordinate + 1) };

<a id="l00239" name="l00239"></a> 239 const longlong3 block_end { min_vector(grid_size, home_coordinate + 3) };

<a id="l00240" name="l00240"></a> 240 

<a id="l00241" name="l00241"></a> 241 // xyz block above

<a id="l00242" name="l00242"></a> 242 dispatch_on_block({ block_diag.x, block_start.y, block_start.z }, block_end);

<a id="l00243" name="l00243"></a> 243 // right side of yz plane

<a id="l00244" name="l00244"></a> 244 dispatch_on_block(

<a id="l00245" name="l00245"></a> 245 { home_coordinate.x, block_diag.y, block_start.z },

<a id="l00246" name="l00246"></a> 246 { block_diag.x, block_end.y, block_end.z });

<a id="l00247" name="l00247"></a> 247 // all the way back to z end on this line

<a id="l00248" name="l00248"></a> 248 dispatch_on_block(

<a id="l00249" name="l00249"></a> 249 { home_coordinate.x, home_coordinate.y, block_diag.z },

<a id="l00250" name="l00250"></a> 250 { block_diag.x, block_diag.y, block_end.z });

<a id="l00251" name="l00251"></a> 251 // same grid cell. Can't skip element ID check to avoid duplicate vicinities.

<a id="l00252" name="l00252"></a> 252 dispatch_on_block(home_coordinate, block_diag, false);

<a id="l00253" name="l00253"></a> 253 }

<a id="l00254" name="l00254"></a> 254 if (n_vicinities != nullptr &amp;&amp; n_thread_vicinities &gt; 0)

<a id="l00255" name="l00255"></a> 255 {

<a id="l00256" name="l00256"></a> 256 Backend::AtomicAdd(n_vicinities, n_thread_vicinities);

<a id="l00257" name="l00257"></a> 257 }

<a id="l00258" name="l00258"></a> 258}

<a id="l00259" name="l00259"></a> 259 

<a id="l00260" name="l00260"></a> 260template &lt;typename Backend, typename NeighborhoodOperator, typename GridStartingIndex&gt;

<a id="l00261" name="l00261"></a> 261__host__ __device__ unsigned int travel_vicinities_insider_full(

<a id="l00262" name="l00262"></a> 262 NeighborhoodOperator&amp; neighborhood,

<a id="l00263" name="l00263"></a> 263 const GridStartingIndex grid_starting_index,

<a id="l00264" name="l00264"></a> 264 const unsigned int* __restrict__ grid_to_original_map,

<a id="l00265" name="l00265"></a> 265 const float cell_size,

<a id="l00266" name="l00266"></a> 266 const float search_radius,

<a id="l00267" name="l00267"></a> 267 const uint3&amp; grid_size,

<a id="l00268" name="l00268"></a> 268 float3 grid_start)

<a id="l00269" name="l00269"></a> 269{

<a id="l00270" name="l00270"></a> 270 const float epsilon = cell_size * 0.01f;

<a id="l00271" name="l00271"></a> 271 const float cell_size_inverse = 1.f / cell_size;

<a id="l00272" name="l00272"></a> 272 const unsigned int steps = std::ceil(search_radius / cell_size);

<a id="l00273" name="l00273"></a> 273 unsigned int n_thread_vicinities = 0;

<a id="l00274" name="l00274"></a> 274 int start, step, stop;

<a id="l00275" name="l00275"></a> 275 

<a id="l00276" name="l00276"></a> 276 unsigned int home_entity_id;

<a id="l00277" name="l00277"></a> 277 auto dispatch_on_block = [&amp;](const uint3&amp; _block_start,

<a id="l00278" name="l00278"></a> 278 const uint3&amp; _block_end,

<a id="l00279" name="l00279"></a> 279 const bool different_cells = true) -&gt; void {

<a id="l00280" name="l00280"></a> 280 n_thread_vicinities += iterate_over_grid_cells_2(

<a id="l00281" name="l00281"></a> 281 neighborhood,

<a id="l00282" name="l00282"></a> 282 grid_starting_index,

<a id="l00283" name="l00283"></a> 283 grid_to_original_map,

<a id="l00284" name="l00284"></a> 284 _block_start,

<a id="l00285" name="l00285"></a> 285 _block_end,

<a id="l00286" name="l00286"></a> 286 grid_size,

<a id="l00287" name="l00287"></a> 287 home_entity_id,

<a id="l00288" name="l00288"></a> 288 different_cells);

<a id="l00289" name="l00289"></a> 289 };

<a id="l00290" name="l00290"></a> 290 

<a id="l00291" name="l00291"></a> 291 Backend::ThreadIteration(start, step, stop, neighborhood.get_home_count());

<a id="l00292" name="l00292"></a> 292 for (home_entity_id = start; home_entity_id &lt; stop; home_entity_id += step)

<a id="l00293" name="l00293"></a> 293 {

<a id="l00294" name="l00294"></a> 294 neighborhood.start_element(home_entity_id);

<a id="l00295" name="l00295"></a> 295 

<a id="l00296" name="l00296"></a> 296 const auto home_coordinate = rocky::vector::cast&lt;uint3&gt;(

<a id="l00297" name="l00297"></a> 297 ((neighborhood.get_home_centroid(home_entity_id) - grid_start) + epsilon)

<a id="l00298" name="l00298"></a> 298 * cell_size_inverse);

<a id="l00299" name="l00299"></a> 299 

<a id="l00300" name="l00300"></a> 300 const uint3 block_start { max_vector(uint3 { steps, steps, steps }, home_coordinate) - steps };

<a id="l00301" name="l00301"></a> 301 const uint3 block_diag { min_vector(grid_size, home_coordinate + 1) };

<a id="l00302" name="l00302"></a> 302 const uint3 block_end { min_vector(grid_size, home_coordinate + (steps + 1)) };

<a id="l00303" name="l00303"></a> 303 

<a id="l00304" name="l00304"></a> 304 // xyz block above

<a id="l00305" name="l00305"></a> 305 dispatch_on_block({ block_diag.x, block_start.y, block_start.z }, block_end);

<a id="l00306" name="l00306"></a> 306 // right side of yz plane

<a id="l00307" name="l00307"></a> 307 dispatch_on_block(

<a id="l00308" name="l00308"></a> 308 { home_coordinate.x, block_diag.y, block_start.z },

<a id="l00309" name="l00309"></a> 309 { block_diag.x, block_end.y, block_end.z });

<a id="l00310" name="l00310"></a> 310 // all the way back to z end on this line

<a id="l00311" name="l00311"></a> 311 dispatch_on_block(

<a id="l00312" name="l00312"></a> 312 { home_coordinate.x, home_coordinate.y, block_diag.z },

<a id="l00313" name="l00313"></a> 313 { block_diag.x, block_diag.y, block_end.z });

<a id="l00314" name="l00314"></a> 314 // same grid cell. Can't skip element ID check to avoid duplicate vicinities.

<a id="l00315" name="l00315"></a> 315 dispatch_on_block(home_coordinate, block_diag, false);

<a id="l00316" name="l00316"></a> 316 

<a id="l00317" name="l00317"></a> 317 neighborhood.finish_element(home_entity_id);

<a id="l00318" name="l00318"></a> 318 }

<a id="l00319" name="l00319"></a> 319 return n_thread_vicinities;

<a id="l00320" name="l00320"></a> 320}