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 model { RockyDeviceModel { this->current_time, this->device_model } }; 49 IRockyVicinityPair 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 model { RockyDeviceModel { this->current_time, this->device_model } }; 74 IRockyVicinityPair 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} IRockyDeviceModelDefinition rocky_model_api.hpp:19 IRockyVicinityPairDefinition rocky_vicinity_pair_api.hpp:12
