## On this page
# api\_device\_particles
Last update: 17.07.2025
1\#pragma once
2
3\#include <rocky20/api/rocky\_model\_api.hpp>
4
5namespace device
6{
7
8template <class backend, typename Callable>
9__global__ void execute_particles_function(
10 SDeviceModel* device_data, double current_time,
11 Callable particle_f_ptr,
12 void* udf_data
13)
14{
15 [IRockyDeviceModel](structIRockyDeviceModel.xhtml) device_model_api;
16 device_model_api._impl.model = device_data;
17 device_model_api._impl.current_time = current_time;
18
19 const cuda_general& general = device_data->get_general();
20 const cuda_particle* particles = device_data->prt;
21
22 int thread_start, thread_step, thread_stop;
23 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.nprt_active);
24 for (int i = thread_start; i < thread_stop; i += thread_step)
25 {
26 if (particles[i].enabled())
27 {
28 [IRockyParticle](structIRockyParticle.xhtml) particle_api(i, device_data);
29 particle_f_ptr(device_model_api, particle_api, udf_data);
30 }
31 }
32}
33
34template <class backend, typename Callable>
35__host__ __device__ void compute_instantaneous_breakage(
36 SDeviceModel* device_data, int device_id,
37 Callable process_instantaneous_breakage_ptr, void* udf_data,
38 cu_real time_curr
39)
40{
41 const cuda_general& general = device_data->get_general();
42
43 if (time_curr < general.breakage_start_time || general.start_sim) return;
44
45 int j = 0;
46 auto* scalars = device_data->particle_scalars;
47
48 int* skip_counters = scalars->get_skip_breakage_counters();
49 double* release_times = scalars->get_release_times();
50 auto* breakage_states = scalars->get_breakage_states();
51
52 int thread_start, thread_step, thread_stop;
53 backend::ThreadIteration(thread_start, thread_step, thread_stop, general.nprt_active);
54 for (int i = thread_start; i < thread_stop; i += thread_step)
55 {
56 [IRockyParticle](structIRockyParticle.xhtml) particle_api(i, device_data);
57 auto& particle = particle_api._impl;
58 auto& p_group = device_data->get_particle_group(particle.get_particle_group_index());
59
60 if (skip_counters[i] < general.nbreak_dt_skip) skip_counters[i]++;
61
62 if (p_group.breakage_model == pbmInstantaneousCustom)
63 {
64 if (breakage_states[i].is_breakable()
65 && time_curr >= release_times[i] + general.breakage_start_release_time
66 && skip_counters[i] >= general.nbreak_dt_skip)
67 {
68 bool broken = false;
69 process_instantaneous_breakage_ptr(particle_api, broken, udf_data);
70
71 if (broken)
72 {
73 backend::AtomicAdd(device_data->nbreak, 1, j);
74 device_data->prt_brk[j].i = i;
75 }
76 }
77 }
78 }
79}
80
81} // namespace device
[IRockyDeviceModel](structIRockyDeviceModel.xhtml)
**Definition** rocky_model_api.hpp:19
[IRockyParticle](structIRockyParticle.xhtml)
**Definition** rocky_particle_api.hpp:27