Corrected GPU smoothing routine

zenghy789 · Jun 13, 2018 · 6dac0c6 · 6dac0c6
1 parent ca16296
commit 6dac0c6
Show file tree

Hide file tree

Showing 2 changed files with 5 additions and 4 deletions.
diff --git a/src/tomography/postprocess_sensitivity_kernels/smooth_cuda.cu b/src/tomography/postprocess_sensitivity_kernels/smooth_cuda.cu
@@ -64,11 +64,12 @@ z_me = zstore_me[NGLL3*ispec + igll ];
 sh_wgll_cube[igll]=wgll_cube[igll];
 __syncthreads();
 
-dat=0;
-normalisation_slice=0;
+dat=0.f;
+normalisation_slice=0.f;
 //We test 125 spectral elements at a time
 for (int i=0;i<n_loop;i++)
 {
+__syncthreads();
 if (NGLL3*i + threadIdx.x < nspec_other){
 x_other = (xstore_other[i*NGLL3*NGLL3 + threadIdx.x*NGLL3 ] + xstore_other[i*NGLL3*NGLL3 + threadIdx.x*NGLL3 + NGLL3 - 1 ])/2;
 y_other = (ystore_other[i*NGLL3*NGLL3 + threadIdx.x*NGLL3 ] + ystore_other[i*NGLL3*NGLL3 + threadIdx.x*NGLL3 + NGLL3 - 1 ])/2;

diff --git a/src/tomography/postprocess_sensitivity_kernels/smooth_sem.F90 b/src/tomography/postprocess_sensitivity_kernels/smooth_sem.F90
@@ -416,7 +416,7 @@ program smooth_sem
   ! frees memory
   deallocate(xstore,ystore,zstore)
   deallocate(xix,xiy,xiz,etax,etay,etaz,gammax,gammay,gammaz)
-  deallocate(ibool)
+  deallocate(ibool,irregular_element_number)
   deallocate(jacobian)
 
   call synchronize_all()
@@ -673,7 +673,7 @@ program smooth_sem
     deallocate(ibool)
 
     ! data file
-    write(prname,'(a,i6.6,a)') trim(input_dir)//'proc',iproc,'_'
+    write(prname,'(a,i6.6,a)') trim(input_dir)//'/proc',iproc,'_'
     local_data_file = trim(prname) // trim(kernel_name) // '.bin'
 
     open(unit = IIN,file = trim(local_data_file),status='old',action='read',form ='unformatted',iostat=ier)