Merge pull request #163 from taiga-project/bugfix-22

Bugfix 22

.gitignore

@@ -2,6 +2,11 @@
 **/.idea/
 **/__pycache__
 
+#preproc data folder
+preproc/renate_od/data/*
+!preproc/renate_od/data/.gitkeep
+
+
 # result folder
 results/*
 !results/.gitkeep

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) {
+    if (cuda_code != cudaSuccess) {
+        fprintf(stderr,"CUDA ERROR in %s (line %d):\n %s\n", filename, line, cudaGetErrorString(cuda_code));
+    }
+}
+
+#endif //TAIGA_CUDA_CUH

include/utils/free.cuh

@@ -0,0 +1,27 @@
+#ifndef TAIGA_FREE_CUH
+#define TAIGA_FREE_CUH
+
+#define HOST 0
+#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);
+void free_common(TaigaCommons *shared_common, int where);
+
+#endif //TAIGA_FREE_CUH

preproc/renate_od/data/.gitkeep

@@ -0,0 +1 @@
+

src/detector/module.cu

@@ -2,20 +2,21 @@
 #include <cuda.h>
 #include <cuda_runtime.h>
 #include <math.h>
+#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){
         double *host_counter;
         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");
         }
     }

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.cuh"
 
 void init_host(TaigaGlobals *host_global, TaigaCommons *host_common){
     host_common->max_step_number = 0;
@@ -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){

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

src/init/fast_mode.cu

@@ -1,11 +1,13 @@
 #include <cuda.h>
 #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));
+    CHECK_ERROR(cudaFree(device_prof));
 }

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));
 }

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

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;