[cig-commits] r22854 - seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda
danielpeter at geodynamics.org
danielpeter at geodynamics.org
Thu Sep 26 06:24:53 PDT 2013
Author: danielpeter
Date: 2013-09-26 06:24:53 -0700 (Thu, 26 Sep 2013)
New Revision: 22854
Modified:
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_scalar_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_vector_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_add_sources_elastic_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_coupling_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_crust_mantle_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_inner_core_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_outer_core_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_kernels_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/mesh_constants_cuda.h
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/transfer_fields_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/update_displacement_cuda.cu
Log:
updates cuda debugging for backward wavefields
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_scalar_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_scalar_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_scalar_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -104,7 +104,7 @@
}
else if(*FORWARD_OR_ADJOINT == 3) {
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_ASSEMBLY();
prepare_boundary_potential_on_device<<<grid,threads>>>(mp->d_b_accel_outer_core,
mp->d_b_send_accel_buffer_outer_core,
@@ -194,7 +194,7 @@
}
else if(*FORWARD_OR_ADJOINT == 3) {
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_ASSEMBLY();
// copies scalar buffer onto GPU
print_CUDA_error_if_any(cudaMemcpy(mp->d_b_send_accel_buffer_outer_core, buffer_recv_scalar,
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_vector_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_vector_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/assemble_MPI_vector_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -110,7 +110,7 @@
}
else if(*FORWARD_OR_ADJOINT == 3) {
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_ASSEMBLY();
prepare_boundary_accel_on_device<<<grid,threads>>>(mp->d_b_accel_crust_mantle,
mp->d_b_send_accel_buffer_crust_mantle,
@@ -157,7 +157,7 @@
}
else if(*FORWARD_OR_ADJOINT == 3) {
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_ASSEMBLY();
prepare_boundary_accel_on_device<<<grid,threads>>>(mp->d_b_accel_inner_core,
mp->d_b_send_accel_buffer_inner_core,
@@ -249,7 +249,7 @@
}
else if(*FORWARD_OR_ADJOINT == 3) {
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_ASSEMBLY();
// copies vector buffer values to GPU
print_CUDA_error_if_any(cudaMemcpy(mp->d_b_send_accel_buffer_crust_mantle, buffer_recv_vector,
@@ -299,7 +299,7 @@
}
else if(*FORWARD_OR_ADJOINT == 3) {
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_ASSEMBLY();
// copies buffer values to GPU
print_CUDA_error_if_any(cudaMemcpy(mp->d_b_send_accel_buffer_inner_core, buffer_recv_vector,
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_add_sources_elastic_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_add_sources_elastic_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_add_sources_elastic_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -138,7 +138,7 @@
double* h_stf_pre_compute) {
TRACE("compute_add_sources_backward_cuda");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_SOURCES();
Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_coupling_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_coupling_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_coupling_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -137,7 +137,7 @@
mp->nspec2D_top_outer_core);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_COUPLING();
// adjoint simulations
compute_coupling_fluid_CMB_kernel<<<grid,threads>>>(mp->d_b_displ_crust_mantle,
@@ -257,7 +257,7 @@
mp->nspec2D_bottom_outer_core);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_COUPLING();
// adjoint simulations
compute_coupling_fluid_ICB_kernel<<<grid,threads>>>(mp->d_b_displ_inner_core,
@@ -387,7 +387,7 @@
mp->nspec2D_bottom_crust_mantle);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_COUPLING();
// adjoint simulations
compute_coupling_CMB_fluid_kernel<<<grid,threads>>>(mp->d_b_displ_crust_mantle,
@@ -520,7 +520,7 @@
mp->nspec2D_top_inner_core);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_COUPLING();
// adjoint simulations
compute_coupling_ICB_fluid_kernel<<<grid,threads>>>(mp->d_b_displ_inner_core,
@@ -634,7 +634,7 @@
mp->d_normal_ocean_load);
}else if( *FORWARD_OR_ADJOINT == 3){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_COUPLING();
// for backward/reconstructed potentials
compute_coupling_ocean_cuda_kernel<<<grid,threads>>>(mp->d_b_accel_crust_mantle,
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_crust_mantle_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_crust_mantle_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_crust_mantle_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -1394,7 +1394,7 @@
mp->NSPEC_CRUST_MANTLE_STRAIN_ONLY);
}else if( FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_FORCES();
Kernel_2_crust_mantle_impl<<< grid,threads>>>(nb_blocks_to_compute,
mp->NGLOB_CRUST_MANTLE,
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_inner_core_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_inner_core_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_inner_core_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -1020,7 +1020,7 @@
mp->NSPEC_INNER_CORE);
}else if( FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_FORCES();
Kernel_2_inner_core_impl<<< grid,threads>>>(nb_blocks_to_compute,
mp->NGLOB_INNER_CORE,
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_outer_core_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_outer_core_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_forces_outer_core_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -523,7 +523,7 @@
mp->NSPEC_OUTER_CORE);
}else if( FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_FORCES();
Kernel_2_outer_core_impl<<<grid,threads>>>(nb_blocks_to_compute,
mp->NGLOB_OUTER_CORE,
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_kernels_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_kernels_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_kernels_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -225,7 +225,7 @@
TRACE("compute_kernels_cm_cuda");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_KERNEL();
Mesh* mp = (Mesh*)(*Mesh_pointer); //get mesh pointer out of fortran integer container
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/mesh_constants_cuda.h
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/mesh_constants_cuda.h 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/mesh_constants_cuda.h 2013-09-26 13:24:53 UTC (rev 22854)
@@ -70,12 +70,24 @@
#define PRINT5(var,offset) // for(i=0;i<10;i++) printf("var(%d)=%f\n",i,var[offset+i]);
#endif
-// daniel debug: run backward simulations with empty arrays to check
+// daniel debug: run backward simulations with/without gpu routines and empty arrays for debugging
#define DEBUG_BACKWARD_SIMULATIONS 0
#if DEBUG_BACKWARD_SIMULATIONS == 1
-#define DEBUG_EMPTY_BACKWARD() return;
+#define DEBUG_BACKWARD_ASSEMBLY() return;
+#define DEBUG_BACKWARD_COUPLING() return;
+#define DEBUG_BACKWARD_FORCES() return;
+#define DEBUG_BACKWARD_KERNEL() return;
+#define DEBUG_BACKWARD_SOURCES() return;
+#define DEBUG_BACKWARD_TRANSFER() return;
+#define DEBUG_BACKWARD_UPDATE() return;
#else
-#define DEBUG_EMPTY_BACKWARD()
+#define DEBUG_BACKWARD_ASSEMBLY()
+#define DEBUG_BACKWARD_COUPLING()
+#define DEBUG_BACKWARD_FORCES()
+#define DEBUG_BACKWARD_KERNEL()
+#define DEBUG_BACKWARD_SOURCES()
+#define DEBUG_BACKWARD_TRANSFER()
+#define DEBUG_BACKWARD_UPDATE()
#endif
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/transfer_fields_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/transfer_fields_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/transfer_fields_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -107,7 +107,7 @@
TRACE("transfer_fields_b_cm_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_crust_mantle,b_displ,sizeof(realw)*(*size),cudaMemcpyHostToDevice),40003);
@@ -124,7 +124,7 @@
TRACE("transfer_fields_b_ic_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_inner_core,b_displ,sizeof(realw)*(*size),cudaMemcpyHostToDevice),40003);
@@ -141,7 +141,7 @@
TRACE("transfer_fields_b_oc_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_outer_core,b_displ,sizeof(realw)*(*size),cudaMemcpyHostToDevice),40003);
@@ -517,7 +517,7 @@
realw* epsilondev_yz) {
TRACE("transfer_b_strain_cm_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
//get mesh pointer out of fortran integer container
Mesh* mp = (Mesh*)(*Mesh_pointer);
@@ -584,7 +584,7 @@
realw* epsilondev_yz) {
TRACE("transfer_b_strain_cm_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
//get mesh pointer out of fortran integer container
Mesh* mp = (Mesh*)(*Mesh_pointer);
@@ -622,7 +622,7 @@
realw* b_R_yz) {
TRACE("transfer_b_Rmemory_cm_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
//get mesh pointer out of fortran integer container
Mesh* mp = (Mesh*)(*Mesh_pointer);
@@ -657,7 +657,7 @@
realw* b_R_yz) {
TRACE("transfer_b_rmemory_ic_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
//get mesh pointer out of fortran integer container
Mesh* mp = (Mesh*)(*Mesh_pointer);
@@ -715,7 +715,7 @@
realw* B_array_rotation) {
TRACE("transfer_b_rotation_to_device");
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_TRANSFER();
//get mesh pointer out of fortran integer container
Mesh* mp = (Mesh*)(*Mesh_pointer);
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/update_displacement_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/update_displacement_cuda.cu 2013-09-26 01:23:42 UTC (rev 22853)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/update_displacement_cuda.cu 2013-09-26 13:24:53 UTC (rev 22854)
@@ -113,7 +113,7 @@
size,deltat,deltatsqover2,deltatover2);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
// kernel for backward fields
UpdateDispVeloc_kernel<<<grid,threads>>>(mp->d_b_displ_inner_core,
@@ -180,7 +180,7 @@
size,deltat,deltatsqover2,deltatover2);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
// kernel for backward fields
UpdateDispVeloc_kernel<<<grid,threads>>>(mp->d_b_displ_crust_mantle,
@@ -278,7 +278,7 @@
size,deltat,deltatsqover2,deltatover2);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
UpdatePotential_kernel<<<grid,threads>>>(mp->d_b_displ_outer_core,
mp->d_b_veloc_outer_core,
@@ -384,7 +384,7 @@
mp->d_rmassz_crust_mantle);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
multiply_accel_elastic_cuda_device<<< grid, threads>>>(mp->d_b_accel_crust_mantle,
mp->d_b_veloc_crust_mantle,
@@ -413,7 +413,7 @@
mp->d_rmassz_inner_core);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
multiply_accel_elastic_cuda_device<<< grid, threads>>>(mp->d_b_accel_inner_core,
mp->d_b_veloc_inner_core,
@@ -567,7 +567,7 @@
mp->d_rmass_outer_core);
}else if( *FORWARD_OR_ADJOINT == 3 ){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
multiply_accel_acoustic_cuda_device<<< grid, threads>>>(mp->d_b_accel_outer_core,
mp->NGLOB_OUTER_CORE,
@@ -630,7 +630,7 @@
deltatover2);
}else if( *FORWARD_OR_ADJOINT == 3){
// debug
- DEBUG_EMPTY_BACKWARD();
+ DEBUG_BACKWARD_UPDATE();
update_veloc_acoustic_cuda_device<<< grid, threads>>>(mp->d_b_veloc_outer_core,
mp->d_b_accel_outer_core,
More information about the CIG-COMMITS
mailing list