bug fix for noise simulations on GPUs

SPECFEM · Jul 8, 2015 · fd75881 · fd75881
1 parent 1b302a5
commit fd75881
Show file tree

Hide file tree

Showing 100 changed files with 1,890 additions and 3,478 deletions.
diff --git a/EXAMPLES/regional_Greece_noise_small/image_alpha_kernel.jpg b/EXAMPLES/regional_Greece_noise_small/image_alpha_kernel.jpg
diff --git a/EXAMPLES/regional_Greece_noise_small/image_beta_kernel.jpg b/EXAMPLES/regional_Greece_noise_small/image_beta_kernel.jpg
diff --git a/EXAMPLES/regional_Greece_noise_small/image_rho_kernel.jpg b/EXAMPLES/regional_Greece_noise_small/image_rho_kernel.jpg
diff --git a/EXAMPLES/regional_Greece_noise_small/state_alpha_kernel.pvsm b/EXAMPLES/regional_Greece_noise_small/state_alpha_kernel.pvsm
diff --git a/EXAMPLES/regional_Greece_noise_small/xcombine_vol_data.sh b/EXAMPLES/regional_Greece_noise_small/xcombine_vol_data.sh
@@ -26,12 +26,15 @@ cp -v ~/SPECFEM3D_GLOBE/utils/Visualization/VTK_ParaView/paraviewpython-example.
 echo
 echo "alpha_kernel"
 echo
-./bin/xcombine_vol_data $slice alpha_kernel $dir $dir OUTPUT_FILES/ $res 1 > tmp.log
-mesh2vtu.pl -i OUTPUT_FILES/reg_1_alpha_kernel.mesh -o OUTPUT_FILES/reg_1_alpha_kernel.vtu >> tmp.log
+./bin/xcombine_vol_data $slice alpha_kernel $dir $dir OUTPUT_FILES/ $res 1 | tee tmp.log
+mesh2vtu.pl -i OUTPUT_FILES/reg_1_alpha_kernel.mesh -o OUTPUT_FILES/reg_1_alpha_kernel.vtu | tee -a tmp.log
 rm -f OUTPUT_FILES/reg_*alpha*.mesh
 min=`grep "min/max" tmp.log | awk '{print $3 }' | sort | head -n 1`
 max=`grep "min/max" tmp.log | awk '{print $4 }' | sort | tail -n 1`
+echo 
+echo "statistics:"
 echo "  alpha_kernel min/max: $min $max"
+echo
 
 ./paraviewpython-example.py state_alpha_kernel.pvsm
 mv image.jpg image_alpha_kernel.jpg
@@ -40,12 +43,15 @@ echo
 echo "beta_kernel"
 echo
 # only for crust_mantle region
-./bin/xcombine_vol_data $slice beta_kernel $dir $dir OUTPUT_FILES/ $res 1 > tmp.log
-mesh2vtu.pl -i OUTPUT_FILES/reg_1_beta_kernel.mesh -o OUTPUT_FILES/reg_1_beta_kernel.vtu >> tmp.log
+./bin/xcombine_vol_data $slice beta_kernel $dir $dir OUTPUT_FILES/ $res 1 | tee tmp.log
+mesh2vtu.pl -i OUTPUT_FILES/reg_1_beta_kernel.mesh -o OUTPUT_FILES/reg_1_beta_kernel.vtu | tee -a tmp.log
 rm -f OUTPUT_FILES/reg_*beta*.mesh
 min=`grep "min/max" tmp.log | awk '{print $3 }' | sort | head -n 1`
 max=`grep "min/max" tmp.log | awk '{print $4 }' | sort | tail -n 1`
+echo
+echo "statistics:"
 echo "  beta_kernel min/max: $min $max"
+echo
 
 sed "s:alpha:beta:g" state_alpha_kernel.pvsm > tmp_beta.pvsm
 ./paraviewpython-example.py tmp_beta.pvsm
@@ -54,12 +60,15 @@ mv image.jpg image_beta_kernel.jpg
 echo
 echo "rho_kernel"
 echo
-./bin/xcombine_vol_data $slice rho_kernel $dir $dir OUTPUT_FILES/ $res 1 > tmp.log
-mesh2vtu.pl -i OUTPUT_FILES/reg_1_rho_kernel.mesh -o OUTPUT_FILES/reg_1_rho_kernel.vtu >> tmp.log
+./bin/xcombine_vol_data $slice rho_kernel $dir $dir OUTPUT_FILES/ $res 1 | tee tmp.log
+mesh2vtu.pl -i OUTPUT_FILES/reg_1_rho_kernel.mesh -o OUTPUT_FILES/reg_1_rho_kernel.vtu | tee -a tmp.log
 rm -f OUTPUT_FILES/reg_*rho*.mesh
 min=`grep "min/max" tmp.log | awk '{print $3 }' | sort | head -n 1`
 max=`grep "min/max" tmp.log | awk '{print $4 }' | sort | tail -n 1`
+echo
+echo "statistics:"
 echo "  rho_kernel min/max: $min $max"
+echo
 
 sed "s:alpha:rho:g" state_alpha_kernel.pvsm > tmp_rho.pvsm
 ./paraviewpython-example.py tmp_rho.pvsm

diff --git a/Makefile.in b/Makefile.in
@@ -61,17 +61,20 @@ FCLINK = $(MPIFCCOMPILE_CHECK)
 ####
 #######################################
 
+# Reduce GPU-register pressure by limited the number of thread spread
+# (GPU for embedded devices are not powerful enough for big kernels)
+# Must match mesh_constants_gpu.h::GPU_ELEM_PER_THREAD
+GPU_ELEM_PER_THREAD := 1
+
+##
+## CUDA
+##
 @COND_CUDA_TRUE@CUDA = yes
 @COND_CUDA_FALSE@CUDA = no
 
 @COND_CUDA5_TRUE@CUDA5 = yes
 @COND_CUDA5_FALSE@CUDA5 = no
 
-@COND_OCL_TRUE@OCL = yes
-@COND_OCL_FALSE@OCL = no
-
-MPI_INCLUDES = @MPI_INCLUDES@
-
 CUDA_FLAGS = @CUDA_FLAGS@
 CUDA_INC = @CUDA_CPPFLAGS@
 CUDA_LINK = @CUDA_LDFLAGS@ @CUDA_LIBS@
@@ -80,30 +83,9 @@ CUDA_DEBUG = --cudart=shared
 @COND_CUDA_TRUE@NVCC = nvcc
 @COND_CUDA_FALSE@NVCC = @CC@
 
-OCL_CPU_FLAGS = @OCL_CPU_FLAGS@
-OCL_GPU_FLAGS = @OCL_GPU_FLAGS@
-
-OCL_INC = @OCL_CFLAGS@
-OCL_LINK = @OCL_LDFLAGS@ @OCL_LIBS@
-
-ifeq ($(OCL), yes)
-  ifeq ($(CUDA), yes)
-    GPU_CUDA_AND_OCL = yes
-  endif
-endif
-
-ifeq ($(OCL), no)
-  ifeq ($(CUDA), no)
-    NO_GPU = yes
-  endif
-endif
-
-ifneq ($(NO_GPU), yes)
-  HAS_GPU = yes
-endif
-
-# GPU architecture
-
+##
+## GPU architecture
+##
 # CUDA architecture / code version
 # Fermi: -gencode=arch=compute_10,code=sm_10 not supported
 # Tesla (default): -gencode=arch=compute_20,code=sm_20
@@ -130,6 +112,37 @@ GENCODE_35 = -gencode=arch=compute_35,code=\"sm_35,compute_35\"
 @COND_CUDA_FALSE@NVCC_FLAGS =
 @COND_CUDA_FALSE@NVCCLINK = $(NVCC) $(NVCC_FLAGS)
 
+##
+## OpenCL
+##
+@COND_OCL_TRUE@OCL = yes
+@COND_OCL_FALSE@OCL = no
+
+OCL_CPU_FLAGS = @OCL_CPU_FLAGS@
+OCL_GPU_FLAGS = @OCL_GPU_FLAGS@
+
+OCL_INC = @OCL_CFLAGS@
+OCL_LINK = @OCL_LDFLAGS@ @OCL_LIBS@
+
+ifeq ($(OCL), yes)
+  ifeq ($(CUDA), yes)
+    GPU_CUDA_AND_OCL = yes
+  endif
+endif
+
+ifeq ($(OCL), no)
+  ifeq ($(CUDA), no)
+    NO_GPU = yes
+  endif
+endif
+
+ifneq ($(NO_GPU), yes)
+  HAS_GPU = yes
+endif
+
+## MPI directories for CUDA / OpenCL
+MPI_INCLUDES = @MPI_INCLUDES@
+
 #######################################
 ####
 #### VTK

diff --git a/src/gpu/boast/noise_add_surface_movie_kernel.rb b/src/gpu/boast/noise_add_surface_movie_kernel.rb
@@ -62,7 +62,7 @@ def BOAST::noise_add_surface_movie_kernel(ref = true, n_dim = 3, n_gllx = 5, n_g
             print eta === eta + noise_surface_movie[INDEX3(ndim,ngll2,indx,igll,iface)]*normal[indx]
           }
 
-          print jacobianw === wgllwgll[k*ngllx+i]*jacobian2D[igll+ngll2*iface]
+          print jacobianw === wgllwgll[j*ngllx+i]*jacobian2D[igll+ngll2*iface]
 
           (0..2).each { |indx|
             print atomicAdd(accel+ iglob*3 + indx, eta*mask_noise[ipoin]*normal[indx]*jacobianw)

diff --git a/src/gpu/boast/noise_transfer_surface_to_host_kernel.rb b/src/gpu/boast/noise_transfer_surface_to_host_kernel.rb
@@ -39,7 +39,7 @@ def BOAST::noise_transfer_surface_to_host_kernel(ref = true, n_dim = 3, n_gllx =
           print i === igll-j*ngllx
           print iglob === ibool[INDEX4(ngllx,ngllx,ngllx,i,j,k,ispec)]-1
           (0..2).each { |indx|
-             print noise_surface_movie[INDEX3(ndim,ngll2,indx,igll,iface)] === displ[iglob*3]
+             print noise_surface_movie[INDEX3(ndim,ngll2,indx,igll,iface)] === displ[iglob*3 + indx]
           }
         }
       close p

diff --git a/src/gpu/boast/references/noise_add_surface_movie_kernel.cu b/src/gpu/boast/references/noise_add_surface_movie_kernel.cu
@@ -43,7 +43,7 @@ __global__ void noise_add_surface_movie_kernel(realw* accel,
                  noise_surface_movie[INDEX3(NDIM,NGLL2,2,igll,iface)]*normal_z);
 
     // weighted jacobian
-    realw jacobianw = wgllwgll[k*NGLLX+i]*jacobian2D[igll+NGLL2*iface];
+    realw jacobianw = wgllwgll[j*NGLLX+i]*jacobian2D[igll+NGLL2*iface];
 
     // note: check error from cuda-memcheck and ddt seems "incorrect", because we
     //          are passing a __constant__ variable pointer around like it was

diff --git a/src/gpu/kernels.gen/assemble_boundary_accel_on_device.cu b/src/gpu/kernels.gen/assemble_boundary_accel_on_device.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/assemble_boundary_accel_on_device_cl.c b/src/gpu/kernels.gen/assemble_boundary_accel_on_device_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/assemble_boundary_potential_on_device.cu b/src/gpu/kernels.gen/assemble_boundary_potential_on_device.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/assemble_boundary_potential_on_device_cl.c b/src/gpu/kernels.gen/assemble_boundary_potential_on_device_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_acoustic_kernel.cu b/src/gpu/kernels.gen/compute_acoustic_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_acoustic_kernel_cl.c b/src/gpu/kernels.gen/compute_acoustic_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_add_sources_adjoint_kernel.cu b/src/gpu/kernels.gen/compute_add_sources_adjoint_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_add_sources_adjoint_kernel_cl.c b/src/gpu/kernels.gen/compute_add_sources_adjoint_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_add_sources_kernel.cu b/src/gpu/kernels.gen/compute_add_sources_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_add_sources_kernel_cl.c b/src/gpu/kernels.gen/compute_add_sources_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_ani_kernel.cu b/src/gpu/kernels.gen/compute_ani_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_ani_kernel_cl.c b/src/gpu/kernels.gen/compute_ani_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_ani_undoatt_kernel.cu b/src/gpu/kernels.gen/compute_ani_undoatt_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_ani_undoatt_kernel_cl.c b/src/gpu/kernels.gen/compute_ani_undoatt_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel.cu b/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel_cl.c b/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel.cu b/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel_cl.c b/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_fluid_CMB_kernel.cu b/src/gpu/kernels.gen/compute_coupling_fluid_CMB_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_fluid_CMB_kernel_cl.c b/src/gpu/kernels.gen/compute_coupling_fluid_CMB_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_fluid_ICB_kernel.cu b/src/gpu/kernels.gen/compute_coupling_fluid_ICB_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_fluid_ICB_kernel_cl.c b/src/gpu/kernels.gen/compute_coupling_fluid_ICB_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_ocean_kernel.cu b/src/gpu/kernels.gen/compute_coupling_ocean_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_coupling_ocean_kernel_cl.c b/src/gpu/kernels.gen/compute_coupling_ocean_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_hess_kernel.cu b/src/gpu/kernels.gen/compute_hess_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_hess_kernel_cl.c b/src/gpu/kernels.gen/compute_hess_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_iso_kernel.cu b/src/gpu/kernels.gen/compute_iso_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_iso_kernel_cl.c b/src/gpu/kernels.gen/compute_iso_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_iso_undoatt_kernel.cu b/src/gpu/kernels.gen/compute_iso_undoatt_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_iso_undoatt_kernel_cl.c b/src/gpu/kernels.gen/compute_iso_undoatt_kernel_cl.c
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*

diff --git a/src/gpu/kernels.gen/compute_rho_kernel.cu b/src/gpu/kernels.gen/compute_rho_kernel.cu
@@ -1,5 +1,5 @@
 //note: please do not modify this file manually!
-//      this file has been generated automatically by BOAST version 1.0.1
+//      this file has been generated automatically by BOAST version 1.0.3
 //      by: make boast_kernels
 
 /*