Merge pull request #1781 from danielpeter/devel

Devel
SPECFEM · Jan 9, 2025 · 8f39006 · 8f39006
2 parents a8182ff + b4d710a
commit 8f39006
Show file tree

Hide file tree

Showing 8 changed files with 688 additions and 318 deletions.
diff --git a/src/gpu/check_fields_cuda.cu b/src/gpu/check_fields_cuda.cu
@@ -222,7 +222,7 @@ void FC_FUNC_(get_norm_acoustic_from_device,
 
   Mesh* mp = (Mesh*)(*Mesh_pointer); //get mesh pointer out of fortran integer container
   realw max = 0.0;
-  realw *d_max;
+  realw *d_max = NULL;
 
   //initializes
   *norm = 0.0f;
@@ -275,6 +275,7 @@ void FC_FUNC_(get_norm_acoustic_from_device,
 
   // on host (allocates & initializes to zero)
   h_max = (realw*) calloc(num_blocks_x*num_blocks_y,sizeof(realw));
+  if (!h_max) { exit_on_error("Error: allocating temporary host h_max array"); }
 
   // allocates memory on device
   gpuMalloc_realw((void**)&d_max,num_blocks_x*num_blocks_y);
@@ -382,7 +383,7 @@ void FC_FUNC_(get_norm_elastic_from_device,
 
   Mesh* mp = (Mesh*)(*Mesh_pointer); //get mesh pointer out of fortran integer container
   realw max,res;
-  realw *d_max;
+  realw *d_max = NULL;
 
   //initializes
   *norm = 0.0f;
@@ -402,6 +403,7 @@ void FC_FUNC_(get_norm_elastic_from_device,
 
   // on host (allocates & initializes to zero)
   h_max = (realw*) calloc(num_blocks_x*num_blocks_y,sizeof(realw));
+  if (!h_max) { exit_on_error("Error: allocating temporary host h_max array"); }
 
   // allocates memory on device
   gpuMalloc_realw((void**)&d_max,num_blocks_x*num_blocks_y);

diff --git a/src/gpu/kernels/wavefield_discontinuity_kernel.cpp b/src/gpu/kernels/wavefield_discontinuity_kernel.cpp
@@ -1 +1,30 @@
+/*
+!=====================================================================
+!
+!                          S p e c f e m 3 D
+!                          -----------------
+!
+!    Main historical authors: Dimitri Komatitsch and Jeroen Tromp
+!                             CNRS, France
+!                      and Princeton University, USA
+!                (there are currently many more authors!)
+!                          (c) October 2017
+!
+! This program is free software; you can redistribute it and/or modify
+! it under the terms of the GNU General Public License as published by
+! the Free Software Foundation; either version 3 of the License, or
+! (at your option) any later version.
+!
+! This program is distributed in the hope that it will be useful,
+! but WITHOUT ANY WARRANTY; without even the implied warranty of
+! MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+! GNU General Public License for more details.
+!
+! You should have received a copy of the GNU General Public License along
+! with this program; if not, write to the Free Software Foundation, Inc.,
+! 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+!
+!=====================================================================
+*/
+
 #include "wavefield_discontinuity_kernel.cu"
diff --git a/src/gpu/kernels/wavefield_discontinuity_kernel.cu b/src/gpu/kernels/wavefield_discontinuity_kernel.cu
@@ -1,5 +1,35 @@
+/*
+!=====================================================================
+!
+!                          S p e c f e m 3 D
+!                          -----------------
+!
+!    Main historical authors: Dimitri Komatitsch and Jeroen Tromp
+!                             CNRS, France
+!                      and Princeton University, USA
+!                (there are currently many more authors!)
+!                          (c) October 2017
+!
+! This program is free software; you can redistribute it and/or modify
+! it under the terms of the GNU General Public License as published by
+! the Free Software Foundation; either version 3 of the License, or
+! (at your option) any later version.
+!
+! This program is distributed in the hope that it will be useful,
+! but WITHOUT ANY WARRANTY; without even the implied warranty of
+! MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+! GNU General Public License for more details.
+!
+! You should have received a copy of the GNU General Public License along
+! with this program; if not, write to the Free Software Foundation, Inc.,
+! 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+!
+!=====================================================================
+*/
+
+
 __global__ void add_acceleration_discontinuity_kernel(
-                                  realw_const_p accel_wd, 
+                                  realw_const_p accel_wd,
                                   realw_const_p mass_in_wd,
                                   const int* boundary_to_iglob_wd,
                                   const int size, realw* accel

diff --git a/src/gpu/wavefield_discontinuity_cuda.cu b/src/gpu/wavefield_discontinuity_cuda.cu
@@ -1,8 +1,37 @@
+/*
+ !=====================================================================
+ !
+ !                         S p e c f e m 3 D
+ !                         -----------------
+ !
+ !     Main historical authors: Dimitri Komatitsch and Jeroen Tromp
+ !                              CNRS, France
+ !                       and Princeton University, USA
+ !                 (there are currently many more authors!)
+ !                           (c) October 2017
+ !
+ ! This program is free software; you can redistribute it and/or modify
+ ! it under the terms of the GNU General Public License as published by
+ ! the Free Software Foundation; either version 3 of the License, or
+ ! (at your option) any later version.
+ !
+ ! This program is distributed in the hope that it will be useful,
+ ! but WITHOUT ANY WARRANTY; without even the implied warranty of
+ ! MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ ! GNU General Public License for more details.
+ !
+ ! You should have received a copy of the GNU General Public License along
+ ! with this program; if not, write to the Free Software Foundation, Inc.,
+ ! 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+ !
+ !=====================================================================
+*/
+
 #include "mesh_constants_gpu.h"
 
 extern EXTERN_LANG
 void FC_FUNC_(wavefield_discontinuity_add_traction_cuda,
-              WAVEFIELD_DISCONTINUITY_ADD_TRACTION_CUDA)(int* size_points, 
+              WAVEFIELD_DISCONTINUITY_ADD_TRACTION_CUDA)(int* size_points,
                                                          int* size_faces,
                                                          long* Mesh_pointer){
   TRACE("wavefield_discontinuity_add_traction_cuda");
@@ -17,7 +46,7 @@ void FC_FUNC_(wavefield_discontinuity_add_traction_cuda,
 
     dim3 grid(num_blocks_x,num_blocks_y);
     dim3 threads(blocksize,1,1);
-    
+
 #ifdef USE_CUDA
     if (run_cuda) {
       add_acceleration_discontinuity_kernel
@@ -29,17 +58,17 @@ void FC_FUNC_(wavefield_discontinuity_add_traction_cuda,
 #endif
 #ifdef USE_HIP
     if (run_hip) {
-      hipLaunchKernelGGL(add_acceleration_discontinuity_kernel, 
+      hipLaunchKernelGGL(add_acceleration_discontinuity_kernel,
        dim3(grid), dim3(threads), 0, mp->compute_stream, mp->d_accel_wd,
                                                   mp->d_mass_in_wd,
-                                                  mp->d_boundary_to_iglob_wd,           
+                                                  mp->d_boundary_to_iglob_wd,
                                                   size, mp->d_accel);
     }
 #endif
-    
+
     size = (*size_faces);
     blocksize = NGLL2;
-    
+
     get_blocks_xy(size,&num_blocks_x,&num_blocks_y);
 
     dim3 grid2(num_blocks_x,num_blocks_y);

diff --git a/src/specfem3D/couple_with_injection.f90 b/src/specfem3D/couple_with_injection.f90
@@ -3328,7 +3328,7 @@ subroutine couple_with_injection_prepare_specfem_files()
   ! check integer size limit: size of reclen must fit onto an 4-byte integer
   if (reclen > int(2147483646.0 / 2)) then
     print *,'reclen needed exceeds integer 4-byte limit: ',reclen
-    print *,'  ',reclen,' custom_real/ndim/npoints_local = ',CUSTOM_REAL, NDIM, npoints_local
+    print *,'  ',reclen,' CUSTOM_REAL/ndim/npoints_local = ',CUSTOM_REAL, NDIM, npoints_local
     print *,'bit size Fortran: ',bit_size(reclen)
     call exit_MPI(myrank,"Error reclen integer limit")
   endif

diff --git a/src/specfem3D/locate_source.F90 b/src/specfem3D/locate_source.F90
@@ -699,7 +699,7 @@ subroutine locate_source()
 
     ! user info for many source points
     if (.not. SHOW_DETAILS_LOCATE_SOURCE .and. NSOURCES >= 10) then
-      write(IMAIN,*) 'for all following ',NSOURCES-1,'sources we will suppress the details to avoid getting to long here...'
+      write(IMAIN,*) 'for all following sources we will suppress the detail infos to avoid getting to long here...'
       write(IMAIN,*)
       write(IMAIN,*) 'In case you wish to see detailed source informations for all source points, consider turning on the'
       write(IMAIN,*) 'parameter SHOW_DETAILS_LOCATE_SOURCE in setup/constants.h'
@@ -759,6 +759,12 @@ subroutine locate_source()
         write(IMAIN,*) '  total moment magnitude Mw = ', Mw
       endif
       write(IMAIN,*)
+      ! statistics
+      write(IMAIN,*) '  source ranges:'
+      write(IMAIN,*) '    depth        : min/max = ',sngl(minval(depth(:))/1000.d0),'/',sngl(maxval(depth(:))/1000.d0),' km'
+      write(IMAIN,*) '    half duration: min/max = ',sngl(minval(hdur(:))),'/',sngl(maxval(hdur(:)))
+      write(IMAIN,*) '    time shift   : min/max = ',sngl(minval(tshift_src(:))),'/',sngl(maxval(tshift_src(:)))
+      write(IMAIN,*)
       write(IMAIN,*) '********************'
       call flush_IMAIN()
     endif

diff --git a/src/specfem3D/wavefield_discontinuity_solver.f90 b/src/specfem3D/wavefield_discontinuity_solver.f90
@@ -166,12 +166,12 @@ end subroutine transfer_wavefield_discontinuity_to_GPU
   subroutine prepare_wavefield_discontinuity_GPU()
   use specfem_par, only: Mesh_pointer
   implicit none
-  call prepare_wavefield_discontinuity_device(Mesh_pointer, ispec_to_elem_wd,&
-                                                nglob_wd, nspec_wd, ibool_wd,&
-                                                boundary_to_iglob_wd,&
-                                                mass_in_wd,&
-                                                nfaces_wd, face_ijk_wd,&
-                                                face_ispec_wd, face_normal_wd,&
+  call prepare_wavefield_discontinuity_device(Mesh_pointer, ispec_to_elem_wd, &
+                                                nglob_wd, nspec_wd, ibool_wd, &
+                                                boundary_to_iglob_wd, &
+                                                mass_in_wd, &
+                                                nfaces_wd, face_ijk_wd, &
+                                                face_ispec_wd, face_normal_wd, &
                                                 face_jacobian2dw_wd)
   end subroutine prepare_wavefield_discontinuity_GPU