Skip to content

Commit

Permalink
hvpvp
Browse files Browse the repository at this point in the history
  • Loading branch information
@icui
icui committed Apr 22, 2020
1 parent 08d504b commit 5329763
Show file tree
Hide file tree
Showing 5 changed files with 303 additions and 31 deletions.
127 changes: 112 additions & 15 deletions src/gpu/boast/compute_hess_kernel.rb
View file
Original file line number Diff line number Diff line change
@@ -1,36 +1,133 @@

module BOAST

def BOAST::compute_hess_kernel(ref = true, n_gll3 = 125)
def BOAST::compute_hess_kernel(ref = true, n_gllx = 5, n_gll2 = 25, n_gll3 = 125, n_gll3_padded = 128)
push_env( :array_start => 0 )
kernel = CKernel::new
function_name = "compute_hess_kernel"
ibool = Int( "ibool", :dir => :in, :dim => [Dim()] )
accel = Real("accel", :dir => :in, :dim => [Dim(3), Dim()] )
b_accel = Real("b_accel", :dir => :in, :dim => [Dim(3), Dim()] )
hess_kl = Real("hess_kl", :dir => :inout,:dim => [Dim()] )
deltat = Real("deltat", :dir => :in)
nspec_ab= Int( "NSPEC_AB", :dir => :in)

v = []
v.push hess_kl = Real("hess_kl", :dir => :inout,:dim => [Dim()] )
v.push nspec = Int( "NSPEC", :dir => :in)
v.push deltat = Real("deltat", :dir => :in)
v.push d_ibool = Int( "d_ibool", :dir => :in, :dim => [Dim()] )
v.push d_b_displ = Real("d_b_displ", :dir => :in, :dim => [Dim(3), Dim()] )
v.push *d_xi = [d_xix = Real("d_xix", :dir => :in, :dim => [Dim()] ), d_xiy = Real("d_xiy", :dir => :in, :dim => [Dim()] ), d_xiz = Real("d_xiz", :dir => :in, :dim => [Dim()] ) ]
v.push *d_eta = [d_etax = Real("d_etax", :dir => :in, :dim => [Dim()] ), d_etay = Real("d_etay",:dir => :in, :dim => [Dim()] ), d_etaz = Real("d_etaz",:dir => :in, :dim => [Dim()] ) ]
v.push *d_gamma = [d_gammax = Real("d_gammax", :dir => :in, :dim => [Dim()] ), d_gammay = Real("d_gammay",:dir => :in, :dim => [Dim()] ), d_gammaz = Real("d_gammaz",:dir => :in, :dim => [Dim()] ) ]
v.push d_hprime_xx = Real("d_hprime_xx", :dir => :in, :dim => [Dim()] )

ngllx = Int("NGLLX", :const => n_gllx)
ngll2 = Int("NGLL2", :const => n_gll2)
ngll3 = Int("NGLL3", :const => n_gll3)
ngll3_padded = Int("NGLL3_PADDED", :const => n_gll3_padded)

p = Procedure(function_name, [ibool, accel, b_accel, hess_kl, deltat, nspec_ab])
p = Procedure(function_name, v)
if (get_lang == CUDA and ref) then
get_output.print File::read("references/#{function_name}.cu")
elsif(get_lang == CL or get_lang == CUDA) then
make_specfem3d_header( :ngll3 => n_gll3 )
make_specfem3d_header( :ngllx => n_gllx, :ngll2 => n_gll2, :ngll3 => n_gll3, :ngll3_padded => n_gll3_padded )
open p

decl ispec = Int("ispec")
decl ijk_ispec = Int("ijk_ispec")
decl tx = Int("tx")
decl iglob = Int("iglob")

decl *s_dummy_loc = ["x", "y", "z"].collect { |a|
Real("s_dummy#{a}_loc", :local => true, :dim => [Dim(ngll3)] )
}
decl sh_hprime_xx = Real("sh_hprime_xx", :local => true, :dim => [Dim(ngll2)] )


# compute div
decl k = Int("K")
decl j = Int("J")
decl i = Int("I")
decl l = Int("l")
decl offset = Int("offset")
tempanl = ["x", "y", "z"].collect { |a|
[ 1, 2, 3 ].collect { |n|
Real("temp#{a}#{n}l")
}
}
decl *(tempanl.flatten)
decl *xil = [ Real("xixl"), Real("xiyl"), Real("xizl") ]
decl *etal = [ Real("etaxl"), Real("etayl"), Real("etazl") ]
decl *gammal= [ Real("gammaxl"),Real("gammayl"),Real("gammazl")]
decl templ= Real("templ")
decl *fac = (1..3).collect { |n| Real("fac#{n}") }

decl hess_kappa = Real('hess_kappa')
# decl hess_mu = Real('hess_mu')
decl hess_tmp = Real('hess_tmp')

print ispec === get_group_id(0) + get_group_id(1)*get_num_groups(0)
print If(ispec < nspec_ab) {
print ijk_ispec === get_local_id(0) + ngll3*ispec
print iglob === ibool[ijk_ispec] - 1
print hess_kl[ijk_ispec] === hess_kl[ijk_ispec] + deltat * ( b_accel[0, iglob] * b_accel[0, iglob]\
+ b_accel[1, iglob] * b_accel[1, iglob]\
+ b_accel[2, iglob] * b_accel[2, iglob])
print ijk_ispec === get_local_id(0) + ngll3*ispec
print tx === get_local_id(0)

# minor threads load hprime
print If(tx < ngll2) {
print sh_hprime_xx[tx] === d_hprime_xx[tx]
}

# loads element displacements
# all threads load their displacement into shared memory
print If(ispec < nspec) {
print iglob === d_ibool[ijk_ispec] - 1
(0..2).each { |indx|
print s_dummy_loc[indx][tx] === d_b_displ[indx,iglob]
}
}

# synchronizes threads
print barrier(:local)

print If(ispec < nspec) {
print k === tx/ngll2
print j === (tx-k*ngll2)/ngllx
print i === tx - k*ngll2 - j*ngllx

tempanl.flatten.each { |t|
print t === 0.0
}
print For(l, 0, ngllx - 1) {
print fac[0] === sh_hprime_xx[l*ngllx + i]
(0..2).each { |indx|
print tempanl[indx][0] === tempanl[indx][0] + s_dummy_loc[indx][k*ngll2 + j*ngllx + l]*fac[0]
}
print fac[1] === sh_hprime_xx[l*ngllx + j]
(0..2).each { |indx|
print tempanl[indx][1] === tempanl[indx][1] + s_dummy_loc[indx][k*ngll2 + l*ngllx + i]*fac[1]
}
print fac[2] === sh_hprime_xx[l*ngllx + k]
(0..2).each { |indx|
print tempanl[indx][2] === tempanl[indx][2] + s_dummy_loc[indx][l*ngll2 + j*ngllx + i]*fac[2]
}
}
print offset === ispec*ngll3_padded + tx
(0..2).each { |indx|
print xil[indx] === d_xi[indx][offset]
print etal[indx] === d_eta[indx][offset]
print gammal[indx] === d_gamma[indx][offset]
}

print hess_kappa === 0.0
# print hess_mu === 0.0

(0..2).each { |indx1|
(0..2).each { |indx2|
if indx1 == indx2 then
print hess_kappa === hess_kappa + xil[indx2]*tempanl[indx1][0] + etal[indx2]*tempanl[indx1][1] + gammal[indx2]*tempanl[indx1][2]
# else
# print hess_mu === hess_mu + xil[indx2]*tempanl[indx1][0] + etal[indx2]*tempanl[indx1][1] + gammal[indx2]*tempanl[indx1][2]
end
}
}

print hess_kappa === hess_kappa * hess_kappa
print hess_tmp === hess_kappa
print hess_kl[ijk_ispec] === hess_kl[ijk_ispec] + deltat * hess_tmp
}
close p
else
Expand Down
13 changes: 8 additions & 5 deletions src/gpu/compute_kernels_gpu.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -686,12 +686,15 @@ void FC_FUNC_ (compute_kernels_hess_gpu,
dim3 grid(num_blocks_x,num_blocks_y);
dim3 threads(blocksize,1,1);

compute_hess_kernel<<<grid,threads,0,mp->compute_stream>>>(mp->d_ibool_crust_mantle.cuda,
mp->d_accel_crust_mantle.cuda,
mp->d_b_accel_crust_mantle.cuda,
mp->d_hess_kl_crust_mantle.cuda,
compute_hess_kernel<<<grid,threads,0,mp->compute_stream>>>(mp->d_hess_kl_crust_mantle.cuda,
mp->NSPEC_CRUST_MANTLE,
deltat,
mp->NSPEC_CRUST_MANTLE);
mp->d_ibool_crust_mantle.cuda,
mp->d_b_displ_crust_mantle.cuda,
mp->d_xix_crust_mantle.cuda,mp->d_xiy_crust_mantle.cuda,mp->d_xiz_crust_mantle.cuda,
mp->d_etax_crust_mantle.cuda,mp->d_etay_crust_mantle.cuda,mp->d_etaz_crust_mantle.cuda,
mp->d_gammax_crust_mantle.cuda,mp->d_gammay_crust_mantle.cuda,mp->d_gammaz_crust_mantle.cuda,
mp->d_hprime_xx.cuda);
}
#endif

Expand Down
96 changes: 91 additions & 5 deletions src/gpu/kernels.gen/compute_hess_kernel.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -84,14 +84,100 @@
#define BLOCKSIZE_TRANSFER 256
#endif

__global__ void compute_hess_kernel(const int * ibool, const float * accel, const float * b_accel, float * hess_kl, const float deltat, const int NSPEC_AB){
__global__ void compute_hess_kernel(float * hess_kl, const int NSPEC, const float deltat, const int * d_ibool, const float * d_b_displ, const float * d_xix, const float * d_xiy, const float * d_xiz, const float * d_etax, const float * d_etay, const float * d_etaz, const float * d_gammax, const float * d_gammay, const float * d_gammaz, const float * d_hprime_xx){
int ispec;
int ijk_ispec;
int tx;
int iglob;
__shared__ float s_dummyx_loc[(NGLL3)];
__shared__ float s_dummyy_loc[(NGLL3)];
__shared__ float s_dummyz_loc[(NGLL3)];
__shared__ float sh_hprime_xx[(NGLL2)];
int K;
int J;
int I;
int l;
int offset;
float tempx1l;
float tempx2l;
float tempx3l;
float tempy1l;
float tempy2l;
float tempy3l;
float tempz1l;
float tempz2l;
float tempz3l;
float xixl;
float xiyl;
float xizl;
float etaxl;
float etayl;
float etazl;
float gammaxl;
float gammayl;
float gammazl;
float templ;
float fac1;
float fac2;
float fac3;
float hess_kappa;
float hess_tmp;
ispec = blockIdx.x + (blockIdx.y) * (gridDim.x);
if (ispec < NSPEC_AB) {
ijk_ispec = threadIdx.x + (NGLL3) * (ispec);
iglob = ibool[ijk_ispec] - (1);
hess_kl[ijk_ispec] = hess_kl[ijk_ispec] + (deltat) * ((b_accel[0 + (3) * (iglob)]) * (b_accel[0 + (3) * (iglob)]) + (b_accel[1 + (3) * (iglob)]) * (b_accel[1 + (3) * (iglob)]) + (b_accel[2 + (3) * (iglob)]) * (b_accel[2 + (3) * (iglob)]));
ijk_ispec = threadIdx.x + (NGLL3) * (ispec);
tx = threadIdx.x;
if (tx < NGLL2) {
sh_hprime_xx[tx] = d_hprime_xx[tx];
}
if (ispec < NSPEC) {
iglob = d_ibool[ijk_ispec] - (1);
s_dummyx_loc[tx] = d_b_displ[0 + (3) * (iglob)];
s_dummyy_loc[tx] = d_b_displ[1 + (3) * (iglob)];
s_dummyz_loc[tx] = d_b_displ[2 + (3) * (iglob)];
}
__syncthreads();
if (ispec < NSPEC) {
K = (tx) / (NGLL2);
J = (tx - ((K) * (NGLL2))) / (NGLLX);
I = tx - ((K) * (NGLL2)) - ((J) * (NGLLX));
tempx1l = 0.0f;
tempx2l = 0.0f;
tempx3l = 0.0f;
tempy1l = 0.0f;
tempy2l = 0.0f;
tempy3l = 0.0f;
tempz1l = 0.0f;
tempz2l = 0.0f;
tempz3l = 0.0f;
for (l = 0; l <= NGLLX - (1); l += 1) {
fac1 = sh_hprime_xx[(l) * (NGLLX) + I];
tempx1l = tempx1l + (s_dummyx_loc[(K) * (NGLL2) + (J) * (NGLLX) + l]) * (fac1);
tempy1l = tempy1l + (s_dummyy_loc[(K) * (NGLL2) + (J) * (NGLLX) + l]) * (fac1);
tempz1l = tempz1l + (s_dummyz_loc[(K) * (NGLL2) + (J) * (NGLLX) + l]) * (fac1);
fac2 = sh_hprime_xx[(l) * (NGLLX) + J];
tempx2l = tempx2l + (s_dummyx_loc[(K) * (NGLL2) + (l) * (NGLLX) + I]) * (fac2);
tempy2l = tempy2l + (s_dummyy_loc[(K) * (NGLL2) + (l) * (NGLLX) + I]) * (fac2);
tempz2l = tempz2l + (s_dummyz_loc[(K) * (NGLL2) + (l) * (NGLLX) + I]) * (fac2);
fac3 = sh_hprime_xx[(l) * (NGLLX) + K];
tempx3l = tempx3l + (s_dummyx_loc[(l) * (NGLL2) + (J) * (NGLLX) + I]) * (fac3);
tempy3l = tempy3l + (s_dummyy_loc[(l) * (NGLL2) + (J) * (NGLLX) + I]) * (fac3);
tempz3l = tempz3l + (s_dummyz_loc[(l) * (NGLL2) + (J) * (NGLLX) + I]) * (fac3);
}
offset = (ispec) * (NGLL3_PADDED) + tx;
xixl = d_xix[offset];
etaxl = d_etax[offset];
gammaxl = d_gammax[offset];
xiyl = d_xiy[offset];
etayl = d_etay[offset];
gammayl = d_gammay[offset];
xizl = d_xiz[offset];
etazl = d_etaz[offset];
gammazl = d_gammaz[offset];
hess_kappa = 0.0f;
hess_kappa = hess_kappa + (xixl) * (tempx1l) + (etaxl) * (tempx2l) + (gammaxl) * (tempx3l);
hess_kappa = hess_kappa + (xiyl) * (tempy1l) + (etayl) * (tempy2l) + (gammayl) * (tempy3l);
hess_kappa = hess_kappa + (xizl) * (tempz1l) + (etazl) * (tempz2l) + (gammazl) * (tempz3l);
hess_kappa = (hess_kappa) * (hess_kappa);
hess_tmp = hess_kappa;
hess_kl[ijk_ispec] = hess_kl[ijk_ispec] + (deltat) * (hess_tmp);
}
}
96 changes: 91 additions & 5 deletions src/gpu/kernels.gen/compute_hess_kernel_cl.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -95,15 +95,101 @@ inline void atomicAdd(volatile __global float *source, const float val) {\n\
#define BLOCKSIZE_TRANSFER 256\n\
#endif\n\
\n\
__kernel void compute_hess_kernel(const __global int * ibool, const __global float * accel, const __global float * b_accel, __global float * hess_kl, const float deltat, const int NSPEC_AB){\n\
__kernel void compute_hess_kernel(__global float * hess_kl, const int NSPEC, const float deltat, const __global int * d_ibool, const __global float * d_b_displ, const __global float * d_xix, const __global float * d_xiy, const __global float * d_xiz, const __global float * d_etax, const __global float * d_etay, const __global float * d_etaz, const __global float * d_gammax, const __global float * d_gammay, const __global float * d_gammaz, const __global float * d_hprime_xx){\n\
int ispec;\n\
int ijk_ispec;\n\
int tx;\n\
int iglob;\n\
__local float s_dummyx_loc[(NGLL3)];\n\
__local float s_dummyy_loc[(NGLL3)];\n\
__local float s_dummyz_loc[(NGLL3)];\n\
__local float sh_hprime_xx[(NGLL2)];\n\
int K;\n\
int J;\n\
int I;\n\
int l;\n\
int offset;\n\
float tempx1l;\n\
float tempx2l;\n\
float tempx3l;\n\
float tempy1l;\n\
float tempy2l;\n\
float tempy3l;\n\
float tempz1l;\n\
float tempz2l;\n\
float tempz3l;\n\
float xixl;\n\
float xiyl;\n\
float xizl;\n\
float etaxl;\n\
float etayl;\n\
float etazl;\n\
float gammaxl;\n\
float gammayl;\n\
float gammazl;\n\
float templ;\n\
float fac1;\n\
float fac2;\n\
float fac3;\n\
float hess_kappa;\n\
float hess_tmp;\n\
ispec = get_group_id(0) + (get_group_id(1)) * (get_num_groups(0));\n\
if (ispec < NSPEC_AB) {\n\
ijk_ispec = get_local_id(0) + (NGLL3) * (ispec);\n\
iglob = ibool[ijk_ispec] - (1);\n\
hess_kl[ijk_ispec] = hess_kl[ijk_ispec] + (deltat) * ((b_accel[0 + (3) * (iglob)]) * (b_accel[0 + (3) * (iglob)]) + (b_accel[1 + (3) * (iglob)]) * (b_accel[1 + (3) * (iglob)]) + (b_accel[2 + (3) * (iglob)]) * (b_accel[2 + (3) * (iglob)]));\n\
ijk_ispec = get_local_id(0) + (NGLL3) * (ispec);\n\
tx = get_local_id(0);\n\
if (tx < NGLL2) {\n\
sh_hprime_xx[tx] = d_hprime_xx[tx];\n\
}\n\
if (ispec < NSPEC) {\n\
iglob = d_ibool[ijk_ispec] - (1);\n\
s_dummyx_loc[tx] = d_b_displ[0 + (3) * (iglob)];\n\
s_dummyy_loc[tx] = d_b_displ[1 + (3) * (iglob)];\n\
s_dummyz_loc[tx] = d_b_displ[2 + (3) * (iglob)];\n\
}\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
if (ispec < NSPEC) {\n\
K = (tx) / (NGLL2);\n\
J = (tx - ((K) * (NGLL2))) / (NGLLX);\n\
I = tx - ((K) * (NGLL2)) - ((J) * (NGLLX));\n\
tempx1l = 0.0f;\n\
tempx2l = 0.0f;\n\
tempx3l = 0.0f;\n\
tempy1l = 0.0f;\n\
tempy2l = 0.0f;\n\
tempy3l = 0.0f;\n\
tempz1l = 0.0f;\n\
tempz2l = 0.0f;\n\
tempz3l = 0.0f;\n\
for (l = 0; l <= NGLLX - (1); l += 1) {\n\
fac1 = sh_hprime_xx[(l) * (NGLLX) + I];\n\
tempx1l = tempx1l + (s_dummyx_loc[(K) * (NGLL2) + (J) * (NGLLX) + l]) * (fac1);\n\
tempy1l = tempy1l + (s_dummyy_loc[(K) * (NGLL2) + (J) * (NGLLX) + l]) * (fac1);\n\
tempz1l = tempz1l + (s_dummyz_loc[(K) * (NGLL2) + (J) * (NGLLX) + l]) * (fac1);\n\
fac2 = sh_hprime_xx[(l) * (NGLLX) + J];\n\
tempx2l = tempx2l + (s_dummyx_loc[(K) * (NGLL2) + (l) * (NGLLX) + I]) * (fac2);\n\
tempy2l = tempy2l + (s_dummyy_loc[(K) * (NGLL2) + (l) * (NGLLX) + I]) * (fac2);\n\
tempz2l = tempz2l + (s_dummyz_loc[(K) * (NGLL2) + (l) * (NGLLX) + I]) * (fac2);\n\
fac3 = sh_hprime_xx[(l) * (NGLLX) + K];\n\
tempx3l = tempx3l + (s_dummyx_loc[(l) * (NGLL2) + (J) * (NGLLX) + I]) * (fac3);\n\
tempy3l = tempy3l + (s_dummyy_loc[(l) * (NGLL2) + (J) * (NGLLX) + I]) * (fac3);\n\
tempz3l = tempz3l + (s_dummyz_loc[(l) * (NGLL2) + (J) * (NGLLX) + I]) * (fac3);\n\
}\n\
offset = (ispec) * (NGLL3_PADDED) + tx;\n\
xixl = d_xix[offset];\n\
etaxl = d_etax[offset];\n\
gammaxl = d_gammax[offset];\n\
xiyl = d_xiy[offset];\n\
etayl = d_etay[offset];\n\
gammayl = d_gammay[offset];\n\
xizl = d_xiz[offset];\n\
etazl = d_etaz[offset];\n\
gammazl = d_gammaz[offset];\n\
hess_kappa = 0.0f;\n\
hess_kappa = hess_kappa + (xixl) * (tempx1l) + (etaxl) * (tempx2l) + (gammaxl) * (tempx3l);\n\
hess_kappa = hess_kappa + (xiyl) * (tempy1l) + (etayl) * (tempy2l) + (gammayl) * (tempy3l);\n\
hess_kappa = hess_kappa + (xizl) * (tempz1l) + (etazl) * (tempz2l) + (gammazl) * (tempz3l);\n\
hess_kappa = (hess_kappa) * (hess_kappa);\n\
hess_tmp = hess_kappa;\n\
hess_kl[ijk_ispec] = hess_kl[ijk_ispec] + (deltat) * (hess_tmp);\n\
}\n\
}\n\
";
2 changes: 1 addition & 1 deletion src/gpu/kernels.gen/kernel_proto.cu.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@ __global__ void compute_coupling_ICB_fluid_kernel(const float * displ_inner_core
__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_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 compute_hess_kernel(const int * ibool, const float * accel, const float * b_accel, float * hess_kl, const float deltat, const int NSPEC_AB);
__global__ void compute_hess_kernel(float * hess_kl, const int NSPEC, const float deltat, const int * d_ibool, const float * d_b_displ, const float * d_xix, const float * d_xiy, const float * d_xiz, const float * d_etax, const float * d_etay, const float * d_etaz, const float * d_gammax, const float * d_gammay, const float * d_gammaz, const float * d_hprime_xx);
__global__ void compute_iso_kernel(const float * epsilondev_xx, const float * epsilondev_yy, const float * epsilondev_xy, const float * epsilondev_xz, const float * epsilondev_yz, const float * epsilon_trace_over_3, const float * b_epsilondev_xx, const float * b_epsilondev_yy, const float * b_epsilondev_xy, const float * b_epsilondev_xz, const float * b_epsilondev_yz, const float * b_epsilon_trace_over_3, float * mu_kl, float * kappa_kl, const int NSPEC, const float deltat);
__global__ void compute_iso_undoatt_kernel(const float * epsilondev_xx, const float * epsilondev_yy, const float * epsilondev_xy, const float * epsilondev_xz, const float * epsilondev_yz, const float * epsilon_trace_over_3, float * mu_kl, float * kappa_kl, const int NSPEC, const float deltat, const int * d_ibool, const float * d_b_displ, const float * d_xix, const float * d_xiy, const float * d_xiz, const float * d_etax, const float * d_etay, const float * d_etaz, const float * d_gammax, const float * d_gammay, const float * d_gammaz, const float * d_hprime_xx);
__global__ void compute_rho_kernel(const int * ibool, const float * accel, const float * b_displ, float * rho_kl, const int NSPEC, const float deltat);
Expand Down

0 comments on commit 5329763

Please sign in to comment.