diff --git a/doc/src/pair_sph_heatconduction.rst b/doc/src/pair_sph_heatconduction.rst index 4716ed54fb..e9004cb5a4 100644 --- a/doc/src/pair_sph_heatconduction.rst +++ b/doc/src/pair_sph_heatconduction.rst @@ -1,8 +1,11 @@ .. index:: pair_style sph/heatconduction +.. index:: pair_style sph/heatconduction/gpu pair_style sph/heatconduction command ===================================== +Accelerator Variants: *sph/heatconduction/gpu* + Syntax """""" diff --git a/lib/gpu/lal_sph_heatconduction.cu b/lib/gpu/lal_sph_heatconduction.cu index c88853e2cf..4cea1433e3 100644 --- a/lib/gpu/lal_sph_heatconduction.cu +++ b/lib/gpu/lal_sph_heatconduction.cu @@ -82,7 +82,6 @@ __kernel void k_sph_heatconduction(const __global numtyp4 *restrict x_, numtyp4 ix; fetch4(ix,i,pos_tex); //x_[i]; int itype=ix.w; numtyp mass_itype = mass[itype]; - numtyp4 iv; fetch4(iv,i,vel_tex); //v_[i]; const numtyp4 extrai = extra[i]; 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]; int jtype=jx.w; - numtyp4 jv; fetch4(jv,j,vel_tex); //v_[j]; // Compute r12 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 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 ihcub = ihsq * ih; numtyp wfd = h - ucl_sqrt(rsq); if (dimension == 3) { @@ -141,7 +138,7 @@ __kernel void k_sph_heatconduction(const __global numtyp4 *restrict x_, } // for nbor } // 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_, @@ -172,7 +169,7 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_, } __syncthreads(); #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 cutsq_p=coeff_in[ONETYPE].z; // cutsq[itype][jtype] #endif @@ -193,8 +190,6 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_, #ifndef ONETYPE int itype=fast_mul((int)MAX_SHARED_TYPES,iw); #endif - numtyp4 iv; fetch4(iv,i,vel_tex); //v_[i]; - int itag=iv.w; const numtyp4 extrai = extra[i]; numtyp rhoi = extrai.x; @@ -214,8 +209,6 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_, int mtype=itype+jx.w; const numtyp cutsq_p=cutsq[mtype]; #endif - numtyp4 jv; fetch4(jv,j,vel_tex); //v_[j]; - int jtag=jv.w; // Compute r12 numtyp delx = ix.x-jx.x; @@ -226,7 +219,7 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_, if (rsqgpu_rank(); 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; if (SPHHeatConductionMF.device->replica_me()==0 && screen) diff --git a/src/GPU/pair_sph_heatconduction_gpu.cpp b/src/GPU/pair_sph_heatconduction_gpu.cpp index c14a6df937..0f0aa079c8 100644 --- a/src/GPU/pair_sph_heatconduction_gpu.cpp +++ b/src/GPU/pair_sph_heatconduction_gpu.cpp @@ -94,7 +94,6 @@ void PairSPHHeatConductionGPU::compute(int eflag, int vflag) double *rho = atom->rho; double *esph = atom->esph; - double *cv = atom->cv; sph_heatconduction_gpu_get_extra_data(rho, esph); if (gpu_mode != GPU_FORCE) { @@ -140,7 +139,7 @@ void PairSPHHeatConductionGPU::compute(int eflag, int vflag) } } else { - auto dE_ptr = (float *)dE_pinned; + auto dE_ptr = (double *)dE_pinned; for (int i = 0; i < nlocal; i++) { desph[i] = dE_ptr[i]; }