api_device_particles

2

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 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 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 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