[cig-commits] [commit] devel,master: stupid commit (9b7ff79)
cig_noreply at geodynamics.org
cig_noreply at geodynamics.org
Thu Nov 6 08:25:54 PST 2014
Repository : https://github.com/geodynamics/specfem3d_globe
On branches: devel,master
Link : https://github.com/geodynamics/specfem3d_globe/compare/bc58e579b3b0838a0968725a076f5904845437ca...be63f20cbb6f462104e949894dbe205d2398cd7f
>---------------------------------------------------------------
commit 9b7ff79e56fa94af04b6422d5644530409c78376
Author: Brice Videau <brice.videau at imag.fr>
Date: Mon Jul 7 14:57:26 2014 +0200
stupid commit
>---------------------------------------------------------------
9b7ff79e56fa94af04b6422d5644530409c78376
src/gpu/kernels.gen/kernel_cuda.mk | 28 ++++++++++++++++++++++++++++
src/gpu/kernels.gen/kernel_inc_cl.c | 28 ++++++++++++++++++++++++++++
src/gpu/kernels.gen/kernel_list.h | 28 ++++++++++++++++++++++++++++
src/gpu/kernels.gen/kernel_proto.cu.h | 28 ++++++++++++++++++++++++++++
4 files changed, 112 insertions(+)
diff --git a/src/gpu/kernels.gen/kernel_cuda.mk b/src/gpu/kernels.gen/kernel_cuda.mk
index 1822fa9..39cc2de 100644
--- a/src/gpu/kernels.gen/kernel_cuda.mk
+++ b/src/gpu/kernels.gen/kernel_cuda.mk
@@ -1,4 +1,32 @@
cuda_kernels_OBJS := \
+ $O/assemble_boundary_accel_on_device.cuda-kernel.o \
+ $O/assemble_boundary_potential_on_device.cuda-kernel.o \
+ $O/prepare_boundary_potential_on_device.cuda-kernel.o \
+ $O/prepare_boundary_accel_on_device.cuda-kernel.o \
+ $O/get_maximum_scalar_kernel.cuda-kernel.o \
+ $O/get_maximum_vector_kernel.cuda-kernel.o \
+ $O/compute_add_sources_adjoint_kernel.cuda-kernel.o \
+ $O/compute_add_sources_kernel.cuda-kernel.o \
+ $O/compute_coupling_fluid_CMB_kernel.cuda-kernel.o \
+ $O/compute_coupling_fluid_ICB_kernel.cuda-kernel.o \
+ $O/compute_coupling_CMB_fluid_kernel.cuda-kernel.o \
+ $O/compute_coupling_ICB_fluid_kernel.cuda-kernel.o \
+ $O/compute_coupling_ocean_kernel.cuda-kernel.o \
+ $O/write_seismograms_transfer_from_device_kernel.cuda-kernel.o \
+ $O/write_seismograms_transfer_strain_from_device_kernel.cuda-kernel.o \
+ $O/noise_transfer_surface_to_host_kernel.cuda-kernel.o \
+ $O/noise_add_source_master_rec_kernel.cuda-kernel.o \
+ $O/noise_add_surface_movie_kernel.cuda-kernel.o \
+ $O/compute_stacey_acoustic_kernel.cuda-kernel.o \
+ $O/compute_stacey_acoustic_backward_kernel.cuda-kernel.o \
+ $O/compute_stacey_elastic_kernel.cuda-kernel.o \
+ $O/compute_stacey_elastic_backward_kernel.cuda-kernel.o \
+ $O/update_disp_veloc_kernel.cuda-kernel.o \
+ $O/update_potential_kernel.cuda-kernel.o \
+ $O/update_accel_elastic_kernel.cuda-kernel.o \
+ $O/update_veloc_elastic_kernel.cuda-kernel.o \
+ $O/update_accel_acoustic_kernel.cuda-kernel.o \
+ $O/update_veloc_acoustic_kernel.cuda-kernel.o \
$O/outer_core_impl_kernel_forward.cuda-kernel.o \
$O/outer_core_impl_kernel_adjoint.cuda-kernel.o \
$O/inner_core_impl_kernel_forward.cuda-kernel.o \
diff --git a/src/gpu/kernels.gen/kernel_inc_cl.c b/src/gpu/kernels.gen/kernel_inc_cl.c
index 73a80ca..683220e 100644
--- a/src/gpu/kernels.gen/kernel_inc_cl.c
+++ b/src/gpu/kernels.gen/kernel_inc_cl.c
@@ -1,3 +1,31 @@
+#include "assemble_boundary_accel_on_device_cl.c"
+#include "assemble_boundary_potential_on_device_cl.c"
+#include "prepare_boundary_potential_on_device_cl.c"
+#include "prepare_boundary_accel_on_device_cl.c"
+#include "get_maximum_scalar_kernel_cl.c"
+#include "get_maximum_vector_kernel_cl.c"
+#include "compute_add_sources_adjoint_kernel_cl.c"
+#include "compute_add_sources_kernel_cl.c"
+#include "compute_coupling_fluid_CMB_kernel_cl.c"
+#include "compute_coupling_fluid_ICB_kernel_cl.c"
+#include "compute_coupling_CMB_fluid_kernel_cl.c"
+#include "compute_coupling_ICB_fluid_kernel_cl.c"
+#include "compute_coupling_ocean_kernel_cl.c"
+#include "write_seismograms_transfer_from_device_kernel_cl.c"
+#include "write_seismograms_transfer_strain_from_device_kernel_cl.c"
+#include "noise_transfer_surface_to_host_kernel_cl.c"
+#include "noise_add_source_master_rec_kernel_cl.c"
+#include "noise_add_surface_movie_kernel_cl.c"
+#include "compute_stacey_acoustic_kernel_cl.c"
+#include "compute_stacey_acoustic_backward_kernel_cl.c"
+#include "compute_stacey_elastic_kernel_cl.c"
+#include "compute_stacey_elastic_backward_kernel_cl.c"
+#include "update_disp_veloc_kernel_cl.c"
+#include "update_potential_kernel_cl.c"
+#include "update_accel_elastic_kernel_cl.c"
+#include "update_veloc_elastic_kernel_cl.c"
+#include "update_accel_acoustic_kernel_cl.c"
+#include "update_veloc_acoustic_kernel_cl.c"
#include "outer_core_impl_kernel_forward_cl.c"
#include "outer_core_impl_kernel_adjoint_cl.c"
#include "inner_core_impl_kernel_forward_cl.c"
diff --git a/src/gpu/kernels.gen/kernel_list.h b/src/gpu/kernels.gen/kernel_list.h
index 1e1f1fa..88b1662 100644
--- a/src/gpu/kernels.gen/kernel_list.h
+++ b/src/gpu/kernels.gen/kernel_list.h
@@ -1,3 +1,31 @@
+BOAST_KERNEL(assemble_boundary_accel_on_device);
+BOAST_KERNEL(assemble_boundary_potential_on_device);
+BOAST_KERNEL(prepare_boundary_potential_on_device);
+BOAST_KERNEL(prepare_boundary_accel_on_device);
+BOAST_KERNEL(get_maximum_scalar_kernel);
+BOAST_KERNEL(get_maximum_vector_kernel);
+BOAST_KERNEL(compute_add_sources_adjoint_kernel);
+BOAST_KERNEL(compute_add_sources_kernel);
+BOAST_KERNEL(compute_coupling_fluid_CMB_kernel);
+BOAST_KERNEL(compute_coupling_fluid_ICB_kernel);
+BOAST_KERNEL(compute_coupling_CMB_fluid_kernel);
+BOAST_KERNEL(compute_coupling_ICB_fluid_kernel);
+BOAST_KERNEL(compute_coupling_ocean_kernel);
+BOAST_KERNEL(write_seismograms_transfer_from_device_kernel);
+BOAST_KERNEL(write_seismograms_transfer_strain_from_device_kernel);
+BOAST_KERNEL(noise_transfer_surface_to_host_kernel);
+BOAST_KERNEL(noise_add_source_master_rec_kernel);
+BOAST_KERNEL(noise_add_surface_movie_kernel);
+BOAST_KERNEL(compute_stacey_acoustic_kernel);
+BOAST_KERNEL(compute_stacey_acoustic_backward_kernel);
+BOAST_KERNEL(compute_stacey_elastic_kernel);
+BOAST_KERNEL(compute_stacey_elastic_backward_kernel);
+BOAST_KERNEL(update_disp_veloc_kernel);
+BOAST_KERNEL(update_potential_kernel);
+BOAST_KERNEL(update_accel_elastic_kernel);
+BOAST_KERNEL(update_veloc_elastic_kernel);
+BOAST_KERNEL(update_accel_acoustic_kernel);
+BOAST_KERNEL(update_veloc_acoustic_kernel);
BOAST_KERNEL(outer_core_impl_kernel_forward);
BOAST_KERNEL(outer_core_impl_kernel_adjoint);
BOAST_KERNEL(inner_core_impl_kernel_forward);
diff --git a/src/gpu/kernels.gen/kernel_proto.cu.h b/src/gpu/kernels.gen/kernel_proto.cu.h
index 40b704f..dd198e9 100644
--- a/src/gpu/kernels.gen/kernel_proto.cu.h
+++ b/src/gpu/kernels.gen/kernel_proto.cu.h
@@ -1,3 +1,31 @@
+__global__ void assemble_boundary_accel_on_device(float * d_accel, const float * d_send_accel_buffer, const int num_interfaces, const int max_nibool_interfaces, const int * d_nibool_interfaces, const int * d_ibool_interfaces);
+__global__ void assemble_boundary_potential_on_device(float * d_potential_dot_dot_acoustic, const float * d_send_potential_dot_dot_buffer, const int num_interfaces, const int max_nibool_interfaces, const int * d_nibool_interfaces, const int * d_ibool_interfaces);
+__global__ void prepare_boundary_potential_on_device(const float * d_potential_dot_dot_acoustic, float * d_send_potential_dot_dot_buffer, const int num_interfaces, const int max_nibool_interfaces, const int * d_nibool_interfaces, const int * d_ibool_interfaces);
+__global__ void prepare_boundary_accel_on_device(const float * d_accel, float * d_send_accel_buffer, const int num_interfaces, const int max_nibool_interfaces, const int * d_nibool_interfaces, const int * d_ibool_interfaces);
+__global__ void get_maximum_scalar_kernel(const float * array, const int size, float * d_max);
+__global__ void get_maximum_vector_kernel(const float * array, const int size, float * d_max);
+__global__ void compute_add_sources_adjoint_kernel(float * accel, const int nrec, const float * adj_sourcearrays, const int * ibool, const int * ispec_selected_rec, const int * pre_computed_irec, const int nadj_rec_local);
+__global__ void compute_add_sources_kernel(float * accel, const int * ibool, const float * sourcearrays, const double * stf_pre_compute, const int myrank, const int * islice_selected_source, const int * ispec_selected_source, const int NSOURCES);
+__global__ void compute_coupling_fluid_CMB_kernel(const float * displ_crust_mantle, float * accel_outer_core, const int * ibool_crust_mantle, const int * ibelm_bottom_crust_mantle, const float * normal_top_outer_core, const float * jacobian2D_top_outer_core, const float * wgllwgll_xy, const int * ibool_outer_core, const int * ibelm_top_outer_core, const int NSPEC2D_TOP_OC);
+__global__ void compute_coupling_fluid_ICB_kernel(const float * displ_inner_core, float * accel_outer_core, const int * ibool_inner_core, const int * ibelm_top_inner_core, const float * normal_bottom_outer_core, const float * jacobian2D_bottom_outer_core, const float * wgllwgll_xy, const int * ibool_outer_core, const int * ibelm_bottom_outer_core, const int NSPEC2D_BOTTOM_OC);
+__global__ void compute_coupling_CMB_fluid_kernel(const float * displ_crust_mantle, float * accel_crust_mantle, const float * accel_outer_core, const int * ibool_crust_mantle, const int * ibelm_bottom_crust_mantle, const float * normal_top_outer_core, const float * jacobian2D_top_outer_core, const float * wgllwgll_xy, const int * ibool_outer_core, const int * ibelm_top_outer_core, const float RHO_TOP_OC, const float minus_g_cmb, const int GRAVITY, const int NSPEC2D_BOTTOM_CM);
+__global__ void compute_coupling_ICB_fluid_kernel(const float * displ_inner_core, float * accel_inner_core, const float * accel_outer_core, const int * ibool_inner_core, const int * ibelm_top_inner_core, const float * normal_bottom_outer_core, const float * jacobian2D_bottom_outer_core, const float * wgllwgll_xy, const int * ibool_outer_core, const int * ibelm_bottom_outer_core, const float RHO_BOTTOM_OC, const float minus_g_icb, const int GRAVITY, const int NSPEC2D_TOP_IC);
+__global__ void compute_coupling_ocean_kernel(float * accel_crust_mantle, const float * rmassx_crust_mantle, const float * rmassy_crust_mantle, const float * rmassz_crust_mantle, const float * rmass_ocean_load, const int npoin_ocean_load, const int * ibool_ocean_load, const float * normal_ocean_load);
+__global__ void write_seismograms_transfer_from_device_kernel(const int * number_receiver_global, const int * ispec_selected_rec, const int * ibool, float * station_seismo_field, const float * d_field, const int nrec_local);
+__global__ void write_seismograms_transfer_strain_from_device_kernel(const int * number_receiver_global, const int * ispec_selected_rec, const int * ibool, float * station_strain_field, const float * d_field, const int nrec_local);
+__global__ void noise_transfer_surface_to_host_kernel(const int * ibelm_top, const int nspec_top, const int * ibool, const float * displ, float * noise_surface_movie);
+__global__ void noise_add_source_master_rec_kernel(const int * ibool, const int * ispec_selected_rec, const int irec_master_noise, float * accel, const float * noise_sourcearray, const int it);
+__global__ void noise_add_surface_movie_kernel(float * accel, const int * ibool, const int * ibelm_top, const int nspec_top, const float * noise_surface_movie, const float * normal_x_noise, const float * normal_y_noise, const float * normal_z_noise, const float * mask_noise, const float * jacobian2D, const float * wgllwgll);
+__global__ void compute_stacey_acoustic_kernel(const float * potential_dot_acoustic, float * potential_dot_dot_acoustic, const int interface_type, const int num_abs_boundary_faces, const int * abs_boundary_ispec, const int * nkmin_xi, const int * nkmin_eta, const int * njmin, const int * njmax, const int * nimin, const int * nimax, const float * abs_boundary_jacobian2D, const float * wgllwgll, const int * ibool, const float * vpstore, const int SAVE_FORWARD, float * b_absorb_potential);
+__global__ void compute_stacey_acoustic_backward_kernel(float * b_potential_dot_dot_acoustic, const float * b_absorb_potential, const int interface_type, const int num_abs_boundary_faces, const int * abs_boundary_ispec, const int * nkmin_xi, const int * nkmin_eta, const int * njmin, const int * njmax, const int * nimin, const int * nimax, const int * ibool);
+__global__ void compute_stacey_elastic_kernel(const float * veloc, float * accel, const int interface_type, const int num_abs_boundary_faces, const int * abs_boundary_ispec, const int * nkmin_xi, const int * nkmin_eta, const int * njmin, const int * njmax, const int * nimin, const int * nimax, const float * abs_boundary_normal, const float * abs_boundary_jacobian2D, const float * wgllwgll, const int * ibool, const float * rho_vp, const float * rho_vs, const int SAVE_FORWARD, float * b_absorb_field);
+__global__ void compute_stacey_elastic_backward_kernel(float * b_accel, const float * b_absorb_field, const int interface_type, const int num_abs_boundary_faces, const int * abs_boundary_ispec, const int * nkmin_xi, const int * nkmin_eta, const int * njmin, const int * njmax, const int * nimin, const int * nimax, const int * ibool);
+__global__ void update_disp_veloc_kernel(float * displ, float * veloc, float * accel, const int size, const float deltat, const float deltatsqover2, const float deltatover2);
+__global__ void update_potential_kernel(float * potential_acoustic, float * potential_dot_acoustic, float * potential_dot_dot_acoustic, const int size, const float deltat, const float deltatsqover2, const float deltatover2);
+__global__ void update_accel_elastic_kernel(float * accel, const float * veloc, const int size, const float two_omega_earth, const float * rmassx, const float * rmassy, const float * rmassz);
+__global__ void update_veloc_elastic_kernel(float * veloc, const float * accel, const int size, const float deltatover2);
+__global__ void update_accel_acoustic_kernel(float * accel, const int size, const float * rmass);
+__global__ void update_veloc_acoustic_kernel(float * veloc, const float * accel, const int size, const float deltatover2);
__global__ void outer_core_impl_kernel_forward(const int nb_blocks_to_compute, const int * d_ibool, const int * d_phase_ispec_inner, const int num_phase_ispec, const int d_iphase, const int use_mesh_coloring_gpu, const float * __restrict__ d_potential, float * d_potential_dot_dot, const float * __restrict__ d_xix, const float * __restrict__ d_xiy, const float * __restrict__ d_xiz, const float * __restrict__ d_etax, const float * __restrict__ d_etay, const float * __restrict__ d_etaz, const float * __restrict__ d_gammax, const float * __restrict__ d_gammay, const float * __restrict__ d_gammaz, const float * __restrict__ d_hprime_xx, const float * __restrict__ d_hprimewgll_xx, const float * __restrict__ wgllwgll_xy, const float * __restrict__ wgllwgll_xz, const float * __restrict__ wgllwgll_yz, const int GRAVITY, const float * __restrict__ d_xstore, const float * __restrict__ d_ystore, const float * __restrict__ d_zstore, const float * __restrict__ d_d_ln_density_dr_table, const float
* __restrict__ d_minus_rho_g_over_kappa_fluid, const float * __restrict__ wgll_cube, const int ROTATION, const float time, const float two_omega_earth, const float deltat, float * d_A_array_rotation, float * d_B_array_rotation, const int NSPEC_OUTER_CORE);
__global__ void outer_core_impl_kernel_adjoint(const int nb_blocks_to_compute, const int * d_ibool, const int * d_phase_ispec_inner, const int num_phase_ispec, const int d_iphase, const int use_mesh_coloring_gpu, const float * __restrict__ d_potential, float * d_potential_dot_dot, const float * __restrict__ d_xix, const float * __restrict__ d_xiy, const float * __restrict__ d_xiz, const float * __restrict__ d_etax, const float * __restrict__ d_etay, const float * __restrict__ d_etaz, const float * __restrict__ d_gammax, const float * __restrict__ d_gammay, const float * __restrict__ d_gammaz, const float * __restrict__ d_hprime_xx, const float * __restrict__ d_hprimewgll_xx, const float * __restrict__ wgllwgll_xy, const float * __restrict__ wgllwgll_xz, const float * __restrict__ wgllwgll_yz, const int GRAVITY, const float * __restrict__ d_xstore, const float * __restrict__ d_ystore, const float * __restrict__ d_zstore, const float * __restrict__ d_d_ln_density_dr_table, const float
* __restrict__ d_minus_rho_g_over_kappa_fluid, const float * __restrict__ wgll_cube, const int ROTATION, const float time, const float two_omega_earth, const float deltat, float * d_A_array_rotation, float * d_B_array_rotation, const int NSPEC_OUTER_CORE);
__global__
More information about the CIG-COMMITS
mailing list