[cig-commits] r19689 - in seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src: cuda specfem3D
joseph.charles at geodynamics.org
joseph.charles at geodynamics.org
Mon Feb 27 16:02:27 PST 2012
Author: joseph.charles
Date: 2012-02-27 16:02:26 -0800 (Mon, 27 Feb 2012)
New Revision: 19689
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/mesh_constants_cuda.h
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/prepare_mesh_constants_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/transfer_fields_cuda.cu
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_acoustic.F90
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_elastic.F90
seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/prepare_timerun.f90
Log:
updates coupling routines
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 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/compute_coupling_cuda.cu 2012-02-28 00:02:26 UTC (rev 19689)
@@ -637,11 +637,9 @@
dim3 grid(num_blocks_x,num_blocks_y);
dim3 threads(5,5,1);
-
-
+
// initializes temporary array to zero
- print_CUDA_error_if_any(cudaMemset(mp->d_updated_dof_ocean_load,0,
- sizeof(int)*mp->NGLOB_AB),88501);
+ print_CUDA_error_if_any(cudaMemset(mp->d_updated_dof_ocean_load,0,sizeof(int)*mp->NGLOB_CRUST_MANTLE_OCEANS),88501);
#ifdef ENABLE_VERY_SLOW_ERROR_CHECKING
exit_on_cuda_error("before kernel compute_coupling_ocean_cuda");
@@ -660,7 +658,7 @@
if( mp->simulation_type == 3 ) {
// re-initializes array
print_CUDA_error_if_any(cudaMemset(mp->d_updated_dof_ocean_load,0,
- sizeof(int)*mp->NGLOB_AB),88502);
+ sizeof(int)*mp->NGLOB_CRUST_MANTLE_OCEANS),88502);
compute_coupling_ocean_cuda_kernel<<<grid,threads>>>(mp->d_b_accel_crust_mantle,
mp->d_rmass_crust_mantle,
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 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/mesh_constants_cuda.h 2012-02-28 00:02:26 UTC (rev 19689)
@@ -169,7 +169,6 @@
// ------------------------------------------------------------------ //
int NSPEC_CRUST_MANTLE;
int NGLOB_CRUST_MANTLE;
- int NGLOB_CRUST_MANTLE_OCEANS;
// interpolators
realw* d_xix_crust_mantle; realw* d_xiy_crust_mantle; realw* d_xiz_crust_mantle;
@@ -419,6 +418,17 @@
int* d_ibelm_top_inner_core;
// ------------------------------------------------------------------ //
+ // oceans
+ // ------------------------------------------------------------------ //
+ int NGLOB_CRUST_MANTLE_OCEANS;
+
+ // model parameter
+ realw* d_rmass_ocean_load;
+
+ // temporary global array: used to synchronize updates on global accel array
+ int* d_updated_dof_ocean_load;
+
+ // ------------------------------------------------------------------ //
// attenuation
// ------------------------------------------------------------------ //
realw* d_alphaval;
@@ -454,6 +464,7 @@
int anisotropic_3D_mantle;
int gravity;
int rotation;
+ int oceans;
int anisotropic_inner_core;
int save_boundary_mesh;
@@ -587,7 +598,6 @@
int noise_tomography;
int nspec_top;
- int* d_ibelm_top_crust_mantle;
realw* d_noise_surface_movie;
realw* d_noise_sourcearray;
@@ -749,11 +759,6 @@
// approximative hessian for preconditioning kernels
realw* d_hess_el_kl;
- // oceans
- realw* d_rmass_ocean_load;
- realw* d_free_surface_normal;
- int* d_updated_dof_ocean_load;
-
// ------------------------------------------------------------------ //
// acoustic wavefield
// ------------------------------------------------------------------ //
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/prepare_mesh_constants_cuda.cu
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/prepare_mesh_constants_cuda.cu 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/prepare_mesh_constants_cuda.cu 2012-02-28 00:02:26 UTC (rev 19689)
@@ -203,6 +203,7 @@
int* NOISE_TOMOGRAPHY,
int* SAVE_FORWARD_f,
int* ABSORBING_CONDITIONS_f,
+ int* OCEANS_f,
int* GRAVITY_f,
int* ROTATION_f,
int* ATTENUATION_f,
@@ -254,6 +255,7 @@
// simulation flags initialization
mp->save_forward = *SAVE_FORWARD_f;
mp->absorbing_conditions = *ABSORBING_CONDITIONS_f;
+ mp->oceans = *OCEANS_f;
mp->gravity = *GRAVITY_f;
mp->rotation = *ROTATION_f;
mp->attenuation = *ATTENUATION_f;
@@ -2017,10 +2019,37 @@
#endif
}
+/* ----------------------------------------------------------------------------------------------- */
+// OCEANS
+/* ----------------------------------------------------------------------------------------------- */
+extern "C"
+void FC_FUNC_(prepare_oceans_device,
+ PREPARE_OCEANS_DEVICE)(long* Mesh_pointer_f,
+ realw* h_rmass_ocean_load) {
+ TRACE("prepare_oceans_device");
+
+ Mesh* mp = (Mesh*)(*Mesh_pointer_f);
+
+ // mass matrix
+ print_CUDA_error_if_any(cudaMalloc((void**)&(mp->d_rmass_ocean_load),
+ sizeof(realw)*mp->NGLOB_CRUST_MANTLE_OCEANS),4501);
+ print_CUDA_error_if_any(cudaMemcpy(mp->d_rmass_ocean_load,h_rmass_ocean_load,
+ sizeof(realw)*mp->NGLOB_CRUST_MANTLE_OCEANS,cudaMemcpyHostToDevice),4502);
+
+ // temporary global array: used to synchronize updates on global accel array
+ print_CUDA_error_if_any(cudaMalloc((void**)&(mp->d_updated_dof_ocean_load),
+ sizeof(int)*mp->NGLOB_CRUST_MANTLE_OCEANS),4502);
+
+
+#ifdef ENABLE_VERY_SLOW_ERROR_CHECKING
+ exit_on_cuda_error("prepare_oceans_device");
+#endif
+}
+
/* ----------------------------------------------------------------------------------------------- */
// for ELASTIC simulations
@@ -2777,6 +2806,11 @@
cudaFree(mp->d_muvstore_inner_core);
cudaFree(mp->d_ibool_inner_core);
+ if( mp->oceans ){
+ cudaFree(mp->d_rmass_ocean_load);
+ cudaFree(mp->d_updated_dof_ocean_load);
+ }
+
if( mp->gravity ){
cudaFree(mp->d_xstore_inner_core);
cudaFree(mp->d_ystore_inner_core);
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 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/cuda/transfer_fields_cuda.cu 2012-02-28 00:02:26 UTC (rev 19689)
@@ -96,57 +96,6 @@
}
-// coupling fluid_CMB_ICB
-extern "C"
-void FC_FUNC_(transfer_coupling_fields_fluid_cmb_icb_to_device,
- TRANSFER_COUPLING_FIELDS_FLUID_CMB_ICB_TO_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* displ_cm, realw* displ_ic, realw* accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_fields_fluid_cmb_icb_to_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_accel_outer_core,accel,sizeof(realw)*(*size_oc),cudaMemcpyHostToDevice),40009);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_displ_crust_mantle,displ_cm,sizeof(realw)*(*size_cm),cudaMemcpyHostToDevice),40010);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_displ_inner_core,displ_ic,sizeof(realw)*(*size_ic),cudaMemcpyHostToDevice),40011);
-
-}
-
-// coupling CMB_ICB_fluid
-extern "C"
-void FC_FUNC_(transfer_coupling_fields_cmb_icb_fluid_to_device,
- TRANSFER_COUPLING_FIELDS_CMB_ICB_FLUID_TO_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* displ_cm, realw* displ_ic, realw* accel_cm, realw* accel_ic, realw* accel_oc, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_fields_cmb_icb_fluid_to_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_accel_outer_core,accel_oc,sizeof(realw)*(*size_oc),cudaMemcpyHostToDevice),40009);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_accel_crust_mantle,accel_cm,sizeof(realw)*(*size_cm),cudaMemcpyHostToDevice),40010);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_accel_inner_core,accel_ic,sizeof(realw)*(*size_ic),cudaMemcpyHostToDevice),40010);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_displ_crust_mantle,displ_cm,sizeof(realw)*(*size_cm),cudaMemcpyHostToDevice),40012);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_displ_inner_core,displ_ic,sizeof(realw)*(*size_ic),cudaMemcpyHostToDevice),40013);
-
-}
-
-// coupling CMB_ocean
-extern "C"
-void FC_FUNC_(transfer_coupling_fields_cmb_ocean_to_device,
- TRANSFER_COUPLING_FIELDS_CMB_OCEAN_TO_DEVICE)(int* size, realw* accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_fields_cmb_ocean_to_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_accel_crust_mantle,accel,sizeof(realw)*(*size),cudaMemcpyHostToDevice),40009);
-
-}
-
/* ----------------------------------------------------------------------------------------------- */
// backward/reconstructed fields
@@ -196,56 +145,6 @@
}
-// coupling fluid_CMB_ICB
-extern "C"
-void FC_FUNC_(transfer_coupling_b_fields_fluid_cmb_icb_to_device,
- TRANSFER_COUPLING_B_FIELDS_FLUID_CMB_ICB_TO_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* b_displ_cm, realw* b_displ_ic, realw* b_accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_b_fields_fluid_cmb_icb_to_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_accel_outer_core,b_accel,sizeof(realw)*(*size_oc),cudaMemcpyHostToDevice),40009);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_crust_mantle,b_displ_cm,sizeof(realw)*(*size_cm),cudaMemcpyHostToDevice),40010);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_inner_core,b_displ_ic,sizeof(realw)*(*size_ic),cudaMemcpyHostToDevice),40011);
-
-}
-
-// coupling CMB_ICB_fluid
-extern "C"
-void FC_FUNC_(transfer_coupling_b_fields_cmb_icb_fluid_to_device,
- TRANSFER_COUPLING_B_FIELDS_CMB_ICB_FLUID_TO_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* b_displ_cm, realw* b_displ_ic, realw* b_accel_cm, realw* b_accel_ic, realw* b_accel_oc, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_b_fields_cmb_icb_fluid_to_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_accel_outer_core,b_accel_oc,sizeof(realw)*(*size_oc),cudaMemcpyHostToDevice),40009);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_accel_crust_mantle,b_accel_cm,sizeof(realw)*(*size_cm),cudaMemcpyHostToDevice),40010);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_accel_inner_core,b_accel_ic,sizeof(realw)*(*size_ic),cudaMemcpyHostToDevice),40011);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_crust_mantle,b_displ_cm,sizeof(realw)*(*size_cm),cudaMemcpyHostToDevice),40012);
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_displ_inner_core,b_displ_ic,sizeof(realw)*(*size_ic),cudaMemcpyHostToDevice),40013);
-}
-
-// coupling CMB_ocean
-extern "C"
-void FC_FUNC_(transfer_coupling_b_fields_cmb_ocean_to_device,
- TRANSFER_COUPLING_B_FIELDS_CMB_OCEAN_TO_DEVICE)(int* size, realw* b_accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_b_fields_cmb_ocean_to_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(mp->d_b_accel_crust_mantle,b_accel,sizeof(realw)*(*size),cudaMemcpyHostToDevice),40009);
-
-}
-
/* ----------------------------------------------------------------------------------------------- */
// transfer memory from GPU device to CPU host
@@ -301,57 +200,6 @@
}
-// coupling fluid_CMB_IMB
-extern "C"
-void FC_FUNC_(transfer_coupling_fields_fluid_cmb_icb_from_device,
- TRANSFER_COUPLING_FIELDS_FLUID_CMB_ICB_FROM_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* displ_cm, realw* displ_ic, realw* accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_fields_fluid_cmb_icb_from_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(accel,mp->d_accel_outer_core,sizeof(realw)*(*size_oc),cudaMemcpyDeviceToHost),40014);
-
- print_CUDA_error_if_any(cudaMemcpy(displ_cm,mp->d_displ_crust_mantle,sizeof(realw)*(*size_cm),cudaMemcpyDeviceToHost),40015);
-
- print_CUDA_error_if_any(cudaMemcpy(displ_ic,mp->d_displ_inner_core,sizeof(realw)*(*size_ic),cudaMemcpyDeviceToHost),40016);
-
-}
-
-// coupling CMB_ICB_fluid
-extern "C"
-void FC_FUNC_(transfer_coupling_fields_cmb_icb_fluid_from_device,
- TRANSFER_COUPLING_FIELDS_CMB_ICB_FLUID_FROM_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* displ_cm, realw* displ_ic, realw* accel_cm, realw* accel_ic, realw* accel_oc, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_fields_cmb_icb_fluid_from_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(accel_oc,mp->d_accel_outer_core,sizeof(realw)*(*size_oc),cudaMemcpyDeviceToHost),40014);
-
- print_CUDA_error_if_any(cudaMemcpy(accel_cm,mp->d_accel_crust_mantle,sizeof(realw)*(*size_cm),cudaMemcpyDeviceToHost),40015);
-
- print_CUDA_error_if_any(cudaMemcpy(accel_ic,mp->d_accel_inner_core,sizeof(realw)*(*size_ic),cudaMemcpyDeviceToHost),40016);
-
- print_CUDA_error_if_any(cudaMemcpy(displ_cm,mp->d_displ_crust_mantle,sizeof(realw)*(*size_cm),cudaMemcpyDeviceToHost),40017);
-
- print_CUDA_error_if_any(cudaMemcpy(displ_ic,mp->d_displ_inner_core,sizeof(realw)*(*size_ic),cudaMemcpyDeviceToHost),40018);
-
-}
-
-// coupling CMB_ocean
-extern "C"
-void FC_FUNC_(transfer_coupling_fields_cmb_ocean_from_device,
- TRANSFER_COUPLING_FIELDS_CMB_OCEAN_FROM_DEVICE)(int* size, realw* accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_fields_cmb_ocean_from_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(accel,mp->d_accel_crust_mantle,sizeof(realw)*(*size),cudaMemcpyDeviceToHost),40014);
-
-}
-
/* ----------------------------------------------------------------------------------------------- */
// backward/reconstructed fields
@@ -407,55 +255,6 @@
}
-// coupling fluid_CMB_ICB
-extern "C"
-void FC_FUNC_(transfer_coupling_b_fields_fluid_cmb_icb_from_device,
- TRANSFER_COUPLING_B_FIELDS_FLUID_CMB_icb_FROM_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* b_displ_cm, realw* b_displ_ic, realw* b_accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_b_fields_fluid_cmb_icb_from_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(b_accel,mp->d_b_accel_outer_core,sizeof(realw)*(*size_oc),cudaMemcpyDeviceToHost),40014);
-
- print_CUDA_error_if_any(cudaMemcpy(b_displ_cm,mp->d_b_displ_crust_mantle,sizeof(realw)*(*size_cm),cudaMemcpyDeviceToHost),40015);
-
- print_CUDA_error_if_any(cudaMemcpy(b_displ_ic,mp->d_b_displ_inner_core,sizeof(realw)*(*size_ic),cudaMemcpyDeviceToHost),40016);
-}
-
-// coupling CMB_fluid
-extern "C"
-void FC_FUNC_(transfer_coupling_b_fields_cmb_icb_fluid_from_device,
- TRANSFER_COUPLING_B_FIELDS_CMB_ICB_FLUID_FROM_DEVICE)(int* size_oc, int* size_cm, int* size_ic, realw* b_displ_cm, realw* b_displ_ic, realw* b_accel_cm, realw* b_accel_ic, realw* b_accel_oc, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_b_fields_cmb_icb_fluid_from_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(b_accel_oc,mp->d_b_accel_outer_core,sizeof(realw)*(*size_oc),cudaMemcpyDeviceToHost),40014);
-
- print_CUDA_error_if_any(cudaMemcpy(b_accel_cm,mp->d_b_accel_crust_mantle,sizeof(realw)*(*size_cm),cudaMemcpyDeviceToHost),40015);
-
- print_CUDA_error_if_any(cudaMemcpy(b_accel_ic,mp->d_b_accel_inner_core,sizeof(realw)*(*size_ic),cudaMemcpyDeviceToHost),40016);
-
- print_CUDA_error_if_any(cudaMemcpy(b_displ_cm,mp->d_b_displ_crust_mantle,sizeof(realw)*(*size_cm),cudaMemcpyDeviceToHost),40017);
-
- print_CUDA_error_if_any(cudaMemcpy(b_displ_ic,mp->d_b_displ_inner_core,sizeof(realw)*(*size_ic),cudaMemcpyDeviceToHost),40018);
-}
-
-// coupling CMB_ocean
-extern "C"
-void FC_FUNC_(transfer_coupling_b_fields_cmb_ocean_from_device,
- TRANSFER_COUPLING_B_FIELDS_CMB_OCEAN_FROM_DEVICE)(int* size, realw* b_accel, long* Mesh_pointer_f) {
-
- TRACE("transfer_coupling_b_fields_cmb_ocean_from_device");
-
- Mesh* mp = (Mesh*)(*Mesh_pointer_f); //get mesh pointer out of fortran integer container
-
- print_CUDA_error_if_any(cudaMemcpy(b_accel,mp->d_b_accel_crust_mantle,sizeof(realw)*(*size),cudaMemcpyDeviceToHost),40014);
-
-}
-
/* ----------------------------------------------------------------------------------------------- */
extern "C"
@@ -580,7 +379,7 @@
extern "C"
void FC_FUNC_(transfer_b_accel_cm_from_device,
- TRNASFER_B_ACCEL_CM_FROM_DEVICE)(int* size, realw* b_accel,long* Mesh_pointer_f) {
+ TRANSFER_B_ACCEL_CM_FROM_DEVICE)(int* size, realw* b_accel,long* Mesh_pointer_f) {
TRACE("transfer_b_accel_cm_from_device");
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_acoustic.F90
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_acoustic.F90 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_acoustic.F90 2012-02-28 00:02:26 UTC (rev 19689)
@@ -134,9 +134,6 @@
!---
!--- couple with mantle at the top of the outer core
!---
- call load_CPU_acoustic()
- call load_CPU_elastic()
-
if(ACTUALLY_COUPLE_FLUID_CMB) &
call compute_coupling_fluid_CMB(displ_crust_mantle,b_displ_crust_mantle, &
ibool_crust_mantle,ibelm_bottom_crust_mantle, &
@@ -156,12 +153,8 @@
wgllwgll_xy,ibool_outer_core,ibelm_bottom_outer_core, &
SIMULATION_TYPE,NSPEC2D_BOTTOM(IREGION_OUTER_CORE))
- call load_GPU_acoustic()
-
else
! on GPU
- call load_GPU_acoustic_coupling_fluid()
-
!---
!--- couple with mantle at the top of the outer core
!---
@@ -173,7 +166,6 @@
if( ACTUALLY_COUPLE_FLUID_ICB ) &
call compute_coupling_fluid_icb_cuda(Mesh_pointer)
- call load_CPU_acoustic_coupling_fluid()
endif
endif ! iphase == 1
@@ -344,87 +336,3 @@
end subroutine compute_forces_ac_update_veloc
-!=====================================================================
-
- subroutine load_GPU_acoustic
-
- use specfem_par
- use specfem_par_outercore
- implicit none
-
- ! daniel: TODO - temporary transfers to the GPU
- call transfer_fields_oc_to_device(NGLOB_OUTER_CORE,displ_outer_core, &
- veloc_outer_core,accel_outer_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_b_fields_oc_to_device(NGLOB_OUTER_CORE,b_displ_outer_core, &
- b_veloc_outer_core,b_accel_outer_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_CPU_acoustic
-
- use specfem_par
- use specfem_par_outercore
- implicit none
-
- ! daniel: TODO - temporary transfers to the CPU
- call transfer_fields_oc_from_device(NGLOB_OUTER_CORE,displ_outer_core, &
- veloc_outer_core,accel_outer_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_b_fields_oc_from_device(NGLOB_OUTER_CORE,b_displ_outer_core, &
- b_veloc_outer_core,b_accel_outer_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_GPU_acoustic_coupling_fluid
-
- use specfem_par
- use specfem_par_outercore,only: accel_outer_core,b_accel_outer_core
- use specfem_par_innercore,only: displ_inner_core,b_displ_inner_core
- use specfem_par_crustmantle,only: displ_crust_mantle,b_displ_crust_mantle
- implicit none
-
- ! daniel: TODO - temporary transfers to the GPU
- call transfer_coupling_fields_fluid_cmb_icb_to_device( &
- NGLOB_OUTER_CORE,NDIM*NGLOB_CRUST_MANTLE,NDIM*NGLOB_INNER_CORE, &
- displ_crust_mantle,displ_inner_core,accel_outer_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_coupling_b_fields_fluid_cmb_icb_to_device( &
- NGLOB_OUTER_CORE_ADJOINT,NDIM*NGLOB_CRUST_MANTLE_ADJOINT,NDIM*NGLOB_INNER_CORE_ADJOINT, &
- b_displ_crust_mantle,b_displ_inner_core,b_accel_outer_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_CPU_acoustic_coupling_fluid
-
- use specfem_par
- use specfem_par_outercore,only: accel_outer_core,b_accel_outer_core
- use specfem_par_innercore,only: displ_inner_core,b_displ_inner_core
- use specfem_par_crustmantle,only: displ_crust_mantle,b_displ_crust_mantle
- implicit none
-
- ! daniel: TODO - temporary transfers to the CPU
- call transfer_coupling_fields_fluid_cmb_icb_from_device( &
- NGLOB_OUTER_CORE,NDIM*NGLOB_CRUST_MANTLE,NDIM*NGLOB_INNER_CORE, &
- displ_crust_mantle,displ_inner_core,accel_outer_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_coupling_b_fields_fluid_cmb_icb_from_device( &
- NGLOB_OUTER_CORE_ADJOINT,NDIM*NGLOB_CRUST_MANTLE_ADJOINT,NDIM*NGLOB_INNER_CORE_ADJOINT, &
- b_displ_crust_mantle,b_displ_inner_core,b_accel_outer_core,Mesh_pointer)
- endif
-
- end subroutine
-
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_elastic.F90
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_elastic.F90 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/compute_forces_elastic.F90 2012-02-28 00:02:26 UTC (rev 19689)
@@ -246,8 +246,6 @@
! only for elements in first matching layer in the solid
if( .not. GPU_MODE ) then
! on CPU
- call load_CPU_elastic()
- call load_CPU_acoustic()
!---
!--- couple with outer core at the bottom of the mantle
!---
@@ -274,12 +272,8 @@
RHO_BOTTOM_OC,minus_g_icb, &
SIMULATION_TYPE,NSPEC2D_TOP(IREGION_INNER_CORE))
- call load_GPU_elastic()
-
else
! on GPU
- call load_GPU_elastic_coupling_fluid()
-
!---
!--- couple with outer core at the bottom of the mantle
!---
@@ -293,7 +287,6 @@
call compute_coupling_icb_fluid_cuda(Mesh_pointer, &
RHO_BOTTOM_OC,minus_g_icb,GRAVITY_VAL)
- call load_CPU_elastic_coupling_fluid()
endif
endif ! iphase == 1
@@ -488,22 +481,18 @@
if(OCEANS_VAL) then
if(.NOT. GPU_MODE) then
! on CPU
- call load_CPU_elastic()
call compute_coupling_ocean(accel_crust_mantle,b_accel_crust_mantle, &
rmass_crust_mantle,rmass_ocean_load,normal_top_crust_mantle, &
ibool_crust_mantle,ibelm_top_crust_mantle, &
updated_dof_ocean_load, &
SIMULATION_TYPE,NSPEC2D_TOP(IREGION_CRUST_MANTLE))
- call load_GPU_elastic()
else
! on GPU
- call load_GPU_elastic_coupling_ocean()
call compute_coupling_ocean_cuda(Mesh_pointer)
- call load_CPU_elastic_coupling_ocean()
endif
endif
@@ -714,143 +703,3 @@
#endif
end subroutine compute_forces_el_update_veloc
-
-!=====================================================================
-
- subroutine load_GPU_elastic
-
- use specfem_par
- use specfem_par_crustmantle
- use specfem_par_innercore
- implicit none
-
- ! daniel: TODO - temporary transfers to the GPU
- call transfer_fields_cm_to_device(NDIM*NGLOB_CRUST_MANTLE,displ_crust_mantle, &
- veloc_crust_mantle,accel_crust_mantle,Mesh_pointer)
- call transfer_fields_ic_to_device(NDIM*NGLOB_INNER_CORE,displ_inner_core, &
- veloc_inner_core,accel_inner_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_b_fields_cm_to_device(NDIM*NGLOB_CRUST_MANTLE,b_displ_crust_mantle, &
- b_veloc_crust_mantle,b_accel_crust_mantle,Mesh_pointer)
- call transfer_b_fields_ic_to_device(NDIM*NGLOB_INNER_CORE,b_displ_inner_core, &
- b_veloc_inner_core,b_accel_inner_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_CPU_elastic
-
- use specfem_par
- use specfem_par_crustmantle
- use specfem_par_innercore
- implicit none
-
- ! daniel: TODO - temporary transfers back to the CPU
- call transfer_fields_cm_from_device(NDIM*NGLOB_CRUST_MANTLE,displ_crust_mantle, &
- veloc_crust_mantle,accel_crust_mantle,Mesh_pointer)
- call transfer_fields_ic_from_device(NDIM*NGLOB_INNER_CORE,displ_inner_core, &
- veloc_inner_core,accel_inner_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_b_fields_cm_from_device(NDIM*NGLOB_CRUST_MANTLE,b_displ_crust_mantle, &
- b_veloc_crust_mantle,b_accel_crust_mantle,Mesh_pointer)
- call transfer_b_fields_ic_from_device(NDIM*NGLOB_INNER_CORE,b_displ_inner_core, &
- b_veloc_inner_core,b_accel_inner_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_GPU_elastic_coupling_fluid
-
- use specfem_par
- use specfem_par_outercore,only: accel_outer_core,b_accel_outer_core
- use specfem_par_innercore,only: displ_inner_core,b_displ_inner_core, &
- accel_inner_core,b_accel_inner_core
- use specfem_par_crustmantle,only: displ_crust_mantle,b_displ_crust_mantle, &
- accel_crust_mantle,b_accel_crust_mantle
- implicit none
-
- ! daniel: TODO - temporary transfers to the GPU
- call transfer_coupling_fields_cmb_icb_fluid_to_device( &
- NGLOB_OUTER_CORE,NDIM*NGLOB_CRUST_MANTLE,NDIM*NGLOB_INNER_CORE, &
- displ_crust_mantle,displ_inner_core,accel_crust_mantle,accel_inner_core, &
- accel_outer_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_coupling_b_fields_cmb_icb_fluid_to_device( &
- NGLOB_OUTER_CORE_ADJOINT,NDIM*NGLOB_CRUST_MANTLE_ADJOINT,NDIM*NGLOB_INNER_CORE_ADJOINT, &
- b_displ_crust_mantle,b_displ_inner_core,b_accel_crust_mantle,b_accel_inner_core, &
- b_accel_outer_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_CPU_elastic_coupling_fluid
-
- use specfem_par
- use specfem_par_outercore,only: accel_outer_core,b_accel_outer_core
- use specfem_par_innercore,only: displ_inner_core,b_displ_inner_core, &
- accel_inner_core,b_accel_inner_core
- use specfem_par_crustmantle,only: displ_crust_mantle,b_displ_crust_mantle, &
- accel_crust_mantle,b_accel_crust_mantle
- implicit none
-
- ! daniel: TODO - temporary transfers to the CPU
- call transfer_coupling_fields_cmb_icb_fluid_from_device( &
- NGLOB_OUTER_CORE,NDIM*NGLOB_CRUST_MANTLE,NDIM*NGLOB_INNER_CORE, &
- displ_crust_mantle,displ_inner_core,accel_crust_mantle,accel_inner_core, &
- accel_outer_core,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_coupling_b_fields_cmb_icb_fluid_from_device( &
- NGLOB_OUTER_CORE_ADJOINT,NDIM*NGLOB_CRUST_MANTLE_ADJOINT,NDIM*NGLOB_INNER_CORE_ADJOINT, &
- b_displ_crust_mantle,b_displ_inner_core,b_accel_crust_mantle,b_accel_inner_core, &
- b_accel_outer_core,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_GPU_elastic_coupling_ocean
-
- use specfem_par
- use specfem_par_crustmantle,only: accel_crust_mantle,b_accel_crust_mantle
- implicit none
-
- ! daniel: TODO - temporary transfers to the GPU
- call transfer_coupling_fields_cmb_ocean_to_device( &
- NDIM*NGLOB_CRUST_MANTLE,accel_crust_mantle,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_coupling_b_fields_cmb_ocean_to_device( &
- NDIM*NGLOB_CRUST_MANTLE_ADJOINT,b_accel_crust_mantle,Mesh_pointer)
- endif
-
- end subroutine
-
-!=====================================================================
-
- subroutine load_CPU_elastic_coupling_ocean
-
- use specfem_par
- use specfem_par_crustmantle,only: accel_crust_mantle,b_accel_crust_mantle
- implicit none
-
- ! daniel: TODO - temporary transfers to the GPU
- call transfer_coupling_fields_cmb_ocean_from_device( &
- NDIM*NGLOB_CRUST_MANTLE,accel_crust_mantle,Mesh_pointer)
-
- if( SIMULATION_TYPE == 3 ) then
- call transfer_coupling_b_fields_cmb_ocean_from_device( &
- NDIM*NGLOB_CRUST_MANTLE_ADJOINT,b_accel_crust_mantle,Mesh_pointer)
- endif
-
- end subroutine
Modified: seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/prepare_timerun.f90
===================================================================
--- seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/prepare_timerun.f90 2012-02-27 20:01:44 UTC (rev 19688)
+++ seismo/3D/SPECFEM3D_GLOBE/branches/SPECFEM3D_GLOBE_SUNFLOWER/src/specfem3D/prepare_timerun.f90 2012-02-28 00:02:26 UTC (rev 19689)
@@ -1075,7 +1075,7 @@
NSPEC_INNER_CORE,NGLOB_INNER_CORE, &
SIMULATION_TYPE,NOISE_TOMOGRAPHY, &
SAVE_FORWARD,ABSORBING_CONDITIONS, &
- GRAVITY_VAL,ROTATION_VAL, &
+ OCEANS_VAL,GRAVITY_VAL,ROTATION_VAL, &
ATTENUATION_VAL,USE_ATTENUATION_MIMIC, &
COMPUTE_AND_STORE_STRAIN, &
ANISOTROPIC_3D_MANTLE_VAL,ANISOTROPIC_INNER_CORE_VAL, &
@@ -1234,6 +1234,14 @@
endif
+ ! prepares oceans arrays
+ if ( OCEANS_VAL ) then
+ if(myrank == 0 ) write(IMAIN,*) " loading oceans arrays"
+
+ call prepare_oceans_device(Mesh_pointer,rmass_ocean_load)
+
+ endif
+
! crust/mantle region
if(myrank == 0 ) write(IMAIN,*) " loading crust/mantle region"
call prepare_crust_mantle_device(Mesh_pointer, &
@@ -1307,7 +1315,6 @@
NSPEC2D_TOP(IREGION_INNER_CORE))
call sync_all()
-
! transfer forward and backward fields to device with initial values
if(myrank == 0 ) write(IMAIN,*) " transfering initial wavefield"
call transfer_fields_cm_to_device(NDIM*NGLOB_CRUST_MANTLE,displ_crust_mantle,veloc_crust_mantle,accel_crust_mantle, &
More information about the CIG-COMMITS
mailing list