Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

New YAKL 2022-02-18 #4804

Merged
merged 6 commits into from
Mar 9, 2022
Merged

New YAKL 2022-02-18 #4804

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
082cd9e
Upgrading to the new YAKL. There was an issue with FFTs. I had to add…
mrnorman Feb 18, 2022
7220dc4
New YAKL
mrnorman Feb 18, 2022
a89589c
New YAKL with fixed unit tests
mrnorman Feb 18, 2022
a9e23a9
bug fix for ESMT FFT init
whannah1 Mar 1, 2022
2a49fb4
This should work around the lbounds and strides linking bugs in IBM a…
mrnorman Mar 4, 2022
5ec41aa
Upgrading YAKL again to enable __host__ __device__ atomics, which rem…
mrnorman Mar 7, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions components/eam/src/physics/crm/samxx/accelerate_crm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,7 +52,7 @@ void accelerate_crm(int nstep, int nstop, bool &ceaseflag) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
// calculate tendency * dtn
yakl::atomicAdd( tbaccel(k,icrm) , t(k,j+offy_s,i+offx_s,icrm) * crm_accel_coef );
yakl::atomicAdd( qtbaccel(k,icrm) , (qcl(k,j,i,icrm) + qci(k,j,i,icrm) + qv(k,j,i,icrm)) * crm_accel_coef );
Expand Down Expand Up @@ -144,7 +144,7 @@ void accelerate_crm(int nstep, int nstop, bool &ceaseflag) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (micro_field(idx_qt,k,j+offy_s,i+offx_s,icrm) < 0.0) {
yakl::atomicAdd( qneg(k,icrm) , micro_field(idx_qt,k,j+offy_s,i+offx_s,icrm) );
}
Expand Down
4 changes: 2 additions & 2 deletions components/eam/src/physics/crm/samxx/advect2_mom_z.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,7 @@ void advect2_mom_z() {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm-1,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm-1,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real dz25=1.0/(4.0*dz(icrm));
int kb = k-1;
real rhoi = dz25 * rhow(k+1,icrm);
Expand All @@ -66,7 +66,7 @@ void advect2_mom_z() {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm-1,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm-1,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real dz25=1.0/(4.0*dz(icrm));
int kb = k-1;
real rhoi = dz25 * rhow(k+1,icrm);
Expand Down
6 changes: 3 additions & 3 deletions components/eam/src/physics/crm/samxx/advect_scalar.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ void advect_scalar(real4d &f, real2d &fadv, real2d &flux) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp = f(k,j+offy_s,i+offx_s,icrm)-f0(k,j+offy_s,i+offx_s,icrm);
yakl::atomicAdd(fadv(k,icrm),tmp);
});
Expand Down Expand Up @@ -87,7 +87,7 @@ void advect_scalar(real5d &f, int ind_f, real2d &fadv, real2d &flux) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp = f(ind_f,k,j+offy_s,i+offx_s,icrm)-f0(k,j+offy_s,i+offx_s,icrm);
yakl::atomicAdd(fadv(k,icrm),tmp);
});
Expand Down Expand Up @@ -134,7 +134,7 @@ void advect_scalar(real5d &f, int ind_f, real3d &fadv, int ind_fadv, real3d &flu
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp = f(ind_f,k,j+offy_s,i+offx_s,icrm)-f0(k,j+offy_s,i+offx_s,icrm);
yakl::atomicAdd(fadv(ind_fadv,k,icrm),tmp);
});
Expand Down
12 changes: 6 additions & 6 deletions components/eam/src/physics/crm/samxx/advect_scalar2D.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,7 @@ void advect_scalar2D(real4d &f, real2d &flux) {
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx+4; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx+4,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx+4,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
if (i >= 2 && i <= nx+1) {
yakl::atomicAdd(flux(k,icrm),www(k,j,i,icrm));
}
Expand Down Expand Up @@ -163,7 +163,7 @@ void advect_scalar2D(real4d &f, real2d &flux) {
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx+1; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx+1,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx+1,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
int ib=i-1;
uuu(k,j,i+offx_uuu,icrm) =
pp2(uuu(k,j,i+offx_uuu,icrm))*min(1.0,min(mx(k,j,i+offx_m,icrm), mn(k,j,ib+offx_m,icrm))) -
Expand Down Expand Up @@ -290,7 +290,7 @@ void advect_scalar2D(real5d &f, int ind_f, real2d &flux) {
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx+4; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx+4,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx+4,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
if (i >= 2 && i <= nx+1) {
yakl::atomicAdd(flux(k,icrm),www(k,j,i,icrm));
}
Expand Down Expand Up @@ -362,7 +362,7 @@ void advect_scalar2D(real5d &f, int ind_f, real2d &flux) {
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx+1; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx+1,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx+1,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
int ib=i-1;
uuu(k,j,i+offx_uuu,icrm)= pp2(uuu(k,j,i+offx_uuu,icrm))*min(1.0,min(mx(k,j,i+offx_m,icrm), mn(k,j,ib+offx_m,icrm))) -
pn2(uuu(k,j,i+offx_uuu,icrm))*min(1.0,min(mx(k,j,ib+offx_m,icrm),mn(k,j,i+offx_m,icrm)));
Expand Down Expand Up @@ -486,7 +486,7 @@ void advect_scalar2D(real5d &f, int ind_f, real3d &flux, int ind_flux) {
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx+4; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx+4,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx+4,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
if (i >= 2 && i <= nx+1) {
yakl::atomicAdd(flux(ind_flux,k,icrm),www(k,j,i,icrm));
}
Expand Down Expand Up @@ -557,7 +557,7 @@ void advect_scalar2D(real5d &f, int ind_f, real3d &flux, int ind_flux) {
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx+1; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx+1,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx+1,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
int ib=i-1;
uuu(k,j,i+offx_uuu,icrm)= pp2(uuu(k,j,i+offx_uuu,icrm))*min(1.0,min(mx(k,j,i+offx_m,icrm), mn(k,j,ib+offx_m,icrm))) -
pn2(uuu(k,j,i+offx_uuu,icrm))*min(1.0,min(mx(k,j,ib+offx_m,icrm),mn(k,j,i+offx_m,icrm)));
Expand Down
12 changes: 6 additions & 6 deletions components/eam/src/physics/crm/samxx/advect_scalar3D.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -142,7 +142,7 @@ void advect_scalar3D(real4d &f, real2d &flux) {
// for (int j=0; j<ny+4; j++) {
// for (int i=0; i<nx+4; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny+4,nx+4,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny+4,nx+4,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (i >= 2 && i <= nx+1 && j >= 2 && j <= ny+1) {
yakl::atomicAdd(flux(k,icrm),www(k,j,i,icrm));
}
Expand Down Expand Up @@ -256,7 +256,7 @@ void advect_scalar3D(real4d &f, real2d &flux) {
// for (int j=0; j<ny+1; j++) {
// for (int i=0; i<nx+1; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny+1,nx+1,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny+1,nx+1,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (j <= ny-1) {
int ib=i-1;
uuu(k,j+offy_uuu,i+offx_uuu,icrm) =
Expand Down Expand Up @@ -439,7 +439,7 @@ void advect_scalar3D(real5d &f, int ind_f, real2d &flux) {
// for (int j=0; j<ny+4; j++) {
// for (int i=0; i<nx+4; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny+4,nx+4,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny+4,nx+4,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (i >= 2 && i <= nx+1 && j >= 2 && j <= ny+1) {
yakl::atomicAdd(flux(k,icrm),www(k,j,i,icrm));
}
Expand Down Expand Up @@ -553,7 +553,7 @@ void advect_scalar3D(real5d &f, int ind_f, real2d &flux) {
// for (int j=0; j<ny+1; j++) {
// for (int i=0; i<nx+1; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny+1,nx+1,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny+1,nx+1,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (j <= ny-1) {
int ib=i-1;
uuu(k,j+offy_uuu,i+offx_uuu,icrm) =
Expand Down Expand Up @@ -738,7 +738,7 @@ void advect_scalar3D(real5d &f, int ind_f, real3d &flux, int ind_flux) {
// for (int j=0; j<ny+4; j++) {
// for (int i=0; i<nx+4; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny+4,nx+4,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny+4,nx+4,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (i >= 2 && i <= nx+1 && j >= 2 && j <= ny+1) {
yakl::atomicAdd(flux(ind_flux,k,icrm),www(k,j,i,icrm));
}
Expand Down Expand Up @@ -879,7 +879,7 @@ void advect_scalar3D(real5d &f, int ind_f, real3d &flux, int ind_flux) {
// for (int j=0; j<ny+1; j++) {
// for (int i=0; i<nx+1; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny+1,nx+1,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny+1,nx+1,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
if (j <= ny-1) {
int ib=i-1;
uuu(k,j+offy_uuu,i+offx_uuu,icrm) =
Expand Down
8 changes: 4 additions & 4 deletions components/eam/src/physics/crm/samxx/crm_variance_transport.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,8 +16,8 @@ void VT_filter(int filter_wn_max, real4d &f_in, real4d &f_out) {

yakl::RealFFT1D<nx> fftx;
yakl::RealFFT1D<fftySize> ffty;
fftx.init(fftx.trig);
ffty.init(ffty.trig);
fftx.init();
ffty.init();

//----------------------------------------------------------------------------
// Forward Fourier transform
Expand Down Expand Up @@ -129,7 +129,7 @@ void VT_diagnose() {
// do j = 1,ny
// do i = 1,nx
// do icrm = 1,ncrms
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
yakl::atomicAdd( t_mean(k,icrm) , t(k,j+offy_s,i+offx_s,icrm) );
yakl::atomicAdd( q_mean(k,icrm) , qv(k,j,i,icrm) + qcl(k,j,i,icrm) + qci(k,j,i,icrm) );
});
Expand Down Expand Up @@ -185,7 +185,7 @@ void VT_diagnose() {
// do j = 1,ny
// do i = 1,nx
// do icrm = 1,ncrms
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
yakl::atomicAdd( t_vt(k,icrm) , t_vt_pert(k,j,i,icrm) * t_vt_pert(k,j,i,icrm) );
yakl::atomicAdd( q_vt(k,icrm) , q_vt_pert(k,j,i,icrm) * q_vt_pert(k,j,i,icrm) );
});
Expand Down
2 changes: 1 addition & 1 deletion components/eam/src/physics/crm/samxx/crmsurface.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ void crmsurface(real1d &bflx) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int j, int i, int icrm) {
parallel_for( SimpleBounds<3>(ny,nx,ncrms) , YAKL_LAMBDA (int j, int i, int icrm) {
real tmp2 = (0.5*(u(0,j+offy_u,i+1+offx_u,icrm)+u(0,j+offy_u,i+offx_u,icrm))+ug);
real tmp3 = (0.5*(v(0,j+YES3D+offy_v,i+offx_v,icrm)+v(0,j+offy_v,i+offx_v,icrm))+vg);
real u_h0 = max(1.0,sqrt(tmp2*tmp2+tmp3*tmp3));
Expand Down
4 changes: 2 additions & 2 deletions components/eam/src/physics/crm/samxx/damping.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ void damping() {
});

// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<2>(nzm,ncrms) , YAKL_DEVICE_LAMBDA (int k, int icrm) {
parallel_for( SimpleBounds<2>(nzm,ncrms) , YAKL_LAMBDA (int k, int icrm) {
if(z(nzm-1,icrm)-z(k,icrm) < fractional_damp_depth*z(nzm-1,icrm)) {
do_damping(k,icrm)=1;
} else {
Expand Down Expand Up @@ -75,7 +75,7 @@ void damping() {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp;

tmp = u(k,offy_u+j,offx_u+i,icrm)/( (real) nx * (real) ny );
Expand Down
4 changes: 2 additions & 2 deletions components/eam/src/physics/crm/samxx/diagnose.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,7 @@ void diagnose() {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real coef1 = rho(k,icrm)*dz(icrm)*adz(k,icrm)*dtfactor;
tabs(k,j,i,icrm) = t(k,j+offy_s,i+offx_s,icrm)-gamaz(k,icrm)+ fac_cond *
(qcl(k,j,i,icrm)+qpl(k,j,i,icrm)) + fac_sub *(qci(k,j,i,icrm) + qpi(k,j,i,icrm));
Expand Down Expand Up @@ -115,7 +115,7 @@ void diagnose() {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real coef1 = rho(k,icrm)*dz(icrm)*adz(k,icrm)*dtfactor;
// Saturated water vapor path with respect to water. Can be used
// with water vapor path (= pw) to compute column-average
Expand Down
4 changes: 2 additions & 2 deletions components/eam/src/physics/crm/samxx/diffuse_mom2D.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -76,7 +76,7 @@ void diffuse_mom2D(real5d &tk) {
// for (int k=0; k<nzm-1; k++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm-1,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm-1,nx,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
int kc=k+1;
real rdz=1.0/dz(icrm);
real rdz2 = rdz*rdz * grdf_z(k,icrm);
Expand All @@ -99,7 +99,7 @@ void diffuse_mom2D(real5d &tk) {

// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<2>(nx,ncrms) , YAKL_DEVICE_LAMBDA (int i, int icrm) {
parallel_for( SimpleBounds<2>(nx,ncrms) , YAKL_LAMBDA (int i, int icrm) {
real rdz=1.0/dz(icrm);
real rdz2 = rdz*rdz * grdf_z(nzm-2,icrm);
real tkz=rdz2*grdf_z(nzm-1,icrm)*tk(0,nzm-1,j+offy_d,i+offx_d,icrm);
Expand Down
4 changes: 2 additions & 2 deletions components/eam/src/physics/crm/samxx/diffuse_mom3D.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -121,7 +121,7 @@ void diffuse_mom3D(real5d &tk) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm-1,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm-1,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
int jb=j-1;
int kc=k+1;
int ib=i-1;
Expand Down Expand Up @@ -149,7 +149,7 @@ void diffuse_mom3D(real5d &tk) {
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int j, int i, int icrm) {
parallel_for( SimpleBounds<3>(ny,nx,ncrms) , YAKL_LAMBDA (int j, int i, int icrm) {
real rdz=1.0/dz(icrm);
real rdz2 = rdz*rdz * grdf_z(nzm-2,icrm);
real tkz=rdz2*grdf_z(nzm-1,icrm)*tk(0,nzm-1,j+offy_d,i+offx_d,icrm);
Expand Down
6 changes: 3 additions & 3 deletions components/eam/src/physics/crm/samxx/diffuse_scalar.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ void diffuse_scalar(real5d &tkh, int ind_tkh, real4d &f, real3d &fluxb, real3d &
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp = f(k,j+offy_s,i+offx_s,icrm)-df(k,j+offy_s,i+offx_s,icrm);
yakl::atomicAdd(fdiff(k,icrm),tmp);
});
Expand Down Expand Up @@ -64,7 +64,7 @@ void diffuse_scalar(real5d &tkh, int ind_tkh, real5d &f, int ind_f, real3d &flux
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp = f(ind_f,k,j+offy_s,i+offx_s,icrm)-df(k,j+offy_s,i+offx_s,icrm);
yakl::atomicAdd(fdiff(k,icrm),tmp);
});
Expand Down Expand Up @@ -99,7 +99,7 @@ void diffuse_scalar(real5d &tkh, int ind_tkh, real5d &f, int ind_f, real4d &flux
// for (int j=0; j<ny; j++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int j, int i, int icrm) {
parallel_for( SimpleBounds<4>(nzm,ny,nx,ncrms) , YAKL_LAMBDA (int k, int j, int i, int icrm) {
real tmp = f(ind_f,k,j+offy_s,i+offx_s,icrm)-df(k,j+offy_s,i+offx_s,icrm);
yakl::atomicAdd(fdiff(ind_fdiff,k,icrm),tmp);
});
Expand Down
6 changes: 3 additions & 3 deletions components/eam/src/physics/crm/samxx/diffuse_scalar2D.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ void diffuse_scalar2D(real4d &field, real3d &fluxb, real3d &fluxt, real5d &tkh,
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
if (k <= nzm-2) {
int kc=k+1;
real rhoi = rhow(kc,icrm)/adzw(kc,icrm);
Expand Down Expand Up @@ -148,7 +148,7 @@ void diffuse_scalar2D(real5d &field, int ind_field, real3d &fluxb, real3d &fluxt
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
if (k <= nzm-2) {
int kc=k+1;
real rhoi = rhow(kc,icrm)/adzw(kc,icrm);
Expand Down Expand Up @@ -237,7 +237,7 @@ void diffuse_scalar2D(real5d &field, int ind_field, real4d &fluxb, int ind_fluxb
// for (int k=0; k<nzm; k++) {
// for (int i=0; i<nx; i++) {
// for (int icrm=0; icrm<ncrms; icrm++) {
parallel_for( SimpleBounds<3>(nzm,nx,ncrms) , YAKL_DEVICE_LAMBDA (int k, int i, int icrm) {
parallel_for( SimpleBounds<3>(nzm,nx,ncrms) , YAKL_LAMBDA (int k, int i, int icrm) {
if (k <= nzm-2) {
int kc=k+1;
real rhoi = rhow(kc,icrm)/adzw(kc,icrm);
Expand Down
Loading