api_cuda_particles

2

5

6namespace rocky::cuda

7{

8

9template <typename Callable>

10void dispatch_particles_kernel(

11 IRockyModel &model,

12 Callable h_particle_f_ptr,

13 void **udf_devices_data

14)

15{

16 if (model.get_number_of_active_particles() <= 0)

17 return;

18

19 for (int i=0; i<model.get_number_of_devices(); i++)

20 {

21 CUDA_SET_DEVICE(model.get_device_id(i));

22

23 void *udf_device_data = udf_devices_data[i];

24 device::execute_particles_function<CudaBackend><<<CD_NBLOCKS, CD_NTHREADS>>>(

25 model.get_device_data(i), model.get_current_time(),

26 h_particle_f_ptr, udf_device_data

27 );

28 }

29

30 // Sync all GPUs

31 sync_devices(model);

32}

33

34template <typename Callable>

35__global__ void execute_instantaneous_breakage_function_kernel(

36 SDeviceModel* device_data, int device_id,

37 Callable instantaneous_breakage_f_ptr,

38 void* udf_device_data,

39 cu_real time_curr

40)

41{

42 device::compute_instantaneous_breakage<CudaBackend>(

43 device_data, device_id,

44 instantaneous_breakage_f_ptr, udf_device_data,

45 time_curr

46 );

47}

48

49template <typename Callable>

50void dispatch_instantaneous_breakage_kernel(

51 IRockyModel& model,

52 Callable h_instantaneous_breakage_f_ptr,

53 void** udf_devices_data

54)

55{

56 for (int i = 0; i < model.get_number_of_devices(); i++)

57 {

58 void* udf_device_data = udf_devices_data[i];

59 SDeviceModel* device_data = model.get_device_data(i);

60

61 CUDA_SET_DEVICE(model.get_device_id(i));

62 rocky::cuda::execute_instantaneous_breakage_function_kernel<<<CD_NBLOCKS, CD_NTHREADS>>>(

63 device_data,

64 i,

65 h_instantaneous_breakage_f_ptr,

66 udf_device_data,

67 model._impl->iteration.current_time);

68 }

69

70 //Sync all GPUs

71 sync_devices(model);

72}

73

74} // namespace rocky::cuda