ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fixed bugs with acquiring depsh from lib/gpu, updated the doc page

Browse Source
This commit is contained in:
Trung Nguyen
2023-12-09 15:39:54 -06:00
parent 3830711dec
commit 267e360bac
4 changed files with 10 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
doc/src/pair_sph_heatconduction.rst
View File

@ -1,8 +1,11 @@
.. index:: pair_style sph/heatconduction .. index:: pair_style sph/heatconduction
.. index:: pair_style sph/heatconduction/gpu
pair_style sph/heatconduction command pair_style sph/heatconduction command
===================================== =====================================
Accelerator Variants: *sph/heatconduction/gpu*
Syntax Syntax
"""""" """"""

18
lib/gpu/lal_sph_heatconduction.cu
View File

@ -82,7 +82,6 @@ __kernel void k_sph_heatconduction(const __global numtyp4 *restrict x_,
numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i]; numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i];
int itype=ix.w; int itype=ix.w;
numtyp mass_itype = mass[itype]; numtyp mass_itype = mass[itype];
numtyp4 iv; fetch4(iv,i,vel_tex); //v_[i];
const numtyp4 extrai = extra[i]; const numtyp4 extrai = extra[i];
numtyp rhoi = extrai.x; numtyp rhoi = extrai.x;
@ -96,7 +95,6 @@ __kernel void k_sph_heatconduction(const __global numtyp4 *restrict x_,
numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j]; numtyp4 jx; fetch4(jx,j,pos_tex); //x_[j];
int jtype=jx.w; int jtype=jx.w;
numtyp4 jv; fetch4(jv,j,vel_tex); //v_[j];
// Compute r12 // Compute r12
numtyp delx = ix.x-jx.x; numtyp delx = ix.x-jx.x;
@ -115,9 +113,8 @@ __kernel void k_sph_heatconduction(const __global numtyp4 *restrict x_,
numtyp esphj = extraj.y; numtyp esphj = extraj.y;
numtyp h = coeffy; // cut[itype][jtype] numtyp h = coeffy; // cut[itype][jtype]
ih = ucl_recip(h); // (numtyp)1.0 / h; numtyp ih = ucl_recip(h); // (numtyp)1.0 / h;
numtyp ihsq = ih * ih; numtyp ihsq = ih * ih;
numtyp ihcub = ihsq * ih;
numtyp wfd = h - ucl_sqrt(rsq); numtyp wfd = h - ucl_sqrt(rsq);
if (dimension == 3) { if (dimension == 3) {
@ -141,7 +138,7 @@ __kernel void k_sph_heatconduction(const __global numtyp4 *restrict x_,
} // for nbor } // for nbor
} // if ii } // if ii
store_drhoE(dEacc,ii,inum,tid,t_per_atom,offset,drhoE); store_drhoE(dEacc,ii,inum,tid,t_per_atom,offset,dE);
} }
__kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_, __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
@ -172,7 +169,7 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
} }
__syncthreads(); __syncthreads();
#else #else
const numtyp coeffx=coeff_in[ONETYPE].x; // viscosity[itype][jtype] const numtyp coeffx=coeff_in[ONETYPE].x; // alpha[itype][jtype]
const numtyp coeffy=coeff_in[ONETYPE].y; // cut[itype][jtype] const numtyp coeffy=coeff_in[ONETYPE].y; // cut[itype][jtype]
const numtyp cutsq_p=coeff_in[ONETYPE].z; // cutsq[itype][jtype] const numtyp cutsq_p=coeff_in[ONETYPE].z; // cutsq[itype][jtype]
#endif #endif
@ -193,8 +190,6 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
#ifndef ONETYPE #ifndef ONETYPE
int itype=fast_mul((int)MAX_SHARED_TYPES,iw); int itype=fast_mul((int)MAX_SHARED_TYPES,iw);
#endif #endif
numtyp4 iv; fetch4(iv,i,vel_tex); //v_[i];
int itag=iv.w;
const numtyp4 extrai = extra[i]; const numtyp4 extrai = extra[i];
numtyp rhoi = extrai.x; numtyp rhoi = extrai.x;
@ -214,8 +209,6 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
int mtype=itype+jx.w; int mtype=itype+jx.w;
const numtyp cutsq_p=cutsq[mtype]; const numtyp cutsq_p=cutsq[mtype];
#endif #endif
numtyp4 jv; fetch4(jv,j,vel_tex); //v_[j];
int jtag=jv.w;
// Compute r12 // Compute r12
numtyp delx = ix.x-jx.x; numtyp delx = ix.x-jx.x;
@ -226,7 +219,7 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
if (rsq<cutsq_p) { if (rsq<cutsq_p) {
numtyp mass_jtype = mass[jtype]; numtyp mass_jtype = mass[jtype];
#ifndef ONETYPE #ifndef ONETYPE
const numtyp coeffx=coeff[mtype].x; // viscosity[itype][jtype] const numtyp coeffx=coeff[mtype].x; // alpha[itype][jtype]
const numtyp coeffy=coeff[mtype].y; // cut[itype][jtype] const numtyp coeffy=coeff[mtype].y; // cut[itype][jtype]
#endif #endif
const numtyp4 extraj = extra[j]; const numtyp4 extraj = extra[j];
@ -234,9 +227,8 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
numtyp esphj = extraj.y; numtyp esphj = extraj.y;
numtyp h = coeffy; // cut[itype][jtype] numtyp h = coeffy; // cut[itype][jtype]
ih = ih = ucl_recip(h); // (numtyp)1.0 / h; numtyp ih = ih = ucl_recip(h); // (numtyp)1.0 / h;
numtyp ihsq = ih * ih; numtyp ihsq = ih * ih;
numtyp ihcub = ihsq * ih;
numtyp wfd = h - ucl_sqrt(rsq); numtyp wfd = h - ucl_sqrt(rsq);
if (dimension == 3) { if (dimension == 3) {

2
lib/gpu/lal_sph_heatconduction_ext.cpp
View File

@ -41,7 +41,7 @@ int sph_heatconduction_gpu_init(const int ntypes, double **cutsq, double** host_
int gpu_rank=SPHHeatConductionMF.device->gpu_rank(); int gpu_rank=SPHHeatConductionMF.device->gpu_rank();
int procs_per_gpu=SPHHeatConductionMF.device->procs_per_gpu(); int procs_per_gpu=SPHHeatConductionMF.device->procs_per_gpu();
SPHHeatConductionMF.device->init_message(screen,"sph_lj",first_gpu,last_gpu); SPHHeatConductionMF.device->init_message(screen,"sph_heatconduction",first_gpu,last_gpu);
bool message=false; bool message=false;
if (SPHHeatConductionMF.device->replica_me()==0 && screen) if (SPHHeatConductionMF.device->replica_me()==0 && screen)

3
src/GPU/pair_sph_heatconduction_gpu.cpp
View File

@ -94,7 +94,6 @@ void PairSPHHeatConductionGPU::compute(int eflag, int vflag)
double *rho = atom->rho; double *rho = atom->rho;
double *esph = atom->esph; double *esph = atom->esph;
double *cv = atom->cv;
sph_heatconduction_gpu_get_extra_data(rho, esph); sph_heatconduction_gpu_get_extra_data(rho, esph);
if (gpu_mode != GPU_FORCE) { if (gpu_mode != GPU_FORCE) {
@ -140,7 +139,7 @@ void PairSPHHeatConductionGPU::compute(int eflag, int vflag)
} }
} else { } else {
auto dE_ptr = (float *)dE_pinned; auto dE_ptr = (double *)dE_pinned;
for (int i = 0; i < nlocal; i++) { for (int i = 0; i < nlocal; i++) {
desph[i] = dE_ptr[i]; desph[i] = dE_ptr[i];
} }