//user function
inline
void flux_calc_kernelz( double *vol_flux_z, const double *zarea,
const double *zvel0, const double *zvel1) {
vol_flux_z[OPS_ACC0(0,0,0)] = 0.125 * dt * (zarea[OPS_ACC1(0,0,0)]) *
( zvel0[OPS_ACC2(0,0,0)] + zvel0[OPS_ACC2(1,0,0)] + zvel0[OPS_ACC2(1,0,0)] + zvel0[OPS_ACC2(1,1,0)] +
zvel1[OPS_ACC3(0,0,0)] + zvel1[OPS_ACC3(1,0,0)] + zvel1[OPS_ACC3(0,1,0)] + zvel1[OPS_ACC3(1,1,0)]);
}
// user function
inline void flux_calc_kernelx(double *vol_flux_x, const double *xarea,
const double *xvel0, const double *xvel1) {
vol_flux_x[OPS_ACC0(0, 0)] =
0.25 * dt * (xarea[OPS_ACC1(0, 0)]) *
((xvel0[OPS_ACC2(0, 0)]) + (xvel0[OPS_ACC2(0, 1)]) +
(xvel1[OPS_ACC3(0, 0)]) + (xvel1[OPS_ACC3(0, 1)]));
}
inline void advec_cell_kernel2_zdir(double *pre_vol, double *post_vol,
const double *volume,
const double *vol_flux_z) {
pre_vol[OPS_ACC0(0, 0, 0)] = volume[OPS_ACC2(0, 0, 0)] +
vol_flux_z[OPS_ACC3(0, 0, 1)] -
vol_flux_z[OPS_ACC3(0, 0, 0)];
post_vol[OPS_ACC1(0, 0, 0)] = volume[OPS_ACC2(0, 0, 0)];
}
// user function
inline void advec_mom_kernel2_x(double *vel1, const double *node_mass_post,
const double *node_mass_pre,
const double *mom_flux) {
vel1[OPS_ACC0(0, 0, 0)] =
(vel1[OPS_ACC0(0, 0, 0)] * node_mass_pre[OPS_ACC2(0, 0, 0)] +
mom_flux[OPS_ACC3(-1, 0, 0)] - mom_flux[OPS_ACC3(0, 0, 0)]) /
node_mass_post[OPS_ACC1(0, 0, 0)];
}
inline void advec_mom_kernel_y2(double *pre_vol, double *post_vol,
const double *volume,
const double *vol_flux_x) {
post_vol[OPS_ACC1(0, 0)] = volume[OPS_ACC2(0, 0)];
pre_vol[OPS_ACC0(0, 0)] = post_vol[OPS_ACC1(0, 0)] +
vol_flux_x[OPS_ACC3(1, 0)] -
vol_flux_x[OPS_ACC3(0, 0)];
}
// user function
inline void flux_calc_kernely(double *vol_flux_y, const double *yarea,
const double *yvel0, const double *yvel1) {
vol_flux_y[OPS_ACC0(0, 0, 0)] =
0.125 * dt * (yarea[OPS_ACC1(0, 0, 0)]) *
(yvel0[OPS_ACC2(0, 0, 0)] + yvel0[OPS_ACC2(1, 0, 0)] +
yvel0[OPS_ACC2(0, 0, 1)] + yvel0[OPS_ACC2(1, 0, 1)] +
yvel1[OPS_ACC3(0, 0, 0)] + yvel1[OPS_ACC3(1, 0, 0)] +
yvel1[OPS_ACC3(0, 0, 1)] + yvel1[OPS_ACC3(1, 0, 1)]);
}
inline void advec_mom_kernel_x1(double *pre_vol, double *post_vol,
const double *volume, const double *vol_flux_x,
const double *vol_flux_y,
const double *vol_flux_z) {
post_vol[OPS_ACC1(0, 0, 0)] =
volume[OPS_ACC2(0, 0, 0)] + vol_flux_y[OPS_ACC4(0, 1, 0)] -
vol_flux_y[OPS_ACC4(0, 0, 0)] + vol_flux_z[OPS_ACC5(0, 0, 1)] -
vol_flux_z[OPS_ACC5(0, 0, 0)];
pre_vol[OPS_ACC0(0, 0, 0)] = post_vol[OPS_ACC1(0, 0, 0)] +
vol_flux_x[OPS_ACC3(1, 0, 0)] -
vol_flux_x[OPS_ACC3(0, 0, 0)];
}
inline void update_halo_kernel1_fr2(double *density0, double *density1,
double *energy0, double *energy1,
double *pressure, double *viscosity,
double *soundspeed , const int* fields) {
if(fields[FIELD_DENSITY0] == 1) density0[OPS_ACC0(0,0,0)] = density0[OPS_ACC0(0,0,-3)];
if(fields[FIELD_DENSITY1] == 1) density1[OPS_ACC1(0,0,0)] = density1[OPS_ACC1(0,0,-3)];
if(fields[FIELD_ENERGY0] == 1) energy0[OPS_ACC2(0,0,0)] = energy0[OPS_ACC2(0,0,-3)];
if(fields[FIELD_ENERGY1] == 1) energy1[OPS_ACC3(0,0,0)] = energy1[OPS_ACC3(0,0,-3)];
if(fields[FIELD_PRESSURE] == 1) pressure[OPS_ACC4(0,0,0)] = pressure[OPS_ACC4(0,0,-3)];
if(fields[FIELD_VISCOSITY] == 1) viscosity[OPS_ACC5(0,0,0)] = viscosity[OPS_ACC5(0,0,-3)];
if(fields[FIELD_SOUNDSPEED] == 1) soundspeed[OPS_ACC6(0,0,0)] = soundspeed[OPS_ACC6(0,0,-3)];
}
// user function
inline void advec_cell_kernel1_ydir(double *pre_vol, double *post_vol,
const double *volume,
const double *vol_flux_x,
const double *vol_flux_y) {
pre_vol[OPS_ACC0(0, 0)] =
volume[OPS_ACC2(0, 0)] +
(vol_flux_y[OPS_ACC4(0, 1)] - vol_flux_y[OPS_ACC4(0, 0)] +
vol_flux_x[OPS_ACC3(1, 0)] - vol_flux_x[OPS_ACC3(0, 0)]);
post_vol[OPS_ACC1(0, 0)] =
pre_vol[OPS_ACC0(0, 0)] -
(vol_flux_y[OPS_ACC4(0, 1)] - vol_flux_y[OPS_ACC4(0, 0)]);
}
//user function
inline void preproc_kernel(const double *u, double *du,
double *ax, double *bx, double *cx, double *ay, double *by, double *cy,
double *az, double *bz, double *cz, int *idx){
double a, b, c, d;
if(idx[0]==0 || idx[0]==nx-1 || idx[1]==0 || idx[1]==ny-1 || idx[2]==0 || idx[2]==nz-1) {
d = 0.0f;
a = 0.0f;
b = 1.0f;
c = 0.0f;
} else {
d = lambda*( u[OPS_ACC0(-1,0,0)] + u[OPS_ACC0(1,0,0)]
+ u[OPS_ACC0(0,-1,0)] + u[OPS_ACC0(0,1,0)]
+ u[OPS_ACC0(0,0,-1)] + u[OPS_ACC0(0,0,1)]
- 6.0f*u[OPS_ACC0(0,0,0)]);
a = -0.5f * lambda;
b = 1.0f + lambda;
c = -0.5f * lambda;
}
du[OPS_ACC1(0,0,0)] = d;
ax[OPS_ACC2(0,0,0)] = a;
bx[OPS_ACC3(0,0,0)] = b;
cx[OPS_ACC4(0,0,0)] = c;
ay[OPS_ACC5(0,0,0)] = a;
by[OPS_ACC6(0,0,0)] = b;
cy[OPS_ACC7(0,0,0)] = c;
az[OPS_ACC8(0,0,0)] = a;
bz[OPS_ACC9(0,0,0)] = b;
cz[OPS_ACC10(0,0,0)] = c;
}
// user function
inline void field_summary_kernel(const double *volume, const double *density0,
const double *energy0, const double *pressure,
const double *xvel0, const double *yvel0,
double *vol, double *mass, double *ie,
double *ke, double *press) {
double vsqrd, cell_vol, cell_mass;
vsqrd = 0.0;
vsqrd = vsqrd +
0.25 * (xvel0[OPS_ACC4(0, 0)] * xvel0[OPS_ACC4(0, 0)] +
yvel0[OPS_ACC5(0, 0)] * yvel0[OPS_ACC5(0, 0)]);
vsqrd = vsqrd +
0.25 * (xvel0[OPS_ACC4(1, 0)] * xvel0[OPS_ACC4(1, 0)] +
yvel0[OPS_ACC5(1, 0)] * yvel0[OPS_ACC5(1, 0)]);
vsqrd = vsqrd +
0.25 * (xvel0[OPS_ACC4(0, 1)] * xvel0[OPS_ACC4(0, 1)] +
yvel0[OPS_ACC5(0, 1)] * yvel0[OPS_ACC5(0, 1)]);
vsqrd = vsqrd +
0.25 * (xvel0[OPS_ACC4(1, 1)] * xvel0[OPS_ACC4(1, 1)] +
yvel0[OPS_ACC5(1, 1)] * yvel0[OPS_ACC5(1, 1)]);
cell_vol = volume[OPS_ACC0(0, 0)];
cell_mass = cell_vol * density0[OPS_ACC1(0, 0)];
*vol = *vol + cell_vol;
*mass = *mass + cell_mass;
*ie = *ie + cell_mass * energy0[OPS_ACC2(0, 0)];
*ke = *ke + cell_mass * 0.5 * vsqrd;
*press = *press + cell_vol * pressure[OPS_ACC3(0, 0)];
}
// user function
inline void calc_dt_kernel_print(const double *xvel0, const double *yvel0,
const double *zvel0, const double *density0,
const double *energy0, const double *pressure,
const double *soundspeed, double *output) {
output[0] = xvel0[OPS_ACC0(0, 0, 0)];
output[1] = yvel0[OPS_ACC1(0, 0, 0)];
output[2] = zvel0[OPS_ACC2(0, 0, 0)];
output[3] = xvel0[OPS_ACC0(1, 0, 0)];
output[4] = yvel0[OPS_ACC1(1, 0, 0)];
output[5] = zvel0[OPS_ACC2(0, 0, 0)];
output[6] = xvel0[OPS_ACC0(1, 1, 0)];
output[7] = yvel0[OPS_ACC1(1, 1, 0)];
output[8] = zvel0[OPS_ACC2(0, 0, 0)];
output[9] = xvel0[OPS_ACC0(0, 1, 0)];
output[10] = yvel0[OPS_ACC1(0, 1, 0)];
output[11] = zvel0[OPS_ACC2(0, 0, 0)];
output[12] = xvel0[OPS_ACC0(0, 0, 1)];
output[13] = yvel0[OPS_ACC1(0, 0, 1)];
output[14] = zvel0[OPS_ACC2(0, 0, 1)];
output[15] = xvel0[OPS_ACC0(1, 0, 1)];
output[16] = yvel0[OPS_ACC1(1, 0, 1)];
output[17] = zvel0[OPS_ACC2(0, 0, 1)];
output[18] = xvel0[OPS_ACC0(1, 1, 1)];
output[19] = yvel0[OPS_ACC1(1, 1, 1)];
output[20] = zvel0[OPS_ACC2(0, 0, 1)];
output[21] = xvel0[OPS_ACC0(0, 1, 1)];
output[22] = yvel0[OPS_ACC1(0, 1, 1)];
output[23] = zvel0[OPS_ACC2(0, 0, 1)];
output[24] = density0[OPS_ACC3(0, 0, 0)];
output[25] = energy0[OPS_ACC4(0, 0, 0)];
output[26] = pressure[OPS_ACC5(0, 0, 0)];
output[27] = soundspeed[OPS_ACC6(0, 0, 0)];
}
inline void update_halo_kernel1_r2(double *density0, double *energy0,
double *energy1, double *u, double *p,
double *sd, const int *fields) {
if (fields[FIELD_DENSITY] == 1)
density0[OPS_ACC0(0, 0)] = density0[OPS_ACC0(-3, 0)];
if (fields[FIELD_ENERGY0] == 1)
energy0[OPS_ACC1(0, 0)] = energy0[OPS_ACC1(-3, 0)];
if (fields[FIELD_ENERGY1] == 1)
energy1[OPS_ACC2(0, 0)] = energy1[OPS_ACC2(-3, 0)];
if (fields[FIELD_U] == 1)
u[OPS_ACC3(0, 0)] = u[OPS_ACC3(-3, 0)];
if (fields[FIELD_P] == 1)
p[OPS_ACC4(0, 0)] = p[OPS_ACC4(-3, 0)];
if (fields[FIELD_SD] == 1)
sd[OPS_ACC5(0, 0)] = sd[OPS_ACC5(-3, 0)];
}
// user function
inline void drhouupdx_kernel(const double *rhou_new, const double *rho_new,
const double *rhoE_new, double *rhou_res) {
double fni =
rhou_new[OPS_ACC0(0)] * rhou_new[OPS_ACC0(0)] / rho_new[OPS_ACC1(0)];
double p = gam1 * (rhoE_new[OPS_ACC2(0)] - 0.5 * fni);
fni = fni + p;
double fnim1 =
rhou_new[OPS_ACC0(-1)] * rhou_new[OPS_ACC0(-1)] / rho_new[OPS_ACC1(-1)];
p = gam1 * (rhoE_new[OPS_ACC2(-1)] - 0.5 * fnim1);
fnim1 = fnim1 + p;
double fnim2 =
rhou_new[OPS_ACC0(-2)] * rhou_new[OPS_ACC0(-2)] / rho_new[OPS_ACC1(-2)];
p = gam1 * (rhoE_new[OPS_ACC2(-2)] - 0.5 * fnim2);
fnim2 = fnim2 + p;
double fnip1 =
rhou_new[OPS_ACC0(1)] * rhou_new[OPS_ACC0(1)] / rho_new[OPS_ACC1(1)];
p = gam1 * (rhoE_new[OPS_ACC2(1)] - 0.5 * fnip1);
fnip1 = fnip1 + p;
double fnip2 =
rhou_new[OPS_ACC0(2)] * rhou_new[OPS_ACC0(2)] / rho_new[OPS_ACC1(2)];
p = gam1 * (rhoE_new[OPS_ACC2(2)] - 0.5 * fnip2);
fnip2 = fnip2 + p;
double deriv = (fnim2 - fnip2 + 8.0 * (fnip1 - fnim1)) / (12.00 * dx);
rhou_res[OPS_ACC3(0)] = deriv;
}
// user function
inline void save_kernel(double *rho_old, double *rhou_old, double *rhoE_old,
const double *rho_new, const double *rhou_new,
const double *rhoE_new) {
rho_old[OPS_ACC0(0)] = rho_new[OPS_ACC3(0)];
rhou_old[OPS_ACC1(0)] = rhou_new[OPS_ACC4(0)];
rhoE_old[OPS_ACC2(0)] = rhoE_new[OPS_ACC5(0)];
}
// user function
inline void calc_dt_kernel(const double *celldx, const double *celldy,
const double *soundspeed, const double *viscosity,
const double *density0, const double *xvel0,
const double *xarea, const double *volume,
const double *yvel0, const double *yarea,
double *dt_min, const double *celldz,
const double *zvel0, const double *zarea) {
double div, ds, dtut, dtvt, dtct, dtwt, dtdivt, cc, dv1, dv2, du1, du2, dw1,
dw2;
ds = MIN(MIN(celldx[OPS_ACC0(0, 0, 0)], celldy[OPS_ACC1(0, 0, 0)]),
celldz[OPS_ACC11(0, 0, 0)]);
ds = 1.0 / (ds * ds);
cc = soundspeed[OPS_ACC2(0, 0, 0)] * soundspeed[OPS_ACC2(0, 0, 0)];
cc = cc + 2.0 * viscosity[OPS_ACC3(0, 0, 0)] / density0[OPS_ACC4(0, 0, 0)];
dtct = ds * cc;
dtct = dtc_safe * 1.0 / MAX(sqrt(dtct), g_small);
du1 = (xvel0[OPS_ACC5(0, 0, 0)] + xvel0[OPS_ACC5(0, 1, 0)] +
xvel0[OPS_ACC5(0, 0, 1)] + xvel0[OPS_ACC5(0, 1, 1)]) *
xarea[OPS_ACC6(0, 0, 0)];
du2 = (xvel0[OPS_ACC5(1, 0, 0)] + xvel0[OPS_ACC5(1, 1, 0)] +
xvel0[OPS_ACC5(1, 0, 1)] + xvel0[OPS_ACC5(1, 1, 1)]) *
xarea[OPS_ACC6(0, 0, 0)];
dtut = dtu_safe * 4.0 * volume[OPS_ACC7(0, 0, 0)] /
MAX(MAX(fabs(du1), fabs(du2)), 1.0e-5 * volume[OPS_ACC7(0, 0, 0)]);
dv1 = (yvel0[OPS_ACC8(0, 0, 0)] + yvel0[OPS_ACC8(1, 0, 0)] +
yvel0[OPS_ACC8(0, 0, 1)] + yvel0[OPS_ACC8(1, 0, 1)]) *
yarea[OPS_ACC9(0, 0, 0)];
dv2 = (yvel0[OPS_ACC8(0, 1, 0)] + yvel0[OPS_ACC8(1, 1, 0)] +
yvel0[OPS_ACC8(0, 1, 1)] + yvel0[OPS_ACC8(1, 1, 1)]) *
yarea[OPS_ACC9(0, 0, 0)];
dtvt = dtv_safe * 4.0 * volume[OPS_ACC7(0, 0, 0)] /
MAX(MAX(fabs(dv1), fabs(dv2)), 1.0e-5 * volume[OPS_ACC7(0, 0, 0)]);
dw1 = (zvel0[OPS_ACC12(0, 0, 0)] + zvel0[OPS_ACC12(0, 1, 0)] +
zvel0[OPS_ACC12(1, 0, 0)] + zvel0[OPS_ACC12(1, 1, 0)]) *
zarea[OPS_ACC13(0, 0, 0)];
dw2 = (zvel0[OPS_ACC12(0, 0, 1)] + zvel0[OPS_ACC12(0, 1, 1)] +
zvel0[OPS_ACC12(1, 0, 1)] + zvel0[OPS_ACC12(1, 1, 1)]) *
zarea[OPS_ACC13(0, 0, 0)];
dtwt = dtw_safe * 4.0 * volume[OPS_ACC7(0, 0, 0)] /
MAX(MAX(fabs(dw1), fabs(dw2)), 1.0e-5 * volume[OPS_ACC7(0, 0, 0)]);
div = du2 - du1 + dv2 - dv1 + dw2 - dw1;
dtdivt = dtdiv_safe * 4.0 * (volume[OPS_ACC7(0, 0, 0)]) /
MAX(volume[OPS_ACC7(0, 0, 0)] * 1.0e-05, fabs(div));
dt_min[OPS_ACC10(0, 0, 0)] =
MIN(MIN(MIN(dtct, dtut), MIN(dtvt, dtdivt)), dtwt);
}
//user function
inline
void poisson_kernel_populate(const int *dispx, const int *dispy, const int *idx, double *u, double *f, double *ref) {
double x = dx * (double)(idx[0]+dispx[0]);
double y = dy * (double)(idx[1]+dispy[0]);
u[OPS_ACC3(0,0)] = sin(M_PI*x)*cos(2.0*M_PI*y);
f[OPS_ACC4(0,0)] = -5.0*M_PI*M_PI*sin(M_PI*x)*cos(2.0*M_PI*y);
ref[OPS_ACC5(0,0)] = sin(M_PI*x)*cos(2.0*M_PI*y);
}
inline void advec_mom_kernel1_y_nonvector(const double *node_flux,
const double *node_mass_pre,
double *mom_flux,
const double *celldy,
const double *vel1) {
double sigma, wind, width;
double vdiffuw, vdiffdw, auw, adw, limiter;
int upwind, donor, downwind, dif;
double advec_vel_temp;
if ((node_flux[OPS_ACC0(0, 0, 0)]) < 0.0) {
upwind = 2;
donor = 1;
downwind = 0;
dif = donor;
} else {
upwind = -1;
donor = 0;
downwind = 1;
dif = upwind;
}
sigma =
fabs(node_flux[OPS_ACC0(0, 0, 0)]) / node_mass_pre[OPS_ACC1(0, donor, 0)];
width = celldy[OPS_ACC3(0, 0, 0)];
vdiffuw = vel1[OPS_ACC4(0, donor, 0)] - vel1[OPS_ACC4(0, upwind, 0)];
vdiffdw = vel1[OPS_ACC4(0, downwind, 0)] - vel1[OPS_ACC4(0, donor, 0)];
limiter = 0.0;
if (vdiffuw * vdiffdw > 0.0) {
auw = fabs(vdiffuw);
adw = fabs(vdiffdw);
wind = 1.0;
if (vdiffdw <= 0.0)
wind = -1.0;
limiter =
wind * MIN(width * ((2.0 - sigma) * adw / width +
(1.0 + sigma) * auw / celldy[OPS_ACC3(0, dif, 0)]) /
6.0,
MIN(auw, adw));
}
advec_vel_temp = vel1[OPS_ACC4(0, donor, 0)] + (1.0 - sigma) * limiter;
mom_flux[OPS_ACC2(0, 0, 0)] = advec_vel_temp * node_flux[OPS_ACC0(0, 0, 0)];
}
// user function
inline void initialise_chunk_kernel_volume(double *volume, const double *celldy,
double *xarea, const double *celldx,
double *yarea, const double *celldz,
double *zarea) {
double d_x, d_y, d_z;
d_x = (grid.xmax - grid.xmin) / (double)grid.x_cells;
d_y = (grid.ymax - grid.ymin) / (double)grid.y_cells;
d_z = (grid.zmax - grid.zmin) / (double)grid.z_cells;
volume[OPS_ACC0(0, 0, 0)] = d_x * d_y * d_z;
xarea[OPS_ACC2(0, 0, 0)] =
celldy[OPS_ACC1(0, 0, 0)] * celldz[OPS_ACC5(0, 0, 0)];
yarea[OPS_ACC4(0, 0, 0)] =
celldx[OPS_ACC3(0, 0, 0)] * celldz[OPS_ACC5(0, 0, 0)];
zarea[OPS_ACC6(0, 0, 0)] =
celldx[OPS_ACC3(0, 0, 0)] * celldy[OPS_ACC1(0, 0, 0)];
}
// user function
inline void calvar_kernel(const double *rho_new, const double *rhou_new,
const double *rhoE_new, double *workarray2,
double *workarray3) {
double p, rhoi, u;
rhoi = 1 / rho_new[OPS_ACC0(0)];
u = rhou_new[OPS_ACC1(0)] * rhoi;
p = gam1 * (rhoE_new[OPS_ACC2(0)] - 0.5 * rho_new[OPS_ACC0(0)] * u * u);
workarray2[OPS_ACC3(0)] = p + rhou_new[OPS_ACC1(0)] * u;
workarray3[OPS_ACC4(0)] = (p + rhoE_new[OPS_ACC2(0)]) * u;
}
// user function
inline void updateRK3_kernel(double *rho_new, double *rhou_new,
double *rhoE_new, double *rho_old,
double *rhou_old, double *rhoE_old,
const double *rho_res, const double *rhou_res,
const double *rhoE_res, const double *a1,
const double *a2) {
rho_new[OPS_ACC0(0)] =
rho_old[OPS_ACC3(0)] + dt * a1[0] * (-rho_res[OPS_ACC6(0)]);
rhou_new[OPS_ACC1(0)] =
rhou_old[OPS_ACC4(0)] + dt * a1[0] * (-rhou_res[OPS_ACC7(0)]);
rhoE_new[OPS_ACC2(0)] =
rhoE_old[OPS_ACC5(0)] + dt * a1[0] * (-rhoE_res[OPS_ACC8(0)]);
rho_old[OPS_ACC3(0)] =
rho_old[OPS_ACC3(0)] + dt * a2[0] * (-rho_res[OPS_ACC6(0)]);
rhou_old[OPS_ACC4(0)] =
rhou_old[OPS_ACC4(0)] + dt * a2[0] * (-rhou_res[OPS_ACC7(0)]);
rhoE_old[OPS_ACC5(0)] =
rhoE_old[OPS_ACC5(0)] + dt * a2[0] * (-rhoE_res[OPS_ACC8(0)]);
}
inline void advec_mom_kernel_post_pre_advec_y( double *node_mass_post, const double *post_vol,
const double *density1, double *node_mass_pre, const double *node_flux) {
node_mass_post[OPS_ACC0(0,0)] = 0.25 * ( density1[OPS_ACC2(0,-1)] * post_vol[OPS_ACC1(0,-1)] +
density1[OPS_ACC2(0,0)] * post_vol[OPS_ACC1(0,0)] +
density1[OPS_ACC2(-1,-1)] * post_vol[OPS_ACC1(-1,-1)] +
density1[OPS_ACC2(-1,0)] * post_vol[OPS_ACC1(-1,0)] );
node_mass_pre[OPS_ACC3(0,0)] = node_mass_post[OPS_ACC0(0,0)] - node_flux[OPS_ACC4(0,-1)] + node_flux[OPS_ACC4(0,0)];
}
// user function
inline void initialize_kernel(double *x, double *rho_new, double *rhou_new,
double *rhoE_new, double *rhoin, int *idx) {
x[OPS_ACC0(0)] = xmin + (idx[0] - 2) * dx;
if (x[OPS_ACC0(0)] >= -4.0) {
rho_new[OPS_ACC1(0)] = 1.0 + eps * sin(lambda * x[OPS_ACC0(0)]);
rhou_new[OPS_ACC2(0)] = ur * rho_new[OPS_ACC1(0)];
rhoE_new[OPS_ACC3(0)] =
(pr / gam1) +
0.5 * pow(rhou_new[OPS_ACC2(0)], 2) / rho_new[OPS_ACC1(0)];
} else {
rho_new[OPS_ACC1(0)] = rhol;
rhou_new[OPS_ACC2(0)] = ul * rho_new[OPS_ACC1(0)];
rhoE_new[OPS_ACC3(0)] =
(pl / gam1) +
0.5 * pow(rhou_new[OPS_ACC2(0)], 2) / rho_new[OPS_ACC1(0)];
}
rhoin[OPS_ACC4(0)] =
gam1 * (rhoE_new[OPS_ACC3(0)] -
0.5 * rhou_new[OPS_ACC2(0)] * rhou_new[OPS_ACC2(0)] /
rho_new[OPS_ACC1(0)]);
}
// user function
inline void
advec_cell_kernel4_zdir(double *density1, double *energy1,
const double *mass_flux_z, const double *vol_flux_z,
const double *pre_vol, const double *post_vol,
double *pre_mass, double *post_mass, double *advec_vol,
double *post_ener, const double *ener_flux) {
pre_mass[OPS_ACC6(0, 0, 0)] =
density1[OPS_ACC0(0, 0, 0)] * pre_vol[OPS_ACC4(0, 0, 0)];
post_mass[OPS_ACC7(0, 0, 0)] = pre_mass[OPS_ACC6(0, 0, 0)] +
mass_flux_z[OPS_ACC2(0, 0, 0)] -
mass_flux_z[OPS_ACC2(0, 0, 1)];
post_ener[OPS_ACC9(0, 0, 0)] =
(energy1[OPS_ACC1(0, 0, 0)] * pre_mass[OPS_ACC6(0, 0, 0)] +
ener_flux[OPS_ACC10(0, 0, 0)] - ener_flux[OPS_ACC10(0, 0, 1)]) /
post_mass[OPS_ACC7(0, 0, 0)];
advec_vol[OPS_ACC8(0, 0, 0)] = pre_vol[OPS_ACC4(0, 0, 0)] +
vol_flux_z[OPS_ACC3(0, 0, 0)] -
vol_flux_z[OPS_ACC3(0, 0, 1)];
density1[OPS_ACC0(0, 0, 0)] =
post_mass[OPS_ACC7(0, 0, 0)] / advec_vol[OPS_ACC8(0, 0, 0)];
energy1[OPS_ACC1(0, 0, 0)] = post_ener[OPS_ACC9(0, 0, 0)];
}
// user function
inline void ideal_gas_kernel(const double *density, const double *energy,
double *pressure, double *soundspeed) {
double sound_speed_squared, v, pressurebyenergy, pressurebyvolume;
v = 1.0 / density[OPS_ACC0(0, 0, 0)];
pressure[OPS_ACC2(0, 0, 0)] =
(1.4 - 1.0) * density[OPS_ACC0(0, 0, 0)] * energy[OPS_ACC1(0, 0, 0)];
pressurebyenergy = (1.4 - 1.0) * density[OPS_ACC0(0, 0, 0)];
pressurebyvolume =
-1.0 * density[OPS_ACC0(0, 0, 0)] * pressure[OPS_ACC2(0, 0, 0)];
sound_speed_squared =
v * v *
(pressure[OPS_ACC2(0, 0, 0)] * pressurebyenergy - pressurebyvolume);
soundspeed[OPS_ACC3(0, 0, 0)] = sqrt(sound_speed_squared);
}
// user function
inline void calc_dt_kernel_print(const double *xvel0, const double *yvel0,
const double *density0, const double *energy0,
const double *pressure,
const double *soundspeed, double *output) {
output[0] = xvel0[OPS_ACC0(1, 0)];
output[1] = yvel0[OPS_ACC1(1, 0)];
output[2] = xvel0[OPS_ACC0(-1, 0)];
output[3] = yvel0[OPS_ACC1(-1, 0)];
output[4] = xvel0[OPS_ACC0(0, 1)];
output[5] = yvel0[OPS_ACC1(0, 1)];
output[6] = xvel0[OPS_ACC0(0, -1)];
output[7] = yvel0[OPS_ACC1(0, -1)];
output[8] = density0[OPS_ACC2(0, 0)];
output[9] = energy0[OPS_ACC3(0, 0)];
output[10] = pressure[OPS_ACC4(0, 0)];
output[11] = soundspeed[OPS_ACC5(0, 0)];
}
// user function
inline void calc_dt_kernel(const double *celldx, const double *celldy,
const double *soundspeed, const double *viscosity,
const double *density0, const double *xvel0,
const double *xarea, const double *volume,
const double *yvel0, const double *yarea,
double *dt_min) {
double div, dsx, dsy, dtut, dtvt, dtct, dtdivt, cc, dv1, dv2;
dsx = celldx[OPS_ACC0(0, 0)];
dsy = celldy[OPS_ACC1(0, 0)];
cc = soundspeed[OPS_ACC2(0, 0)] * soundspeed[OPS_ACC2(0, 0)];
cc = cc + 2.0 * viscosity[OPS_ACC3(0, 0)] / density0[OPS_ACC4(0, 0)];
cc = MAX(sqrt(cc), g_small);
dtct = dtc_safe * MIN(dsx, dsy) / cc;
div = 0.0;
dv1 = (xvel0[OPS_ACC5(0, 0)] + xvel0[OPS_ACC5(0, 1)]) * xarea[OPS_ACC6(0, 0)];
dv2 = (xvel0[OPS_ACC5(1, 0)] + xvel0[OPS_ACC5(1, 1)]) * xarea[OPS_ACC6(1, 0)];
div = div + dv2 - dv1;
dtut = dtu_safe * 2.0 * volume[OPS_ACC7(0, 0)] /
MAX(MAX(fabs(dv1), fabs(dv2)), g_small * volume[OPS_ACC7(0, 0)]);
dv1 = (yvel0[OPS_ACC8(0, 0)] + yvel0[OPS_ACC8(1, 0)]) * yarea[OPS_ACC9(0, 0)];
dv2 = (yvel0[OPS_ACC8(0, 1)] + yvel0[OPS_ACC8(1, 1)]) * yarea[OPS_ACC9(0, 1)];
div = div + dv2 - dv1;
dtvt = dtv_safe * 2.0 * volume[OPS_ACC7(0, 0)] /
MAX(MAX(fabs(dv1), fabs(dv2)), g_small * volume[OPS_ACC7(0, 0)]);
div = div / (2.0 * volume[OPS_ACC7(0, 0)]);
if (div < -g_small)
dtdivt = dtdiv_safe * (-1.0 / div);
else
dtdivt = g_big;
dt_min[OPS_ACC10(0, 0)] = MIN(MIN(dtct, dtut), MIN(dtvt, dtdivt));
}
// user function
inline void tea_leaf_cg_calc_w_reduce_kernel(double *w, const double *Kx,
const double *Ky, const double *p,
const double *rx, const double *ry,
double *pw) {
w[OPS_ACC0(0, 0)] = (1.0 + (*ry) * (Ky[OPS_ACC2(0, 1)] + Ky[OPS_ACC2(0, 0)]) +
(*rx) * (Kx[OPS_ACC1(1, 0)] + Kx[OPS_ACC1(0, 0)])) *
p[OPS_ACC3(0, 0)] -
(*ry) * (Ky[OPS_ACC2(0, 1)] * p[OPS_ACC3(0, 1)] +
Ky[OPS_ACC2(0, 0)] * p[OPS_ACC3(0, -1)]) -
(*rx) * (Kx[OPS_ACC1(1, 0)] * p[OPS_ACC3(1, 0)] +
Kx[OPS_ACC1(0, 0)] * p[OPS_ACC3(-1, 0)]);
*pw = *pw + w[OPS_ACC0(0, 0)] * p[OPS_ACC3(0, 0)];
}
// host stub function
void ops_par_loop_advec_cell_kernel2_ydir_execute(ops_kernel_descriptor *desc) {
ops_block block = desc->block;
int dim = desc->dim;
int *range = desc->range;
ops_arg arg0 = desc->args[0];
ops_arg arg1 = desc->args[1];
ops_arg arg2 = desc->args[2];
ops_arg arg3 = desc->args[3];
// Timing
double t1, t2, c1, c2;
ops_arg args[4] = {arg0, arg1, arg2, arg3};
#ifdef CHECKPOINTING
if (!ops_checkpointing_before(args, 4, range, 66))
return;
#endif
if (OPS_diags > 1) {
OPS_kernels[66].count++;
ops_timers_core(&c2, &t2);
}
// compute locally allocated range for the sub-block
int start[2];
int end[2];
for (int n = 0; n < 2; n++) {
start[n] = range[2 * n];
end[n] = range[2 * n + 1];
}
#ifdef OPS_DEBUG
ops_register_args(args, "advec_cell_kernel2_ydir");
#endif
// set up initial pointers and exchange halos if necessary
int base0 = args[0].dat->base_offset;
double *__restrict__ pre_vol = (double *)(args[0].data + base0);
int base1 = args[1].dat->base_offset;
double *__restrict__ post_vol = (double *)(args[1].data + base1);
int base2 = args[2].dat->base_offset;
const double *__restrict__ volume = (double *)(args[2].data + base2);
int base3 = args[3].dat->base_offset;
const double *__restrict__ vol_flux_y = (double *)(args[3].data + base3);
// initialize global variable with the dimension of dats
int xdim0_advec_cell_kernel2_ydir = args[0].dat->size[0];
int xdim1_advec_cell_kernel2_ydir = args[1].dat->size[0];
int xdim2_advec_cell_kernel2_ydir = args[2].dat->size[0];
int xdim3_advec_cell_kernel2_ydir = args[3].dat->size[0];
if (OPS_diags > 1) {
ops_timers_core(&c1, &t1);
OPS_kernels[66].mpi_time += t1 - t2;
}
#pragma omp parallel for
for (int n_y = start[1]; n_y < end[1]; n_y++) {
#ifdef intel
#pragma loop_count(10000)
#pragma omp simd aligned(pre_vol, post_vol, volume, vol_flux_y)
#else
#pragma simd
#endif
for (int n_x = start[0]; n_x < end[0]; n_x++) {
pre_vol[OPS_ACC0(0, 0)] = volume[OPS_ACC2(0, 0)] +
vol_flux_y[OPS_ACC3(0, 1)] -
vol_flux_y[OPS_ACC3(0, 0)];
post_vol[OPS_ACC1(0, 0)] = volume[OPS_ACC2(0, 0)];
}
}
if (OPS_diags > 1) {
ops_timers_core(&c2, &t2);
OPS_kernels[66].time += t2 - t1;
}
if (OPS_diags > 1) {
// Update kernel record
ops_timers_core(&c1, &t1);
OPS_kernels[66].mpi_time += t1 - t2;
OPS_kernels[66].transfer += ops_compute_transfer(dim, start, end, &arg0);
OPS_kernels[66].transfer += ops_compute_transfer(dim, start, end, &arg1);
OPS_kernels[66].transfer += ops_compute_transfer(dim, start, end, &arg2);
OPS_kernels[66].transfer += ops_compute_transfer(dim, start, end, &arg3);
}
}
// host stub function
void ops_par_loop_update_halo_kernel1_b2_execute(ops_kernel_descriptor *desc) {
ops_block block = desc->block;
int dim = desc->dim;
int *range = desc->range;
ops_arg arg0 = desc->args[0];
ops_arg arg1 = desc->args[1];
ops_arg arg2 = desc->args[2];
ops_arg arg3 = desc->args[3];
ops_arg arg4 = desc->args[4];
ops_arg arg5 = desc->args[5];
ops_arg arg6 = desc->args[6];
ops_arg arg7 = desc->args[7];
// Timing
double t1, t2, c1, c2;
ops_arg args[8] = {arg0, arg1, arg2, arg3, arg4, arg5, arg6, arg7};
#ifdef CHECKPOINTING
if (!ops_checkpointing_before(args, 8, range, 9))
return;
#endif
if (OPS_diags > 1) {
OPS_kernels[9].count++;
ops_timers_core(&c2, &t2);
}
// compute locally allocated range for the sub-block
int start[2];
int end[2];
for (int n = 0; n < 2; n++) {
start[n] = range[2 * n];
end[n] = range[2 * n + 1];
}
#ifdef OPS_DEBUG
ops_register_args(args, "update_halo_kernel1_b2");
#endif
// set up initial pointers and exchange halos if necessary
int base0 = args[0].dat->base_offset;
double *__restrict__ density0 = (double *)(args[0].data + base0);
int base1 = args[1].dat->base_offset;
double *__restrict__ density1 = (double *)(args[1].data + base1);
int base2 = args[2].dat->base_offset;
double *__restrict__ energy0 = (double *)(args[2].data + base2);
int base3 = args[3].dat->base_offset;
double *__restrict__ energy1 = (double *)(args[3].data + base3);
int base4 = args[4].dat->base_offset;
double *__restrict__ pressure = (double *)(args[4].data + base4);
int base5 = args[5].dat->base_offset;
double *__restrict__ viscosity = (double *)(args[5].data + base5);
int base6 = args[6].dat->base_offset;
double *__restrict__ soundspeed = (double *)(args[6].data + base6);
const int *__restrict__ fields = (int *)args[7].data;
// initialize global variable with the dimension of dats
int xdim0_update_halo_kernel1_b2 = args[0].dat->size[0];
int xdim1_update_halo_kernel1_b2 = args[1].dat->size[0];
int xdim2_update_halo_kernel1_b2 = args[2].dat->size[0];
int xdim3_update_halo_kernel1_b2 = args[3].dat->size[0];
int xdim4_update_halo_kernel1_b2 = args[4].dat->size[0];
int xdim5_update_halo_kernel1_b2 = args[5].dat->size[0];
int xdim6_update_halo_kernel1_b2 = args[6].dat->size[0];
if (OPS_diags > 1) {
ops_timers_core(&c1, &t1);
OPS_kernels[9].mpi_time += t1 - t2;
}
#pragma omp parallel for
for (int n_y = start[1]; n_y < end[1]; n_y++) {
#ifdef intel
#pragma loop_count(10000)
#pragma omp simd aligned(density0, density1, energy0, energy1, pressure, \
viscosity, soundspeed)
#else
#pragma simd
#endif
for (int n_x = start[0]; n_x < end[0]; n_x++) {
if (fields[FIELD_DENSITY0] == 1)
density0[OPS_ACC0(0, 0)] = density0[OPS_ACC0(0, 3)];
if (fields[FIELD_DENSITY1] == 1)
density1[OPS_ACC1(0, 0)] = density1[OPS_ACC1(0, 3)];
if (fields[FIELD_ENERGY0] == 1)
energy0[OPS_ACC2(0, 0)] = energy0[OPS_ACC2(0, 3)];
if (fields[FIELD_ENERGY1] == 1)
energy1[OPS_ACC3(0, 0)] = energy1[OPS_ACC3(0, 3)];
if (fields[FIELD_PRESSURE] == 1)
pressure[OPS_ACC4(0, 0)] = pressure[OPS_ACC4(0, 3)];
if (fields[FIELD_VISCOSITY] == 1)
viscosity[OPS_ACC5(0, 0)] = viscosity[OPS_ACC5(0, 3)];
if (fields[FIELD_SOUNDSPEED] == 1)
soundspeed[OPS_ACC6(0, 0)] = soundspeed[OPS_ACC6(0, 3)];
}
}
if (OPS_diags > 1) {
ops_timers_core(&c2, &t2);
OPS_kernels[9].time += t2 - t1;
}
if (OPS_diags > 1) {
// Update kernel record
ops_timers_core(&c1, &t1);
OPS_kernels[9].mpi_time += t1 - t2;
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg0);
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg1);
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg2);
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg3);
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg4);
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg5);
OPS_kernels[9].transfer += ops_compute_transfer(dim, start, end, &arg6);
}
}
