## On this page
# api\_device\_vicinity
Last update: 17.07.2025
1\#pragma once
2
3\#include <rocky20/api/device/api\_device\_grid\_search.hpp>
4\#include <rocky20/api/rocky\_model\_api.hpp>
5\#include <rocky20/api/rocky\_vicinity\_pair\_api.hpp>
6
7struct BaseNeighborhoodOperator
8{
9 SDeviceModel* device_model;
10 const double current_time;
11 const size\_t n_total_particles;
12 const int device_id;
13 void* user_data;
14
15 inline ROCKY_FUNCTIONS size\_t get_near_count() const
16 {
17 return this->n_total_particles;
18 }
19 inline ROCKY_FUNCTIONS double3 get_near_centroid(size\_t id) const
20 {
21 return this->device_model->prt[id].get_centroid();
22 }
23};
24
25template <typename Callable>
26struct ParticlesNeighborhoodOperator : BaseNeighborhoodOperator
27{
28 Callable& kernel;
29
30 inline ROCKY_FUNCTIONS size\_t get_home_count() const
31 {
32 return this->n_total_particles;
33 }
34 inline ROCKY_FUNCTIONS double3 get_home_centroid(size\_t id) const
35 {
36 return this->device_model->prt[id].get_centroid();
37 }
38 inline ROCKY_FUNCTIONS bool is_valid(size\_t home_id, size\_t near_id, bool different_cells)
39 {
40 /\* First, if particles are from same cell, we need to check their IDs to avoid duplicates.
41 \* Then we check if any of them is local to this device, to avoid halo x halo vicinities. \*/
42 return (different_cells || (near_id > home_id))
43 && ((int(this->device_model->prt[home_id].gpu_id) == this->device_id)
44 || (int(this->device_model->prt[near_id].gpu_id) == this->device_id));
45 }
46 inline ROCKY_FUNCTIONS void operate(size\_t home_id, size\_t near_id)
47 {
48 [IRockyDeviceModel](structIRockyDeviceModel.xhtml) model { RockyDeviceModel { this->current_time, this->device_model } };
49 [IRockyVicinityPair](structIRockyVicinityPair.xhtml) vicinity { this->device_model, home_id, near_id, true };
50 this->kernel(model, vicinity, this->user_data);
51 }
52};
53
54template <typename Callable>
55struct TrianglesNeighborhoodOperator : BaseNeighborhoodOperator
56{
57 Callable& kernel;
58
59 inline ROCKY_FUNCTIONS size\_t get_home_count() const
60 {
61 return this->device_model->gen->ntrn_active;
62 }
63 inline ROCKY_FUNCTIONS double3 get_home_centroid(size\_t id) const
64 {
65 return RockyBoundaryTriangle(this->device_model, id).get_centroid();
66 }
67 inline ROCKY_FUNCTIONS bool is_valid(size\_t home_id, size\_t near_id, bool different_cells) const
68 {
69 return near_id < this->device_model->gen->nprt_active;
70 }
71 inline ROCKY_FUNCTIONS void operate(size\_t home_id, size\_t near_id)
72 {
73 [IRockyDeviceModel](structIRockyDeviceModel.xhtml) model { RockyDeviceModel { this->current_time, this->device_model } };
74 [IRockyVicinityPair](structIRockyVicinityPair.xhtml) vicinity { this->device_model, near_id, home_id, false };
75 this->kernel(model, vicinity, this->user_data);
76 }
77};
78
79template <typename Backend, typename Callable>
80__global__ void dispatch_travel_vicinities_particles(
81 Callable kernel, GridSearchParameters params, double current_time, void* user_data)
82{
83 ParticlesNeighborhoodOperator<Callable> neighborhood {
84 params.device_model, current_time, params.n_total_particles,
85 params.device_id, user_data, kernel
86 };
87
88 travel_vicinities_insider<Backend>(
89 neighborhood,
90 params.grid_starting_index,
91 params.grid_to_morton,
92 params.cell_size,
93 params.max_search_radius,
94 params.grid_size,
95 params.grid_start,
96 params.n_vicinities);
97}
98
99template <typename Backend, typename Callable>
100__global__ void dispatch_travel_vicinities_triangles(
101 Callable kernel, GridSearchParameters params, double current_time, void* user_data)
102{
103 TrianglesNeighborhoodOperator<Callable> neighborhood {
104 params.device_model, current_time, params.n_total_particles,
105 params.device_id, user_data, kernel
106 };
107
108 travel_vicinities_outsider<Backend>(
109 neighborhood,
110 params.grid_starting_index,
111 params.grid_to_morton,
112 params.cell_size,
113 params.max_search_radius,
114 params.grid_size,
115 params.grid_start,
116 params.n_vicinities);
117}
[IRockyDeviceModel](structIRockyDeviceModel.xhtml)
**Definition** rocky_model_api.hpp:19
[IRockyVicinityPair](structIRockyVicinityPair.xhtml)
**Definition** rocky_vicinity_pair_api.hpp:12