[cig-commits] [commit] devel: working nopn regression testing (c1cff1e)

cig_noreply at geodynamics.org cig_noreply at geodynamics.org
Tue Jul 29 09:57:01 PDT 2014


Repository : https://github.com/geodynamics/specfem3d_globe

On branch  : devel
Link       : https://github.com/geodynamics/specfem3d_globe/compare/e21ea2ffe05127e06dc547229fa9617f7f56559c...9b7489516553b79bc155e080df413bcf6900932f

>---------------------------------------------------------------

commit c1cff1e9106897541af62ec1f2c85ac5b5cf8987
Author: Brice Videau <brice.videau at imag.fr>
Date:   Mon Jul 7 14:56:16 2014 +0200

    working nopn regression testing


>---------------------------------------------------------------

c1cff1e9106897541af62ec1f2c85ac5b5cf8987
 src/gpu/boast/compute_coupling_fluid_CMB_kernel.rb |  4 ++--
 .../compute_coupling_CMB_fluid_kernel.cu           |  2 +-
 .../compute_coupling_CMB_fluid_kernel_cl.c         |  2 +-
 .../compute_coupling_ICB_fluid_kernel.cu           |  2 +-
 .../compute_coupling_ICB_fluid_kernel_cl.c         |  2 +-
 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 ----------------------
 9 files changed, 6 insertions(+), 118 deletions(-)

diff --git a/src/gpu/boast/compute_coupling_fluid_CMB_kernel.rb b/src/gpu/boast/compute_coupling_fluid_CMB_kernel.rb
index 9a45a99..866dcee 100644
--- a/src/gpu/boast/compute_coupling_fluid_CMB_kernel.rb
+++ b/src/gpu/boast/compute_coupling_fluid_CMB_kernel.rb
@@ -41,7 +41,7 @@ module BOAST
       ibelm_1               = Int( "ibelm_top_outer_core",      :dir => :in,    :dim => [ Dim() ])
       rho_oc                = Real("RHO_TOP_OC",                :dir => :in)
       minus_g               = Real("minus_g_cmb",               :dir => :in)
-      gravity               = Int( "GRAVITY")
+      gravity               = Int( "GRAVITY",                   :dir => :in)
       nspec2D               = Int( "NSPEC2D_BOTTOM_CM",         :dir => :in)
     elsif type == :ICB_fluid then
       function_name = "compute_coupling_ICB_fluid_kernel"
@@ -56,7 +56,7 @@ module BOAST
       ibelm_1               = Int( "ibelm_bottom_outer_core",   :dir => :in,    :dim => [ Dim() ])
       rho_oc                = Real("RHO_BOTTOM_OC",             :dir => :in)
       minus_g               = Real("minus_g_icb",               :dir => :in)
-      gravity               = Int( "GRAVITY")
+      gravity               = Int( "GRAVITY",                   :dir => :in)
       nspec2D               = Int( "NSPEC2D_TOP_IC",            :dir => :in)
     else
      raise "Unsupported coupling_fluid_type!"
diff --git a/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel.cu b/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel.cu
index 89010db..cc6a8c0 100644
--- a/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel.cu
+++ b/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel.cu
@@ -49,7 +49,7 @@
 #ifndef BLOCKSIZE_TRANSFER
 #define BLOCKSIZE_TRANSFER 256
 #endif
-__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, int GRAVITY, const int NSPEC2D_BOTTOM_CM){
+__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){
   int i;
   int j;
   int k;
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
index 1b4b1a1..eb5d6bd 100644
--- a/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel_cl.c
+++ b/src/gpu/kernels.gen/compute_coupling_CMB_fluid_kernel_cl.c
@@ -60,7 +60,7 @@ inline void atomicAdd(volatile __global float *source, const float val) {\n\
 #ifndef BLOCKSIZE_TRANSFER\n\
 #define BLOCKSIZE_TRANSFER 256\n\
 #endif\n\
-__kernel void compute_coupling_CMB_fluid_kernel(const __global float * displ_crust_mantle, __global float * accel_crust_mantle, const __global float * accel_outer_core, const __global int * ibool_crust_mantle, const __global int * ibelm_bottom_crust_mantle, const __global float * normal_top_outer_core, const __global float * jacobian2D_top_outer_core, const __global float * wgllwgll_xy, const __global int * ibool_outer_core, const __global int * ibelm_top_outer_core, const float RHO_TOP_OC, const float minus_g_cmb, int GRAVITY, const int NSPEC2D_BOTTOM_CM){\n\
+__kernel void compute_coupling_CMB_fluid_kernel(const __global float * displ_crust_mantle, __global float * accel_crust_mantle, const __global float * accel_outer_core, const __global int * ibool_crust_mantle, const __global int * ibelm_bottom_crust_mantle, const __global float * normal_top_outer_core, const __global float * jacobian2D_top_outer_core, const __global float * wgllwgll_xy, const __global int * ibool_outer_core, const __global int * ibelm_top_outer_core, const float RHO_TOP_OC, const float minus_g_cmb, const int GRAVITY, const int NSPEC2D_BOTTOM_CM){\n\
   int i;\n\
   int j;\n\
   int k;\n\
diff --git a/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel.cu b/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel.cu
index dec29c7..b7d65a6 100644
--- a/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel.cu
+++ b/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel.cu
@@ -49,7 +49,7 @@
 #ifndef BLOCKSIZE_TRANSFER
 #define BLOCKSIZE_TRANSFER 256
 #endif
-__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, int GRAVITY, const int NSPEC2D_TOP_IC){
+__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){
   int i;
   int j;
   int k;
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
index ffec8e8..d5f8df8 100644
--- a/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel_cl.c
+++ b/src/gpu/kernels.gen/compute_coupling_ICB_fluid_kernel_cl.c
@@ -60,7 +60,7 @@ inline void atomicAdd(volatile __global float *source, const float val) {\n\
 #ifndef BLOCKSIZE_TRANSFER\n\
 #define BLOCKSIZE_TRANSFER 256\n\
 #endif\n\
-__kernel void compute_coupling_ICB_fluid_kernel(const __global float * displ_inner_core, __global float * accel_inner_core, const __global float * accel_outer_core, const __global int * ibool_inner_core, const __global int * ibelm_top_inner_core, const __global float * normal_bottom_outer_core, const __global float * jacobian2D_bottom_outer_core, const __global float * wgllwgll_xy, const __global int * ibool_outer_core, const __global int * ibelm_bottom_outer_core, const float RHO_BOTTOM_OC, const float minus_g_icb, int GRAVITY, const int NSPEC2D_TOP_IC){\n\
+__kernel void compute_coupling_ICB_fluid_kernel(const __global float * displ_inner_core, __global float * accel_inner_core, const __global float * accel_outer_core, const __global int * ibool_inner_core, const __global int * ibelm_top_inner_core, const __global float * normal_bottom_outer_core, const __global float * jacobian2D_bottom_outer_core, const __global float * wgllwgll_xy, const __global int * ibool_outer_core, const __global int * ibelm_bottom_outer_core, const float RHO_BOTTOM_OC, const float minus_g_icb, const int GRAVITY, const int NSPEC2D_TOP_IC){\n\
   int i;\n\
   int j;\n\
   int k;\n\
diff --git a/src/gpu/kernels.gen/kernel_cuda.mk b/src/gpu/kernels.gen/kernel_cuda.mk
index 39cc2de..1822fa9 100644
--- a/src/gpu/kernels.gen/kernel_cuda.mk
+++ b/src/gpu/kernels.gen/kernel_cuda.mk
@@ -1,32 +1,4 @@
 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 683220e..73a80ca 100644
--- a/src/gpu/kernels.gen/kernel_inc_cl.c
+++ b/src/gpu/kernels.gen/kernel_inc_cl.c
@@ -1,31 +1,3 @@
-#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 88b1662..1e1f1fa 100644
--- a/src/gpu/kernels.gen/kernel_list.h
+++ b/src/gpu/kernels.gen/kernel_list.h
@@ -1,31 +1,3 @@
-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 90e6434..40b704f 100644
--- a/src/gpu/kernels.gen/kernel_proto.cu.h
+++ b/src/gpu/kernels.gen/kernel_proto.cu.h
@@ -1,31 +1,3 @@
-__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, 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, 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