From 660ef263bedff4ae75d5759423a1262b713a7698 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 16:29:17 +0200 Subject: [PATCH 01/15] #22 Refactor: move header to include --- {src => include}/utils/cuda_basic_functions.cuh | 0 src/utils/cuda_basic_functions.cu | 2 +- 2 files changed, 1 insertion(+), 1 deletion(-) rename {src => include}/utils/cuda_basic_functions.cuh (100%) diff --git a/src/utils/cuda_basic_functions.cuh b/include/utils/cuda_basic_functions.cuh similarity index 100% rename from src/utils/cuda_basic_functions.cuh rename to include/utils/cuda_basic_functions.cuh diff --git a/src/utils/cuda_basic_functions.cu b/src/utils/cuda_basic_functions.cu index 943b245..d8a1660 100644 --- a/src/utils/cuda_basic_functions.cu +++ b/src/utils/cuda_basic_functions.cu @@ -1,4 +1,4 @@ -#include "cuda_basic_functions.cuh" +#include "utils/cuda_basic_functions.cuh" inline void cErrorCheck(const char *file, int line){ cudaThreadSynchronize(); From 9c737d251461676c568fb51a0e7368c62509debb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 17:06:01 +0200 Subject: [PATCH 02/15] #162 Add preproc data folder --- .gitignore | 5 +++++ preproc/renate_od/data/.gitkeep | 1 + 2 files changed, 6 insertions(+) create mode 100644 preproc/renate_od/data/.gitkeep diff --git a/.gitignore b/.gitignore index 633e6dd..9743456 100644 --- a/.gitignore +++ b/.gitignore @@ -2,6 +2,11 @@ **/.idea/ **/__pycache__ +#preproc data folder +preproc/renate_od/data/* +!preproc/renate_od/data/.gitkeep + + # result folder results/* !results/.gitkeep diff --git a/preproc/renate_od/data/.gitkeep b/preproc/renate_od/data/.gitkeep new file mode 100644 index 0000000..8b13789 --- /dev/null +++ b/preproc/renate_od/data/.gitkeep @@ -0,0 +1 @@ + From 3915faf7de04c68fbe771dc40e706c97c0587dfc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 17:19:08 +0200 Subject: [PATCH 03/15] #22 Add CUDA memory free function --- include/init/device/init.cuh | 1 + src/init/device/init.cu | 7 +++++++ 2 files changed, 8 insertions(+) diff --git a/include/init/device/init.cuh b/include/init/device/init.cuh index 8f95822..931f902 100644 --- a/include/init/device/init.cuh +++ b/include/init/device/init.cuh @@ -3,6 +3,7 @@ void init_host(TaigaGlobals *host_global, TaigaCommons *host_common); void init_grid(ShotProp shot, RunProp run, TaigaCommons *host_common, TaigaCommons *shared_common); +void free_grid(TaigaCommons *shared_common); void init_device_structs(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals *shared_global, TaigaCommons *shared_common); void set_particle_number(RunProp *run, TaigaGlobals *host_global, TaigaGlobals *shared_global); diff --git a/src/init/device/init.cu b/src/init/device/init.cu index 9659ddd..3aa9773 100644 --- a/src/init/device/init.cu +++ b/src/init/device/init.cu @@ -54,6 +54,13 @@ void init_grid(ShotProp shot, RunProp run, TaigaCommons *host_common, TaigaCommo } } +void free_grid(TaigaCommons *shared_common){ + cudaFree(shared_common->grid_size); + cudaFree(shared_common->spline_rgrid); + cudaFree(shared_common->spline_zgrid); +} + + void init_device_structs(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals *shared_global, TaigaCommons *shared_common){ shared_global->particle_number = run.particle_number; shared_common->max_step_number = run.step_device; From 84903c838c1386a4b9b6175b203f0a6ce522f136 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 17:47:29 +0200 Subject: [PATCH 04/15] #22 Add error management, move cuda functions --- include/utils/cuda.cuh | 14 ++++++++++++++ include/utils/cuda_basic_functions.cuh | 7 ------- src/utils/{cuda_basic_functions.cu => cuda.cu} | 12 +----------- 3 files changed, 15 insertions(+), 18 deletions(-) create mode 100644 include/utils/cuda.cuh delete mode 100644 include/utils/cuda_basic_functions.cuh rename src/utils/{cuda_basic_functions.cu => cuda.cu} (83%) diff --git a/include/utils/cuda.cuh b/include/utils/cuda.cuh new file mode 100644 index 0000000..155b9ee --- /dev/null +++ b/include/utils/cuda.cuh @@ -0,0 +1,14 @@ +#ifndef TAIGA_CUDA_CUH +#define TAIGA_CUDA_CUH + +#define CHECK_ERROR(cuda_code) { manage_cuda_error((cuda_code), __FILE__, __LINE__); } + +void set_cuda(int debug_flag); + +inline void manage_cuda_error(cudaError_t cuda_code, const char *filename, int line, bool abort=true) { + if (cuda_code != cudaSuccess) { + fprintf(stderr,"CUDA ERROR in %s (line %d):\n %s\n", filename, line, cudaGetErrorString(code)); + } +} + +#endif //TAIGA_CUDA_CUH diff --git a/include/utils/cuda_basic_functions.cuh b/include/utils/cuda_basic_functions.cuh deleted file mode 100644 index d1eea31..0000000 --- a/include/utils/cuda_basic_functions.cuh +++ /dev/null @@ -1,7 +0,0 @@ -#ifndef CUDA_BASIC_FUNCTIONS_CUH -#define CUDA_BASIC_FUNCTIONS_CUH - -void set_cuda(int debug_flag); -inline void cErrorCheck(const char *file, int line); - -#endif //CUDA_BASIC_FUNCTIONS_CUH diff --git a/src/utils/cuda_basic_functions.cu b/src/utils/cuda.cu similarity index 83% rename from src/utils/cuda_basic_functions.cu rename to src/utils/cuda.cu index d8a1660..c487e87 100644 --- a/src/utils/cuda_basic_functions.cu +++ b/src/utils/cuda.cu @@ -1,14 +1,4 @@ -#include "utils/cuda_basic_functions.cuh" - -inline void cErrorCheck(const char *file, int line){ - cudaThreadSynchronize(); - cudaError_t err = cudaGetLastError(); - if (err != cudaSuccess){ - printf("Error: %s\n", cudaGetErrorString(err)); - printf(" @ %s: %d\n", file, line); - exit(1); - } -} +#include "utils/cuda.cuh" void set_cuda(int debug_flag){ int num_devices, device_i, active_device=0; From 9a2d0bb33522d50755a089ca2831756c44b8fcc7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 17:47:52 +0200 Subject: [PATCH 05/15] #22 Add error management, move cuda functions --- src/main.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/main.cu b/src/main.cu index 0d3e1b0..fbc7ce7 100644 --- a/src/main.cu +++ b/src/main.cu @@ -15,7 +15,7 @@ #include "utils/debug_functions.c" #include "utils/basic_functions.h" #include "utils/dataio/dir_functions.c" -#include "utils/cuda_basic_functions.cuh" +#include "utils/cuda.cuh" #include "utils/physics.c" #include "utils/dataio/data_import.c" From c9bb92f043a7a9b4f797ec1a1a1e24b2da801af6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 17:48:16 +0200 Subject: [PATCH 06/15] #22 Add CUDA memory free function --- src/init/device/init.cu | 2 +- src/init/thomson.cu | 6 ++++++ 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/src/init/device/init.cu b/src/init/device/init.cu index 3aa9773..17c070a 100644 --- a/src/init/device/init.cu +++ b/src/init/device/init.cu @@ -3,7 +3,7 @@ #include "utils/taiga_constants.h" #include "utils/physics.h" #include "utils/basic_functions.c" -#include "utils/cuda_basic_functions.cu" +#include "utils/cuda.cu" void init_host(TaigaGlobals *host_global, TaigaCommons *host_common){ host_common->max_step_number = 0; diff --git a/src/init/thomson.cu b/src/init/thomson.cu index a2cf085..9805475 100644 --- a/src/init/thomson.cu +++ b/src/init/thomson.cu @@ -24,3 +24,9 @@ void set_thomson_profiles(ShotProp shot, TaigaCommons *host_common, TaigaCommons shared_common->ts_density = shared_ts_density; shared_common->ts_temperature = shared_ts_temperature; } + +void free_thomson_profiles(TaigaCommons *shared_common) { + cudaFree(shared_common->ts_psi); + cudaFree(shared_common->ts_density); + cudaFree(shared_common->ts_temperature); +} From 5e15ed73c43c965a243a37e3a1d0db4235670ee3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 17:50:19 +0200 Subject: [PATCH 07/15] #22 Add error management --- src/init/device/init.cu | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/src/init/device/init.cu b/src/init/device/init.cu index 17c070a..e0b30e8 100644 --- a/src/init/device/init.cu +++ b/src/init/device/init.cu @@ -36,17 +36,17 @@ void init_grid(ShotProp shot, RunProp run, TaigaCommons *host_common, TaigaCommo size_t size_Z = host_common->grid_size[1] * sizeof(double); memcpy(shared_common, host_common, size_commons); - - cudaMalloc((void **) &shared_grid_size, size_grid_dim); - cudaMemcpy(shared_grid_size, host_common->grid_size, size_grid_dim, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMalloc((void **) &shared_grid_size, size_grid_dim)); + CHECK_ERROR(cudaMemcpy(shared_grid_size, host_common->grid_size, size_grid_dim, cudaMemcpyHostToDevice)); shared_common->grid_size = shared_grid_size; - - cudaMalloc((void **) &shared_rgrid, size_R); - cudaMemcpy(shared_rgrid, host_common->spline_rgrid, size_R, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMalloc((void **) &shared_rgrid, size_R)); + CHECK_ERROR(cudaMemcpy(shared_rgrid, host_common->spline_rgrid, size_R, cudaMemcpyHostToDevice)); shared_common->spline_rgrid = shared_rgrid; - - cudaMalloc((void **) &shared_zgrid, size_Z); - cudaMemcpy(shared_zgrid, host_common->spline_zgrid, size_Z, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMalloc((void **) &shared_zgrid, size_Z)); + CHECK_ERROR(cudaMemcpy(shared_zgrid, host_common->spline_zgrid, size_Z, cudaMemcpyHostToDevice)); shared_common->spline_zgrid = shared_zgrid; if (run.debug){ @@ -55,9 +55,9 @@ void init_grid(ShotProp shot, RunProp run, TaigaCommons *host_common, TaigaCommo } void free_grid(TaigaCommons *shared_common){ - cudaFree(shared_common->grid_size); - cudaFree(shared_common->spline_rgrid); - cudaFree(shared_common->spline_zgrid); + CHECK_ERROR(cudaFree(shared_common->grid_size)); + CHECK_ERROR(cudaFree(shared_common->spline_rgrid)); + CHECK_ERROR(cudaFree(shared_common->spline_zgrid)); } From 2f73533099ec13da392ad1b182472cbd02ca76df Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 18:01:25 +0200 Subject: [PATCH 08/15] #22 Add error management --- src/init/device/init.cu | 2 +- src/init/device/sync.cu | 23 +++++++------ src/init/fast_mode.cu | 3 +- src/init/structures/beam.cu | 61 +++++++++++++++++---------------- src/init/structures/detector.cu | 25 ++++++++------ src/init/thomson.cu | 19 +++++----- 6 files changed, 70 insertions(+), 63 deletions(-) diff --git a/src/init/device/init.cu b/src/init/device/init.cu index e0b30e8..e45362f 100644 --- a/src/init/device/init.cu +++ b/src/init/device/init.cu @@ -3,7 +3,7 @@ #include "utils/taiga_constants.h" #include "utils/physics.h" #include "utils/basic_functions.c" -#include "utils/cuda.cu" +#include "utils/cuda.cuh" void init_host(TaigaGlobals *host_global, TaigaCommons *host_common){ host_common->max_step_number = 0; diff --git a/src/init/device/sync.cu b/src/init/device/sync.cu index a17f66b..2ebe849 100644 --- a/src/init/device/sync.cu +++ b/src/init/device/sync.cu @@ -1,10 +1,11 @@ #include "cuda.h" +#include "utils/cuda.cuh" void sync_device_structs(TaigaGlobals *g_device, TaigaGlobals *g_shared, TaigaCommons *c_device, TaigaCommons *c_shared, bool is_all_io){ - cudaMemcpy(c_device, c_shared, sizeof(TaigaCommons), cudaMemcpyHostToDevice); + CHECK_ERROR(cudaMemcpy(c_device, c_shared, sizeof(TaigaCommons), cudaMemcpyHostToDevice)); if (is_all_io){ - cudaMemcpy(g_device, g_shared, sizeof(TaigaGlobals), cudaMemcpyHostToDevice); + CHECK_ERROR(cudaMemcpy(g_device, g_shared, sizeof(TaigaGlobals), cudaMemcpyHostToDevice)); } } @@ -22,22 +23,22 @@ void coord_memcopy_back(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals double* host_time_of_flight =(double*)malloc(size_coord); int* host_detcellid =(int*)malloc(size_detcellid); - cudaMemcpy(host_rad, g_shared->rad, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_rad, g_shared->rad, size_coord, cudaMemcpyDeviceToHost)); g_host->rad = host_rad; - cudaMemcpy(host_z, g_shared->z, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_z, g_shared->z, size_coord, cudaMemcpyDeviceToHost)); g_host->z = host_z; - cudaMemcpy(host_tor, g_shared->tor, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_tor, g_shared->tor, size_coord, cudaMemcpyDeviceToHost)); g_host->tor = host_tor; - cudaMemcpy(host_vrad, g_shared->vrad, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_vrad, g_shared->vrad, size_coord, cudaMemcpyDeviceToHost)); g_host->vrad = host_vrad; - cudaMemcpy(host_vz, g_shared->vz, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_vz, g_shared->vz, size_coord, cudaMemcpyDeviceToHost)); g_host->vz = host_vz; - cudaMemcpy(host_vtor, g_shared->vtor, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_vtor, g_shared->vtor, size_coord, cudaMemcpyDeviceToHost)); g_host->vtor = host_vtor; - cudaMemcpy(host_intensity, g_shared->intensity, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_intensity, g_shared->intensity, size_coord, cudaMemcpyDeviceToHost)); g_host->intensity = host_intensity; - cudaMemcpy(host_time_of_flight, g_shared->time_of_flight, size_coord, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_time_of_flight, g_shared->time_of_flight, size_coord, cudaMemcpyDeviceToHost)); g_host->time_of_flight = host_time_of_flight; - cudaMemcpy(host_detcellid, g_shared->detcellid, size_detcellid, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_detcellid, g_shared->detcellid, size_detcellid, cudaMemcpyDeviceToHost)); g_host->detcellid = host_detcellid; } \ No newline at end of file diff --git a/src/init/fast_mode.cu b/src/init/fast_mode.cu index 14114b1..82af3a4 100644 --- a/src/init/fast_mode.cu +++ b/src/init/fast_mode.cu @@ -1,11 +1,12 @@ #include #include "core/init/generate_coords.cuh" #include "utils/physics.h" +#include "utils/cuda.cuh" void init_fastmode(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals *device_global){ BeamProfile *device_prof; size_t size_prof = sizeof(BeamProfile); - cudaMalloc((void **) &device_prof, size_prof); + CHECK_ERROR(cudaMalloc((void **) &device_prof, size_prof)); init_beam_profile(device_prof, shot); generate_coords <<< run.block_number, run.block_size >>> (device_global, beam, device_prof, get_mass(beam.species, beam.charge)); } \ No newline at end of file diff --git a/src/init/structures/beam.cu b/src/init/structures/beam.cu index c8c0e1e..fef9bcb 100644 --- a/src/init/structures/beam.cu +++ b/src/init/structures/beam.cu @@ -1,4 +1,5 @@ #include "interface/data_import/beam.h" +#include "utils/cuda.cuh" void init_coords(BeamProp *beam, ShotProp *shot, RunProp *run, TaigaGlobals *g_host, TaigaGlobals *g_shared) { size_t size_coord = run->block_size * run->block_number * sizeof(double); @@ -40,26 +41,26 @@ void init_coords(BeamProp *beam, ShotProp *shot, RunProp *run, TaigaGlobals *g_h memcpy(g_shared, g_host, size_globals); } - cudaMalloc((void **) &shared_rad, size_coord); - cudaMalloc((void **) &shared_z, size_coord); - cudaMalloc((void **) &shared_tor, size_coord); - cudaMalloc((void **) &shared_vrad, size_coord); - cudaMalloc((void **) &shared_vz, size_coord); - cudaMalloc((void **) &shared_vtor, size_coord); - cudaMalloc((void **) &shared_intensity, size_coord); - cudaMalloc((void **) &shared_time_of_flight, size_coord); - cudaMalloc((void **) &shared_detcellid, size_detcellid); + CHECK_ERROR(cudaMalloc((void **) &shared_rad, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_z, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_tor, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_vrad, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_vz, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_vtor, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_intensity, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_time_of_flight, size_coord)); + CHECK_ERROR(cudaMalloc((void **) &shared_detcellid, size_detcellid)); if (run->mode == ALL_IO){ - cudaMemcpy(shared_rad, g_host->rad, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_z, g_host->z, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_tor, g_host->tor, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_vrad, g_host->vrad, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_vz, g_host->vz, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_vtor, g_host->vtor, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_intensity, g_host->intensity, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_time_of_flight, g_host->time_of_flight, size_coord, cudaMemcpyHostToDevice); - cudaMemcpy(shared_detcellid, g_host->detcellid, size_detcellid, cudaMemcpyHostToDevice); + CHECK_ERROR(cudaMemcpy(shared_rad, g_host->rad, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_z, g_host->z, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_tor, g_host->tor, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_vrad, g_host->vrad, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_vz, g_host->vz, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_vtor, g_host->vtor, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_intensity, g_host->intensity, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_time_of_flight, g_host->time_of_flight, size_coord, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_detcellid, g_host->detcellid, size_detcellid, cudaMemcpyHostToDevice)); } g_shared->rad = shared_rad; @@ -86,21 +87,21 @@ void init_beam_profile(BeamProfile *device_prof, ShotProp shot){ double *shared_radial_profile; double *shared_cross_grid; double *shared_cross_profile; - - cudaMalloc((void **) &shared_radial_grid, size_rad_prof); - cudaMalloc((void **) &shared_radial_profile, size_rad_prof); - cudaMalloc((void **) &shared_cross_grid, size_cross_prof); - cudaMalloc((void **) &shared_cross_profile, size_cross_prof); - - cudaMemcpy(shared_radial_grid, host_prof->radial_grid, size_rad_prof, cudaMemcpyHostToDevice); - cudaMemcpy(shared_radial_profile, host_prof->radial_profile, size_rad_prof, cudaMemcpyHostToDevice); - cudaMemcpy(shared_cross_grid, host_prof->cross_grid, size_cross_prof, cudaMemcpyHostToDevice); - cudaMemcpy(shared_cross_profile, host_prof->cross_profile, size_cross_prof, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMalloc((void **) &shared_radial_grid, size_rad_prof)); + CHECK_ERROR(cudaMalloc((void **) &shared_radial_profile, size_rad_prof)); + CHECK_ERROR(cudaMalloc((void **) &shared_cross_grid, size_cross_prof)); + CHECK_ERROR(cudaMalloc((void **) &shared_cross_profile, size_cross_prof)); + + CHECK_ERROR(cudaMemcpy(shared_radial_grid, host_prof->radial_grid, size_rad_prof, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_radial_profile, host_prof->radial_profile, size_rad_prof, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_cross_grid, host_prof->cross_grid, size_cross_prof, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_cross_profile, host_prof->cross_profile, size_cross_prof, cudaMemcpyHostToDevice)); shared_prof->radial_grid = shared_radial_grid; shared_prof->radial_profile = shared_radial_profile; shared_prof->cross_grid = shared_cross_grid; shared_prof->cross_profile = shared_cross_profile; - - cudaMemcpy(device_prof, shared_prof, size_prof, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMemcpy(device_prof, shared_prof, size_prof, cudaMemcpyHostToDevice)); } diff --git a/src/init/structures/detector.cu b/src/init/structures/detector.cu index 7d3e7a3..e508a1b 100644 --- a/src/init/structures/detector.cu +++ b/src/init/structures/detector.cu @@ -1,3 +1,6 @@ +#include "init/structures/detector.cuh" +#include "utils/cuda.cuh" + void init_detector(DetectorProp* shared_detector, DetectorProp *device_detector, ShotProp shot){ double *host_detector_xgrid, *device_detector_xgrid; double *host_detector_ygrid, *device_detector_ygrid; @@ -15,18 +18,18 @@ void init_detector(DetectorProp* shared_detector, DetectorProp *device_detector, size_t size_detector_ygrid = 2 * shared_detector->length_ygrid * sizeof(double); size_t size_counter = shared_detector->number_of_detector_cells * sizeof(double); size_t size_detector = sizeof(DetectorProp); - - cudaMalloc((void **) &device_detector_xgrid, size_detector_xgrid); - cudaMalloc((void **) &device_detector_ygrid, size_detector_ygrid); - cudaMalloc((void **) &device_counter, size_counter); - - cudaMemcpy(device_detector_xgrid, host_detector_xgrid, size_detector_xgrid, cudaMemcpyHostToDevice); - cudaMemcpy(device_detector_ygrid, host_detector_ygrid, size_detector_ygrid, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMalloc((void **) &device_detector_xgrid, size_detector_xgrid)); + CHECK_ERROR(cudaMalloc((void **) &device_detector_ygrid, size_detector_ygrid)); + CHECK_ERROR(cudaMalloc((void **) &device_counter, size_counter)); + + CHECK_ERROR(cudaMemcpy(device_detector_xgrid, host_detector_xgrid, size_detector_xgrid, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(device_detector_ygrid, host_detector_ygrid, size_detector_ygrid, cudaMemcpyHostToDevice)); shared_detector->xgrid = device_detector_xgrid; shared_detector->ygrid = device_detector_ygrid; shared_detector->counter = device_counter; - - cudaMemcpy(device_detector, shared_detector, size_detector, cudaMemcpyHostToDevice); + + CHECK_ERROR(cudaMemcpy(device_detector, shared_detector, size_detector, cudaMemcpyHostToDevice)); }else{ printf("===============================\n"); printf("Detector postprocessor module: OFF\n"); @@ -55,8 +58,8 @@ void set_detector_geometry(ShotProp shot, TaigaCommons *host_common, TaigaCommon tokaniser = strtok(NULL, ","); host_common->detector_geometry[4] = (strtod(tokaniser, NULL) * PI / 180.0); - cudaMalloc((void **) &shared_detector_geometry, size_detector); - cudaMemcpy(shared_detector_geometry, host_common->detector_geometry, size_detector, cudaMemcpyHostToDevice); + CHECK_ERROR(cudaMalloc((void **) &shared_detector_geometry, size_detector)); + CHECK_ERROR(cudaMemcpy(shared_detector_geometry, host_common->detector_geometry, size_detector, cudaMemcpyHostToDevice)); shared_common->detector_geometry = shared_detector_geometry; } } \ No newline at end of file diff --git a/src/init/thomson.cu b/src/init/thomson.cu index 9805475..d5d7fa5 100644 --- a/src/init/thomson.cu +++ b/src/init/thomson.cu @@ -1,4 +1,5 @@ #include "init/thomson.cuh" +#include "utils/cuda.cuh" void import_thomson_profiles(ShotProp shot, TaigaCommons *c) { c->ts_length = read_vector(&c->ts_psi, "input/tsProf", shot.name, "flux.prof"); @@ -11,13 +12,13 @@ void set_thomson_profiles(ShotProp shot, TaigaCommons *host_common, TaigaCommons import_thomson_profiles(shot, host_common); size_t size_ts = host_common->ts_length * sizeof(double); - cudaMalloc((void **) &shared_ts_psi, size_ts); - cudaMalloc((void **) &shared_ts_density, size_ts); - cudaMalloc((void **) &shared_ts_temperature, size_ts); + CHECK_ERROR(cudaMalloc((void **) &shared_ts_psi, size_ts)); + CHECK_ERROR(cudaMalloc((void **) &shared_ts_density, size_ts)); + CHECK_ERROR(cudaMalloc((void **) &shared_ts_temperature, size_ts)); - cudaMemcpy(shared_ts_psi, host_common->ts_psi, size_ts, cudaMemcpyHostToDevice); - cudaMemcpy(shared_ts_density, host_common->ts_density, size_ts, cudaMemcpyHostToDevice); - cudaMemcpy(shared_ts_temperature, host_common->ts_temperature, size_ts, cudaMemcpyHostToDevice); + CHECK_ERROR(cudaMemcpy(shared_ts_psi, host_common->ts_psi, size_ts, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_ts_density, host_common->ts_density, size_ts, cudaMemcpyHostToDevice)); + CHECK_ERROR(cudaMemcpy(shared_ts_temperature, host_common->ts_temperature, size_ts, cudaMemcpyHostToDevice)); shared_common->ts_length = host_common->ts_length; shared_common->ts_psi = shared_ts_psi; @@ -26,7 +27,7 @@ void set_thomson_profiles(ShotProp shot, TaigaCommons *host_common, TaigaCommons } void free_thomson_profiles(TaigaCommons *shared_common) { - cudaFree(shared_common->ts_psi); - cudaFree(shared_common->ts_density); - cudaFree(shared_common->ts_temperature); + CHECK_ERROR(cudaFree(shared_common->ts_psi)); + CHECK_ERROR(cudaFree(shared_common->ts_density)); + CHECK_ERROR(cudaFree(shared_common->ts_temperature)); } From 261cf67805c565a1b653d9084b7a1c6716e1cc19 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 18:03:02 +0200 Subject: [PATCH 09/15] #22 Add CUDA memory free function --- src/init/fast_mode.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/src/init/fast_mode.cu b/src/init/fast_mode.cu index 82af3a4..e8bae88 100644 --- a/src/init/fast_mode.cu +++ b/src/init/fast_mode.cu @@ -9,4 +9,5 @@ void init_fastmode(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals *devi CHECK_ERROR(cudaMalloc((void **) &device_prof, size_prof)); init_beam_profile(device_prof, shot); generate_coords <<< run.block_number, run.block_size >>> (device_global, beam, device_prof, get_mass(beam.species, beam.charge)); + CHECK_ERROR(cudaFree(device_prof)); } \ No newline at end of file From 10d7478de67354687d4eaa7ec48424f87a91a90d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 20:14:57 +0200 Subject: [PATCH 10/15] #22 Fix error manager --- include/utils/cuda.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/utils/cuda.cuh b/include/utils/cuda.cuh index 155b9ee..4e8b44f 100644 --- a/include/utils/cuda.cuh +++ b/include/utils/cuda.cuh @@ -7,7 +7,7 @@ void set_cuda(int debug_flag); inline void manage_cuda_error(cudaError_t cuda_code, const char *filename, int line, bool abort=true) { if (cuda_code != cudaSuccess) { - fprintf(stderr,"CUDA ERROR in %s (line %d):\n %s\n", filename, line, cudaGetErrorString(code)); + fprintf(stderr,"CUDA ERROR in %s (line %d):\n %s\n", filename, line, cudaGetErrorString(cuda_code)); } } From 3b5551225fb3d597fa73f0c06de23a2e247e4cb4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 20:18:40 +0200 Subject: [PATCH 11/15] #22 Fix error manager --- src/main.cu | 2 +- tests/test_field.cu | 1 + 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/src/main.cu b/src/main.cu index fbc7ce7..58455ff 100644 --- a/src/main.cu +++ b/src/main.cu @@ -15,7 +15,7 @@ #include "utils/debug_functions.c" #include "utils/basic_functions.h" #include "utils/dataio/dir_functions.c" -#include "utils/cuda.cuh" +#include "utils/cuda.cu" #include "utils/physics.c" #include "utils/dataio/data_import.c" diff --git a/tests/test_field.cu b/tests/test_field.cu index 108ff2d..2425246 100644 --- a/tests/test_field.cu +++ b/tests/test_field.cu @@ -7,6 +7,7 @@ #include "utils/prop.c" #include "utils/basic_functions.h" #include "utils/dataio/data_import.c" +#include "utils/cuda.cu" #include "interface/data_import/field_import.cu" #include "interface/parameter_reader.c" #include "init/device/sync.cu" From 9003e246c484ae163a431c9b9417516c2bc17c9a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Fri, 20 May 2022 22:01:43 +0200 Subject: [PATCH 12/15] #22 Fix error manager --- include/utils/cuda.cuh | 2 +- src/main.cu | 27 +++++++++++++++------------ 2 files changed, 16 insertions(+), 13 deletions(-) diff --git a/include/utils/cuda.cuh b/include/utils/cuda.cuh index 4e8b44f..16c5c32 100644 --- a/include/utils/cuda.cuh +++ b/include/utils/cuda.cuh @@ -5,7 +5,7 @@ void set_cuda(int debug_flag); -inline void manage_cuda_error(cudaError_t cuda_code, const char *filename, int line, bool abort=true) { +inline void manage_cuda_error(cudaError_t cuda_code, const char *filename, int line) { if (cuda_code != cudaSuccess) { fprintf(stderr,"CUDA ERROR in %s (line %d):\n %s\n", filename, line, cudaGetErrorString(cuda_code)); } diff --git a/src/main.cu b/src/main.cu index 58455ff..5b3577a 100644 --- a/src/main.cu +++ b/src/main.cu @@ -56,6 +56,8 @@ #include "detector/postproc.cu" #include "detector/sum.cu" +void free_taiga(); + void input_init_taiga(int argc, char *argv[], ShotProp *shot, BeamProp *beam, RunProp *run){ char *input; @@ -165,10 +167,10 @@ int main(int argc, char *argv[]){ host_common = (TaigaCommons*)malloc(size_commons); shared_common = (TaigaCommons*)malloc(size_commons); shared_detector = (DetectorProp*)malloc(size_detector_prop); - - cudaMalloc((void **) &device_global, size_global); - cudaMalloc((void **) &device_common, size_commons); - cudaMalloc((void **) &device_detector, size_detector_prop); + + CHECK_ERROR(cudaMalloc((void **) &device_global, size_global)); + CHECK_ERROR(cudaMalloc((void **) &device_common, size_commons)); + CHECK_ERROR(cudaMalloc((void **) &device_detector, size_detector_prop)); init_host(host_global, host_common); @@ -212,8 +214,8 @@ int main(int argc, char *argv[]){ } host_service_array[4] = 55555.55555; - cudaMalloc((void **) &device_service_array, dimService); - cudaMemcpy(device_service_array, host_service_array, dimService, cudaMemcpyHostToDevice); + CHECK_ERROR(cudaMalloc((void **) &device_service_array, dimService)); + CHECK_ERROR(cudaMemcpy(device_service_array, host_service_array, dimService, cudaMemcpyHostToDevice)); // if (run.mode == ALL_IO){ @@ -236,7 +238,7 @@ int main(int argc, char *argv[]){ taiga <<< run.block_number, run.block_size >>> (device_global, device_common, device_service_array); if (step_i == 0) cudaEventRecord(cuda_event_core_end, 0); - cudaEventSynchronize(cuda_event_core_end); + CHECK_ERROR(cudaEventSynchronize(cuda_event_core_end)); if (run.mode == ALL_IO){ // ION COORDS (device2HOST) @@ -257,9 +259,9 @@ int main(int argc, char *argv[]){ // Get CUDA timer cudaEventElapsedTime(&cuda_event_core, cuda_event_core_start, cuda_event_core_end); cudaEventElapsedTime(&cuda_event_copy, cuda_event_copy_start, cuda_event_copy_end); - if (run.mode == ALL_IO) run.cpu_time_copy = ((double) (4.0+run.step_host)*(cpu_event_copy_end - cpu_event_copy_start)) / CLOCKS_PER_SEC; - run.cuda_time_copy = (double) (1.0+run.step_host)*cuda_event_copy/1000.0; - run.cuda_time_core = run.step_host*cuda_event_core/1000.0; + if (run.mode == ALL_IO) run.cpu_time_copy = ((double) (4+run.step_host)*(double)(cpu_event_copy_end - cpu_event_copy_start)) / CLOCKS_PER_SEC; + run.cuda_time_copy = (double) (1+run.step_host)*cuda_event_copy/1000.0; + run.cuda_time_core = (double)run.step_host*cuda_event_core/1000.0; printf("===============================\n"); printf ("CUDA kernel runtime: %lf s\n", run.cuda_time_core); printf ("CUDA memcopy time: %lf s\n", run.cuda_time_copy); @@ -267,7 +269,7 @@ int main(int argc, char *argv[]){ printf("===============================\n"); //! MEMCOPY (device2HOST) - cudaMemcpy(host_service_array, device_service_array, dimService, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_service_array, device_service_array, dimService, cudaMemcpyDeviceToHost)); if(host_service_array[0] != 42.24){ printf("\n +----------------------------+\n | Fatal error in running. | \n | The CUDA did not run well. |\n | Service value: %11lf |\n +----------------------------+\n\n", host_service_array[0]); }else{ @@ -292,8 +294,9 @@ int main(int argc, char *argv[]){ //! FREE host_service_array variables (RAM, cuda) free(host_service_array); cudaFree(device_service_array); + //free_taiga(); - cudaThreadExit(); + CHECK_ERROR(cudaThreadExit()); printf("Ready.\n\n"); } } From 26a1e9f6ffce66a36d025ea96dd4cec7c72091e5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Sat, 21 May 2022 03:16:07 +0200 Subject: [PATCH 13/15] #22 Collect free together --- include/init/device/init.cuh | 1 - include/utils/free.cuh | 7 +++++ src/detector/module.cu | 5 ++-- src/init/device/init.cu | 7 ----- src/init/thomson.cu | 6 ----- src/main.cu | 1 + src/utils/free.cu | 52 ++++++++++++++++++++++++++++++++++++ 7 files changed, 63 insertions(+), 16 deletions(-) create mode 100644 include/utils/free.cuh create mode 100644 src/utils/free.cu diff --git a/include/init/device/init.cuh b/include/init/device/init.cuh index 931f902..8f95822 100644 --- a/include/init/device/init.cuh +++ b/include/init/device/init.cuh @@ -3,7 +3,6 @@ void init_host(TaigaGlobals *host_global, TaigaCommons *host_common); void init_grid(ShotProp shot, RunProp run, TaigaCommons *host_common, TaigaCommons *shared_common); -void free_grid(TaigaCommons *shared_common); void init_device_structs(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals *shared_global, TaigaCommons *shared_common); void set_particle_number(RunProp *run, TaigaGlobals *host_global, TaigaGlobals *shared_global); diff --git a/include/utils/free.cuh b/include/utils/free.cuh new file mode 100644 index 0000000..6f7ca72 --- /dev/null +++ b/include/utils/free.cuh @@ -0,0 +1,7 @@ +#ifndef TAIGA_FREE_CUH +#define TAIGA_FREE_CUH + + +void free_common(TaigaCommons *shared_common); + +#endif //TAIGA_FREE_CUH diff --git a/src/detector/module.cu b/src/detector/module.cu index 377ecc8..d02d249 100644 --- a/src/detector/module.cu +++ b/src/detector/module.cu @@ -2,6 +2,7 @@ #include #include #include +#include "utils/cuda.cuh" void export_detector(DetectorProp* shared_detector, DetectorProp *device_detector, TaigaGlobals *shared_global, ShotProp shot, RunProp run){ if (shared_detector->detector_module_on){ @@ -9,13 +10,13 @@ void export_detector(DetectorProp* shared_detector, DetectorProp *device_detecto size_t size_counter = shared_detector->number_of_detector_cells * sizeof(double); host_counter = (double*)malloc(size_counter); - cudaMemcpy(host_counter, shared_detector->counter, size_counter, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_counter, shared_detector->counter, size_counter, cudaMemcpyDeviceToHost)); export_data(host_counter, shared_detector->number_of_detector_cells, run.folder_out, run.runnumber, "detector", "cellcounter.dat", shared_detector->length_xgrid); if (run.mode == ALL_IO){ size_t size_detcellid = run.block_size * run.block_number * sizeof(int); int* host_detcellid =(int*)malloc(size_detcellid); - cudaMemcpy(host_detcellid, shared_global->detcellid, size_detcellid, cudaMemcpyDeviceToHost); + CHECK_ERROR(cudaMemcpy(host_detcellid, shared_global->detcellid, size_detcellid, cudaMemcpyDeviceToHost)); export_data(host_detcellid, shared_global->particle_number, run.folder_out, run.runnumber, "detector", "cellid.dat"); } } diff --git a/src/init/device/init.cu b/src/init/device/init.cu index e45362f..1e46307 100644 --- a/src/init/device/init.cu +++ b/src/init/device/init.cu @@ -54,13 +54,6 @@ void init_grid(ShotProp shot, RunProp run, TaigaCommons *host_common, TaigaCommo } } -void free_grid(TaigaCommons *shared_common){ - CHECK_ERROR(cudaFree(shared_common->grid_size)); - CHECK_ERROR(cudaFree(shared_common->spline_rgrid)); - CHECK_ERROR(cudaFree(shared_common->spline_zgrid)); -} - - void init_device_structs(BeamProp beam, ShotProp shot, RunProp run, TaigaGlobals *shared_global, TaigaCommons *shared_common){ shared_global->particle_number = run.particle_number; shared_common->max_step_number = run.step_device; diff --git a/src/init/thomson.cu b/src/init/thomson.cu index d5d7fa5..13acab9 100644 --- a/src/init/thomson.cu +++ b/src/init/thomson.cu @@ -25,9 +25,3 @@ void set_thomson_profiles(ShotProp shot, TaigaCommons *host_common, TaigaCommons shared_common->ts_density = shared_ts_density; shared_common->ts_temperature = shared_ts_temperature; } - -void free_thomson_profiles(TaigaCommons *shared_common) { - CHECK_ERROR(cudaFree(shared_common->ts_psi)); - CHECK_ERROR(cudaFree(shared_common->ts_density)); - CHECK_ERROR(cudaFree(shared_common->ts_temperature)); -} diff --git a/src/main.cu b/src/main.cu index 5b3577a..21eb1ce 100644 --- a/src/main.cu +++ b/src/main.cu @@ -17,6 +17,7 @@ #include "utils/dataio/dir_functions.c" #include "utils/cuda.cu" #include "utils/physics.c" +#include "utils/free.cu" #include "utils/dataio/data_import.c" #include "interface/data_import/field_import.cu" diff --git a/src/utils/free.cu b/src/utils/free.cu new file mode 100644 index 0000000..a6b161d --- /dev/null +++ b/src/utils/free.cu @@ -0,0 +1,52 @@ +#include "utils/free.cuh" + + + + +void free_taiga() { + +} + + +void free_detector(DetectorProp *detector){ + CHECK_ERROR(cudaFree(detector->counter)); + CHECK_ERROR(cudaFree(detector->xgrid)); + CHECK_ERROR(cudaFree(detector->ygrid)); +} + +void free_beam(BeamProfile *beam_profile){ + CHECK_ERROR(cudaFree(beam_profile->radial_grid)); + CHECK_ERROR(cudaFree(beam_profile->radial_profile)); + CHECK_ERROR(cudaFree(beam_profile->cross_grid)); + CHECK_ERROR(cudaFree(beam_profile->cross_profile)); +} + +void free_global(TaigaGlobals *taiga_global){ + CHECK_ERROR(cudaFree(taiga_global->rad)); + CHECK_ERROR(cudaFree(taiga_global->z)); + CHECK_ERROR(cudaFree(taiga_global->tor)); + CHECK_ERROR(cudaFree(taiga_global->vrad)); + CHECK_ERROR(cudaFree(taiga_global->vz)); + CHECK_ERROR(cudaFree(taiga_global->vtor)); + CHECK_ERROR(cudaFree(taiga_global->detcellid)); + CHECK_ERROR(cudaFree(taiga_global->intensity)); + CHECK_ERROR(cudaFree(taiga_global->time_of_flight)); +} + +void free_common(TaigaCommons *shared_common){ + CHECK_ERROR(cudaFree(shared_common->grid_size)); + CHECK_ERROR(cudaFree(shared_common->spline_rgrid)); + CHECK_ERROR(cudaFree(shared_common->spline_zgrid)); + CHECK_ERROR(cudaFree(shared_common->brad)); + CHECK_ERROR(cudaFree(shared_common->bz)); + CHECK_ERROR(cudaFree(shared_common->btor)); + CHECK_ERROR(cudaFree(shared_common->erad)); + CHECK_ERROR(cudaFree(shared_common->ez)); + CHECK_ERROR(cudaFree(shared_common->etor)); + CHECK_ERROR(cudaFree(shared_common->psi_n)); + CHECK_ERROR(cudaFree(shared_common->detector_geometry)); + CHECK_ERROR(cudaFree(shared_common->ts_psi)); + CHECK_ERROR(cudaFree(shared_common->ts_density)); + CHECK_ERROR(cudaFree(shared_common->ts_temperature)); +} + From 0deb88000893b359780a82f946a8cf4e976ba7ac Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Sat, 21 May 2022 03:34:41 +0200 Subject: [PATCH 14/15] #22 Collect free together --- include/utils/free.cuh | 11 ++++++- src/utils/free.cu | 71 ++++++++++++++++++++---------------------- 2 files changed, 44 insertions(+), 38 deletions(-) diff --git a/include/utils/free.cuh b/include/utils/free.cuh index 6f7ca72..9db4843 100644 --- a/include/utils/free.cuh +++ b/include/utils/free.cuh @@ -1,7 +1,16 @@ #ifndef TAIGA_FREE_CUH #define TAIGA_FREE_CUH +#define HOST 0 +#define DEVICE 1 +#define FREE(where, p) { \ + if (where == DEVICE) {CHECK_ERROR(cudaFree(p));} \ + else {free(p);} \ +} -void free_common(TaigaCommons *shared_common); +void free_detector(DetectorProp *detector, int where); +void free_beam(BeamProfile *beam_profile, int where); +void free_global(TaigaGlobals *taiga_global, int where); +void free_common(TaigaCommons *shared_common, int where); #endif //TAIGA_FREE_CUH diff --git a/src/utils/free.cu b/src/utils/free.cu index a6b161d..d43dc30 100644 --- a/src/utils/free.cu +++ b/src/utils/free.cu @@ -1,52 +1,49 @@ #include "utils/free.cuh" - - - void free_taiga() { } -void free_detector(DetectorProp *detector){ - CHECK_ERROR(cudaFree(detector->counter)); - CHECK_ERROR(cudaFree(detector->xgrid)); - CHECK_ERROR(cudaFree(detector->ygrid)); +void free_detector(DetectorProp *detector, int where){ + FREE(where, detector->counter); + FREE(where, detector->xgrid); + FREE(where, detector->ygrid); } -void free_beam(BeamProfile *beam_profile){ - CHECK_ERROR(cudaFree(beam_profile->radial_grid)); - CHECK_ERROR(cudaFree(beam_profile->radial_profile)); - CHECK_ERROR(cudaFree(beam_profile->cross_grid)); - CHECK_ERROR(cudaFree(beam_profile->cross_profile)); +void free_beam(BeamProfile *beam_profile, int where){ + FREE(where, beam_profile->radial_grid); + FREE(where, beam_profile->radial_profile); + FREE(where, beam_profile->cross_grid); + FREE(where, beam_profile->cross_profile); } -void free_global(TaigaGlobals *taiga_global){ - CHECK_ERROR(cudaFree(taiga_global->rad)); - CHECK_ERROR(cudaFree(taiga_global->z)); - CHECK_ERROR(cudaFree(taiga_global->tor)); - CHECK_ERROR(cudaFree(taiga_global->vrad)); - CHECK_ERROR(cudaFree(taiga_global->vz)); - CHECK_ERROR(cudaFree(taiga_global->vtor)); - CHECK_ERROR(cudaFree(taiga_global->detcellid)); - CHECK_ERROR(cudaFree(taiga_global->intensity)); - CHECK_ERROR(cudaFree(taiga_global->time_of_flight)); +void free_global(TaigaGlobals *taiga_global, int where){ + FREE(where, taiga_global->rad); + FREE(where, taiga_global->z); + FREE(where, taiga_global->tor); + FREE(where, taiga_global->vrad); + FREE(where, taiga_global->vz); + FREE(where, taiga_global->vtor); + FREE(where, taiga_global->detcellid); + FREE(where, taiga_global->intensity); + FREE(where, taiga_global->time_of_flight); } -void free_common(TaigaCommons *shared_common){ - CHECK_ERROR(cudaFree(shared_common->grid_size)); - CHECK_ERROR(cudaFree(shared_common->spline_rgrid)); - CHECK_ERROR(cudaFree(shared_common->spline_zgrid)); - CHECK_ERROR(cudaFree(shared_common->brad)); - CHECK_ERROR(cudaFree(shared_common->bz)); - CHECK_ERROR(cudaFree(shared_common->btor)); - CHECK_ERROR(cudaFree(shared_common->erad)); - CHECK_ERROR(cudaFree(shared_common->ez)); - CHECK_ERROR(cudaFree(shared_common->etor)); - CHECK_ERROR(cudaFree(shared_common->psi_n)); - CHECK_ERROR(cudaFree(shared_common->detector_geometry)); - CHECK_ERROR(cudaFree(shared_common->ts_psi)); - CHECK_ERROR(cudaFree(shared_common->ts_density)); - CHECK_ERROR(cudaFree(shared_common->ts_temperature)); +void free_common(TaigaCommons *shared_common, int where){ + FREE(where, shared_common->grid_size); + FREE(where, shared_common->spline_rgrid); + FREE(where, shared_common->spline_zgrid); + FREE(where, shared_common->brad); + FREE(where, shared_common->bz); + FREE(where, shared_common->btor); + FREE(where, shared_common->erad); + FREE(where, shared_common->ez); + FREE(where, shared_common->etor); + FREE(where, shared_common->psi_n); + FREE(where, shared_common->detector_geometry); + FREE(where, shared_common->ts_psi); + FREE(where, shared_common->ts_density); + FREE(where, shared_common->ts_temperature); } From caf236b2382f758c361aea009799423e4df22cdd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?M=C3=A1ty=C3=A1s=20Aradi?= <20679946+matyasaradi@users.noreply.github.com> Date: Sat, 21 May 2022 03:58:43 +0200 Subject: [PATCH 15/15] #22 Collect free together --- include/utils/free.cuh | 13 +++++++- src/main.cu | 12 +++---- src/utils/free.cu | 74 +++++++++++++++++++++++++++--------------- 3 files changed, 65 insertions(+), 34 deletions(-) diff --git a/include/utils/free.cuh b/include/utils/free.cuh index 9db4843..e83fc36 100644 --- a/include/utils/free.cuh +++ b/include/utils/free.cuh @@ -2,12 +2,23 @@ #define TAIGA_FREE_CUH #define HOST 0 -#define DEVICE 1 +#define SHARED 1 +#define DEVICE 2 #define FREE(where, p) { \ + if (where == HOST) {free(p);} \ + else {CHECK_ERROR(cudaFree(p));} \ +} +#define FREE_MAIN(where, p) { \ if (where == DEVICE) {CHECK_ERROR(cudaFree(p));} \ else {free(p);} \ } +void free_taiga(TaigaGlobals *host_global, TaigaGlobals *shared_global, TaigaGlobals *device_global, + TaigaCommons *host_common, TaigaCommons *shared_common, TaigaCommons *device_common, + DetectorProp *shared_detector, DetectorProp *device_detector, + ShotProp *shot, BeamProp *beam, RunProp *run, + double *host_service_array, double *device_service_array); + void free_detector(DetectorProp *detector, int where); void free_beam(BeamProfile *beam_profile, int where); void free_global(TaigaGlobals *taiga_global, int where); diff --git a/src/main.cu b/src/main.cu index 21eb1ce..6ef5ec3 100644 --- a/src/main.cu +++ b/src/main.cu @@ -57,10 +57,7 @@ #include "detector/postproc.cu" #include "detector/sum.cu" -void free_taiga(); - void input_init_taiga(int argc, char *argv[], ShotProp *shot, BeamProp *beam, RunProp *run){ - char *input; for (int i=1; icounter); FREE(where, detector->xgrid); FREE(where, detector->ygrid); + FREE_MAIN(where, detector); } void free_beam(BeamProfile *beam_profile, int where){ @@ -16,34 +33,37 @@ void free_beam(BeamProfile *beam_profile, int where){ FREE(where, beam_profile->radial_profile); FREE(where, beam_profile->cross_grid); FREE(where, beam_profile->cross_profile); + FREE_MAIN(where, beam_profile); } -void free_global(TaigaGlobals *taiga_global, int where){ - FREE(where, taiga_global->rad); - FREE(where, taiga_global->z); - FREE(where, taiga_global->tor); - FREE(where, taiga_global->vrad); - FREE(where, taiga_global->vz); - FREE(where, taiga_global->vtor); - FREE(where, taiga_global->detcellid); - FREE(where, taiga_global->intensity); - FREE(where, taiga_global->time_of_flight); +void free_global(TaigaGlobals *global, int where){ + FREE(where, global->rad); + FREE(where, global->z); + FREE(where, global->tor); + FREE(where, global->vrad); + FREE(where, global->vz); + FREE(where, global->vtor); + FREE(where, global->detcellid); + FREE(where, global->intensity); + FREE(where, global->time_of_flight); + FREE_MAIN(where, global); } -void free_common(TaigaCommons *shared_common, int where){ - FREE(where, shared_common->grid_size); - FREE(where, shared_common->spline_rgrid); - FREE(where, shared_common->spline_zgrid); - FREE(where, shared_common->brad); - FREE(where, shared_common->bz); - FREE(where, shared_common->btor); - FREE(where, shared_common->erad); - FREE(where, shared_common->ez); - FREE(where, shared_common->etor); - FREE(where, shared_common->psi_n); - FREE(where, shared_common->detector_geometry); - FREE(where, shared_common->ts_psi); - FREE(where, shared_common->ts_density); - FREE(where, shared_common->ts_temperature); +void free_common(TaigaCommons *common, int where){ + FREE(where, common->grid_size); + FREE(where, common->spline_rgrid); + FREE(where, common->spline_zgrid); + FREE(where, common->brad); + FREE(where, common->bz); + FREE(where, common->btor); + FREE(where, common->erad); + FREE(where, common->ez); + FREE(where, common->etor); + FREE(where, common->psi_n); + FREE(where, common->detector_geometry); + FREE(where, common->ts_psi); + FREE(where, common->ts_density); + FREE(where, common->ts_temperature); + FREE_MAIN(where, common); }