From 377cd52b2e98efa558237b51a22ffbebda5ccff6 Mon Sep 17 00:00:00 2001 From: sjplimp Date: Mon, 5 Mar 2012 15:28:58 +0000 Subject: [PATCH] git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@7903 f3b2605a-c512-4ea7-a41b-209d697bcdaa --- lib/cuda/cuda_pair.cu | 9 +++++---- lib/cuda/cuda_pair_kernel.cu | 11 ++++------- lib/cuda/pair_sw_cuda_kernel_nc.cu | 19 ++++++++++--------- lib/cuda/pair_tersoff_cuda_kernel_nc.cu | 19 ++++++++++--------- 4 files changed, 29 insertions(+), 29 deletions(-) diff --git a/lib/cuda/cuda_pair.cu b/lib/cuda/cuda_pair.cu index f84e6129af..0685ce4490 100644 --- a/lib/cuda/cuda_pair.cu +++ b/lib/cuda/cuda_pair.cu @@ -794,6 +794,7 @@ void Cuda_Pair_PreKernel_AllStyles(cuda_shared_data* sdata, cuda_shared_neighlis grid.x = layout.x; grid.y = layout.y; grid.z = 1; int size=(unsigned)(layout.y*layout.x)*sharedperproc*sizeof(ENERGY_FLOAT); + if(sdata->pair.collect_forces_later) size+=(unsigned)(sdata->atom.nmax*3*sizeof(F_FLOAT)); Cuda_UpdateBuffer(sdata,size); if(sdata->pair.use_block_per_atom) @@ -843,10 +844,10 @@ void Cuda_Pair_PostKernel_AllStyles(cuda_shared_data* sdata, dim3& grid, int& sh #include "pair_buck_coul_cut_cuda.cu" #include "pair_buck_coul_long_cuda.cu" #include "pair_buck_cuda.cu" -#include "pair_cg_cmm_cuda.cu" -#include "pair_cg_cmm_coul_cut_cuda.cu" -#include "pair_cg_cmm_coul_debye_cuda.cu" -#include "pair_cg_cmm_coul_long_cuda.cu" +#include "pair_lj_sdk_cuda.cu" +#include "pair_lj_sdk_coul_cut_cuda.cu" +#include "pair_lj_sdk_coul_debye_cuda.cu" +#include "pair_lj_sdk_coul_long_cuda.cu" #include "pair_gran_hooke_cuda.cu" #include "pair_lj_charmm_coul_charmm_implicit_cuda.cu" #include "pair_lj_charmm_coul_charmm_cuda.cu" diff --git a/lib/cuda/cuda_pair_kernel.cu b/lib/cuda/cuda_pair_kernel.cu index 35a0ef1f1a..72c0e0aa25 100644 --- a/lib/cuda/cuda_pair_kernel.cu +++ b/lib/cuda/cuda_pair_kernel.cu @@ -28,9 +28,6 @@ #define A4 -1.453152027 #define A5 1.061405429 -inline __device__ int sbmask(int j) { - return j >> SBBITS & 3; -} template __global__ void Pair_Kernel_TpA(int eflag, int vflag,int eflag_atom,int vflag_atom) @@ -132,7 +129,7 @@ __global__ void Pair_Kernel_TpA(int eflag, int vflag,int eflag_atom,int vflag_at fpair += PairBuckCuda_Eval(rsq,itype * _cuda_ntypes + jtype,factor_lj,eflag,evdwl); break; case PAIR_CG_CMM: - fpair += PairCGCMMCuda_Eval(rsq,itype * _cuda_ntypes + jtype,factor_lj,eflag,evdwl); + fpair += PairLJSDKCuda_Eval(rsq,itype * _cuda_ntypes + jtype,factor_lj,eflag,evdwl); break; case PAIR_LJ_CHARMM: fpair += PairLJCharmmCuda_Eval(rsq,itype * _cuda_ntypes + jtype,factor_lj,eflag,evdwl); @@ -417,7 +414,7 @@ template (); + if(vflagm&&eflag) PairVirialCompute_A_Kernel_Template<1,1>(); + else if(eflag) PairVirialCompute_A_Kernel_Template<1,0>(); + else if(vflagm) PairVirialCompute_A_Kernel_Template<0,1>(); #undef fxtmp #undef fytmp #undef fztmp diff --git a/lib/cuda/pair_tersoff_cuda_kernel_nc.cu b/lib/cuda/pair_tersoff_cuda_kernel_nc.cu index f94f35a587..125c73855a 100644 --- a/lib/cuda/pair_tersoff_cuda_kernel_nc.cu +++ b/lib/cuda/pair_tersoff_cuda_kernel_nc.cu @@ -994,7 +994,6 @@ __global__ void Pair_Tersoff_Kernel_TpA_RIJ()//F_FLOAT4* _glob_r_ij,int* _glob_n } } } - __syncthreads(); } __syncthreads(); @@ -1034,19 +1033,21 @@ __global__ void Pair_Tersoff_Kernel_TpA_RIJ()//F_FLOAT4* _glob_r_ij,int* _glob_n } if(eflag_atom && i<_nlocal) { - _eatom[i] += evdwl; + _eatom[i] = ENERGY_F(0.5) * evdwl; } if(vflag_atom && i<_nlocal) { - _vatom[i] += ENERGY_F(0.5) * sharedV[0 * blockDim.x]; - _vatom[i+_nmax] += ENERGY_F(0.5) * sharedV[1 * blockDim.x]; - _vatom[i+2*_nmax] += ENERGY_F(0.5) * sharedV[2 * blockDim.x]; - _vatom[i+3*_nmax] += ENERGY_F(0.5) * sharedV[3 * blockDim.x]; - _vatom[i+4*_nmax] += ENERGY_F(0.5) * sharedV[4 * blockDim.x]; - _vatom[i+5*_nmax] += ENERGY_F(0.5) * sharedV[5 * blockDim.x]; + _vatom[i] = ENERGY_F(0.5) * sharedV[0 * blockDim.x]; + _vatom[i+_nmax] = ENERGY_F(0.5) * sharedV[1 * blockDim.x]; + _vatom[i+2*_nmax] = ENERGY_F(0.5) * sharedV[2 * blockDim.x]; + _vatom[i+3*_nmax] = ENERGY_F(0.5) * sharedV[3 * blockDim.x]; + _vatom[i+4*_nmax] = ENERGY_F(0.5) * sharedV[4 * blockDim.x]; + _vatom[i+5*_nmax] = ENERGY_F(0.5) * sharedV[5 * blockDim.x]; } - if(vflagm||eflag) PairVirialCompute_A_Kernel_Template<1,1>(); + if(vflagm&&eflag) PairVirialCompute_A_Kernel_Template<1,1>(); + else if(eflag) PairVirialCompute_A_Kernel_Template<1,0>(); + else if(vflagm) PairVirialCompute_A_Kernel_Template<0,1>(); #undef fxtmp #undef fytmp #undef fztmp