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

git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@12589 f3b2605a-c512-4ea7-a41b-209d697bcdaa

Browse Source
This commit is contained in:
sjplimp
2014-10-06 22:59:18 +00:00
parent 621fa7d600
commit 352e1e9df4
65 changed files with 376 additions and 376 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
src/USER-CUDA/atom_vec_angle_cuda.cpp
View File

@ -56,7 +56,7 @@ using namespace LAMMPS_NS;
#define BUFEXTRA 1000
#define NCUDAEXCHANGE 12 //nextra x y z vx vy vz tag type mask image molecule
#define BUF_FLOAT double
#define BUF_CFLOAT double
/* ---------------------------------------------------------------------- */
AtomVecAngleCuda::AtomVecAngleCuda(LAMMPS *lmp) : AtomVecAngle(lmp)
@ -145,8 +145,8 @@ int AtomVecAngleCuda::pack_comm(int n, int* iswap, double *buf,
return AtomVecAngle::pack_comm(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackComm(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}
@ -157,8 +157,8 @@ int AtomVecAngleCuda::pack_comm_vel(int n, int* iswap, double *buf,
return AtomVecAngle::pack_comm_vel(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackCommVel(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}

10
src/USER-CUDA/atom_vec_atomic_cuda.cpp
View File

@ -56,7 +56,7 @@ using namespace LAMMPS_NS;
#define NCUDAEXCHANGE 11 //nextra x y z vx vy vz tag type mask image
#define BUF_FLOAT double
#define BUF_CFLOAT double
/* ---------------------------------------------------------------------- */
AtomVecAtomicCuda::AtomVecAtomicCuda(LAMMPS *lmp) : AtomVecAtomic(lmp)
@ -141,8 +141,8 @@ int AtomVecAtomicCuda::pack_comm(int n, int* iswap, double *buf,
return AtomVecAtomic::pack_comm(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackComm(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}
@ -153,8 +153,8 @@ int AtomVecAtomicCuda::pack_comm_vel(int n, int* iswap, double *buf,
return AtomVecAtomic::pack_comm_vel(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackCommVel(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}
/* ---------------------------------------------------------------------- */

10
src/USER-CUDA/atom_vec_charge_cuda.cpp
View File

@ -55,7 +55,7 @@ using namespace LAMMPS_NS;
#define BUFEXTRA 1000
#define NCUDAEXCHANGE 12 //nextra x y z vx vy vz tag type mask image q
#define BUF_FLOAT double
#define BUF_CFLOAT double
/* ---------------------------------------------------------------------- */
AtomVecChargeCuda::AtomVecChargeCuda(LAMMPS *lmp) : AtomVecCharge(lmp)
@ -140,8 +140,8 @@ int AtomVecChargeCuda::pack_comm(int n, int* iswap, double *buf,
return AtomVecCharge::pack_comm(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackComm(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}
@ -152,8 +152,8 @@ int AtomVecChargeCuda::pack_comm_vel(int n, int* iswap, double *buf,
return AtomVecCharge::pack_comm_vel(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackCommVel(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}

10
src/USER-CUDA/atom_vec_full_cuda.cpp
View File

@ -56,7 +56,7 @@ using namespace LAMMPS_NS;
#define BUFEXTRA 1000
#define NCUDAEXCHANGE 13 //nextra x y z vx vy vz tag type mask image q molecule
#define BUF_FLOAT double
#define BUF_CFLOAT double
/* ---------------------------------------------------------------------- */
AtomVecFullCuda::AtomVecFullCuda(LAMMPS *lmp) :
@ -146,8 +146,8 @@ int AtomVecFullCuda::pack_comm(int n, int* iswap, double *buf,
return AtomVecFull::pack_comm(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackComm(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}
@ -158,8 +158,8 @@ int AtomVecFullCuda::pack_comm_vel(int n, int* iswap, double *buf,
return AtomVecFull::pack_comm_vel(n,iswap,buf,pbc_flag,pbc);
int m = Cuda_CommCuda_PackCommVel(&cuda->shared_data,n,*iswap,(void*) buf,pbc,pbc_flag);
if((sizeof(X_FLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && m)
m=(m+1)*sizeof(X_CFLOAT)/sizeof(double);
return m;
}

52
src/USER-CUDA/comm_cuda.cpp
View File

@ -121,11 +121,11 @@ void CommCuda::init()
cu_pbc->upload();
delete cu_slablo;
cu_slablo = new cCudaData<double, X_FLOAT,x>(slablo,cuda->shared_data.comm.maxswap);
cu_slablo = new cCudaData<double, X_CFLOAT,x>(slablo,cuda->shared_data.comm.maxswap);
cu_slablo->upload();
delete cu_slabhi;
cu_slabhi = new cCudaData<double, X_FLOAT,x>(slabhi,cuda->shared_data.comm.maxswap);
cu_slabhi = new cCudaData<double, X_CFLOAT,x>(slabhi,cuda->shared_data.comm.maxswap);
cu_slabhi->upload();
cuda->shared_data.comm.pbc.dev_data=cu_pbc->dev_data();
@ -214,8 +214,8 @@ void CommCuda::forward_comm_cuda()
int size_forward_recv_now=0;
if((sizeof(X_FLOAT)!=sizeof(double)) && size_forward_recv[iswap]) //some complicated way to safe some transfer size if single precision is used
size_forward_recv_now=(size_forward_recv[iswap]+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && size_forward_recv[iswap]) //some complicated way to safe some transfer size if single precision is used
size_forward_recv_now=(size_forward_recv[iswap]+1)*sizeof(X_CFLOAT)/sizeof(double);
else
size_forward_recv_now=size_forward_recv[iswap];
my_gettime(CLOCK_REALTIME,&time1);
@ -226,8 +226,8 @@ my_gettime(CLOCK_REALTIME,&time1);
my_gettime(CLOCK_REALTIME,&time2);
if((sizeof(X_FLOAT)!=sizeof(double)) && n) //some complicated way to safe some transfer size if single precision is used
n=(n+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && n) //some complicated way to safe some transfer size if single precision is used
n=(n+1)*sizeof(X_CFLOAT)/sizeof(double);
//printf("RecvSize: %i SendSize: %i\n",size_forward_recv_now,n);
MPI_Send(buf_send,n,MPI_DOUBLE,sendproc[iswap],0,world);
@ -285,8 +285,8 @@ cuda->shared_data.cuda_timings.comm_forward_mpi_lower+=
{
n = Cuda_CommCuda_PackComm_Self(&cuda->shared_data,sendnum[iswap],iswap,firstrecv[iswap],pbc[iswap],pbc_flag[iswap]);
if(n<0) error->all(FLERR," # CUDA ERRROR on PackComm_Self");
if((sizeof(X_FLOAT)!=sizeof(double)) && n)
n=(n+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && n)
n=(n+1)*sizeof(X_CFLOAT)/sizeof(double);
}
}
else if (ghost_velocity)
@ -346,8 +346,8 @@ my_gettime(CLOCK_REALTIME,&time1);
my_gettime(CLOCK_REALTIME,&time2);
if((sizeof(X_FLOAT)!=sizeof(double)) && n) //some complicated way to safe some transfer size if single precision is used
n=(n+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && n) //some complicated way to safe some transfer size if single precision is used
n=(n+1)*sizeof(X_CFLOAT)/sizeof(double);
cuda->shared_data.comm.send_size[iswap]=n;
}
else if (ghost_velocity)
@ -358,8 +358,8 @@ my_gettime(CLOCK_REALTIME,&time1);
my_gettime(CLOCK_REALTIME,&time2);
if((sizeof(X_FLOAT)!=sizeof(double)) && n) //some complicated way to safe some transfer size if single precision is used
n=(n+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && n) //some complicated way to safe some transfer size if single precision is used
n=(n+1)*sizeof(X_CFLOAT)/sizeof(double);
cuda->shared_data.comm.send_size[iswap]=n;
}
else
@ -388,8 +388,8 @@ my_gettime(CLOCK_REALTIME,&time2);
{
n = Cuda_CommCuda_PackComm_Self(&cuda->shared_data,sendnum[iswap],iswap,firstrecv[iswap],pbc[iswap],pbc_flag[iswap]);
if(n<0) error->all(FLERR," # CUDA ERRROR on PackComm_Self");
if((sizeof(X_FLOAT)!=sizeof(double)) && n)
n=(n+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && n)
n=(n+1)*sizeof(X_CFLOAT)/sizeof(double);
}
}
else if (ghost_velocity)
@ -441,8 +441,8 @@ void CommCuda::forward_comm_transfer_cuda()
int size_forward_recv_now=0;
if((sizeof(X_FLOAT)!=sizeof(double)) && size_forward_recv[iswap]) //some complicated way to safe some transfer size if single precision is used
size_forward_recv_now=(size_forward_recv[iswap]+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && size_forward_recv[iswap]) //some complicated way to safe some transfer size if single precision is used
size_forward_recv_now=(size_forward_recv[iswap]+1)*sizeof(X_CFLOAT)/sizeof(double);
else
size_forward_recv_now=size_forward_recv[iswap];
@ -485,8 +485,8 @@ cuda->shared_data.cuda_timings.comm_forward_mpi_upper+=
{
/* int size_forward_recv_now=0;
if((sizeof(X_FLOAT)!=sizeof(double)) && size_forward_recv[iswap]) //some complicated way to safe some transfer size if single precision is used
size_forward_recv_now=(size_forward_recv[iswap]+1)*sizeof(X_FLOAT)/sizeof(double);
if((sizeof(X_CFLOAT)!=sizeof(double)) && size_forward_recv[iswap]) //some complicated way to safe some transfer size if single precision is used
size_forward_recv_now=(size_forward_recv[iswap]+1)*sizeof(X_CFLOAT)/sizeof(double);
else
size_forward_recv_now=size_forward_recv[iswap];
@ -693,8 +693,8 @@ void CommCuda::reverse_comm()
if (comm_f_only) {
int size_recv_now=size_reverse_recv[iswap];
if((sizeof(F_FLOAT)!=sizeof(double))&& size_reverse_recv[iswap])
size_recv_now=(size_recv_now+1)*sizeof(F_FLOAT)/sizeof(double);
if((sizeof(F_CFLOAT)!=sizeof(double))&& size_reverse_recv[iswap])
size_recv_now=(size_recv_now+1)*sizeof(F_CFLOAT)/sizeof(double);
MPI_Irecv(buf_recv,size_recv_now,MPI_DOUBLE,
sendproc[iswap],0,world,&request);
@ -705,8 +705,8 @@ void CommCuda::reverse_comm()
}
else buf=NULL;
int size_reverse_send_now=size_reverse_send[iswap];
if((sizeof(F_FLOAT)!=sizeof(double))&& size_reverse_send[iswap])
size_reverse_send_now=(size_reverse_send_now+1)*sizeof(F_FLOAT)/sizeof(double);
if((sizeof(F_CFLOAT)!=sizeof(double))&& size_reverse_send[iswap])
size_reverse_send_now=(size_reverse_send_now+1)*sizeof(F_CFLOAT)/sizeof(double);
MPI_Send(buf,size_reverse_send_now,MPI_DOUBLE,
recvproc[iswap],0,world);
MPI_Wait(&request,&status);
@ -1371,8 +1371,8 @@ void CommCuda::allocate_swap(int n)
if(cu_sendlist)
{
cu_pbc=new cCudaData<int, int, xy> ((int*)pbc,n,6);
cu_slablo = new cCudaData<double, X_FLOAT,x>(slablo,n);
cu_slabhi = new cCudaData<double, X_FLOAT,x>(slabhi,n);
cu_slablo = new cCudaData<double, X_CFLOAT,x>(slablo,n);
cu_slabhi = new cCudaData<double, X_CFLOAT,x>(slabhi,n);
cuda->shared_data.comm.pbc.dev_data=cu_pbc->dev_data();
cuda->shared_data.comm.slablo.dev_data=cu_slablo->dev_data();
@ -1400,8 +1400,8 @@ void CommCuda::allocate_multi(int n)
delete cu_multilo;
delete cu_multihi;
cu_multilo = new cCudaData<double, X_FLOAT,xy>(slablo,n,atom->ntypes+1);
cu_multihi = new cCudaData<double, X_FLOAT,xy>(slabhi,n,atom->ntypes+1);
cu_multilo = new cCudaData<double, X_CFLOAT,xy>(slablo,n,atom->ntypes+1);
cu_multihi = new cCudaData<double, X_CFLOAT,xy>(slabhi,n,atom->ntypes+1);
cuda->shared_data.comm.multilo.dev_data=cu_multilo->dev_data();
cuda->shared_data.comm.multihi.dev_data=cu_multihi->dev_data();

8
src/USER-CUDA/comm_cuda.h
View File

@ -48,10 +48,10 @@ public:
protected:
class Cuda *cuda;
cCudaData<int, int, xy>* cu_pbc;
cCudaData<double, X_FLOAT, x>* cu_slablo;
cCudaData<double, X_FLOAT, x>* cu_slabhi;
cCudaData<double, X_FLOAT, xy>* cu_multilo;
cCudaData<double, X_FLOAT, xy>* cu_multihi;
cCudaData<double, X_CFLOAT, x>* cu_slablo;
cCudaData<double, X_CFLOAT, x>* cu_slabhi;
cCudaData<double, X_CFLOAT, xy>* cu_multilo;
cCudaData<double, X_CFLOAT, xy>* cu_multihi;
cCudaData<int, int, xy>* cu_sendlist;
virtual void grow_send(int,int); // reallocate send buffer

12
src/USER-CUDA/compute_temp_cuda.cpp
View File

@ -112,10 +112,10 @@ double ComputeTempCuda::compute_scalar()
{
if(cuda->begin_setup)
{
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_FLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_FLOAT, x> (&t_scalar,1);
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_CFLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_CFLOAT, x> (&t_scalar,1);
invoked_scalar = update->ntimestep;
Cuda_ComputeTempCuda_Scalar(&cuda->shared_data,groupbit,(ENERGY_FLOAT*) cu_t_scalar->dev_data());
Cuda_ComputeTempCuda_Scalar(&cuda->shared_data,groupbit,(ENERGY_CFLOAT*) cu_t_scalar->dev_data());
cu_t_scalar->download();
}
else
@ -170,12 +170,12 @@ void ComputeTempCuda::compute_vector()
int i;
if(cuda->begin_setup)
{
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_FLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_FLOAT, x> (&t_scalar,1);
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_CFLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_CFLOAT, x> (&t_scalar,1);
invoked_vector = update->ntimestep;
Cuda_ComputeTempCuda_Vector(&cuda->shared_data,groupbit,(ENERGY_FLOAT*) cu_t_vector->dev_data());
Cuda_ComputeTempCuda_Vector(&cuda->shared_data,groupbit,(ENERGY_CFLOAT*) cu_t_vector->dev_data());
cu_t_vector->download();
}
else

4
src/USER-CUDA/compute_temp_cuda.h
View File

@ -65,8 +65,8 @@ class ComputeTempCuda : public Compute {
void dof_compute();
double t_vector[6];
double t_scalar;
cCudaData<double , ENERGY_FLOAT , x>* cu_t_scalar;
cCudaData<double , ENERGY_FLOAT , x>* cu_t_vector;
cCudaData<double , ENERGY_CFLOAT , x>* cu_t_scalar;
cCudaData<double , ENERGY_CFLOAT , x>* cu_t_vector;
};

14
src/USER-CUDA/compute_temp_partial_cuda.cpp
View File

@ -135,10 +135,10 @@ double ComputeTempPartialCuda::compute_scalar()
{
if(cuda->begin_setup)
{
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_FLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_FLOAT, x> (&t_scalar,1);
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_CFLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_CFLOAT, x> (&t_scalar,1);
invoked_scalar = update->ntimestep;
Cuda_ComputeTempPartialCuda_Scalar(&cuda->shared_data,groupbit,(ENERGY_FLOAT*) cu_t_scalar->dev_data(),xflag,yflag,zflag);
Cuda_ComputeTempPartialCuda_Scalar(&cuda->shared_data,groupbit,(ENERGY_CFLOAT*) cu_t_scalar->dev_data(),xflag,yflag,zflag);
cu_t_scalar->download();
}
else
@ -193,12 +193,12 @@ void ComputeTempPartialCuda::compute_vector()
int i;
if(cuda->begin_setup)
{
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_FLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_FLOAT, x> (&t_scalar,1);
if(not cu_t_vector) cu_t_vector = new cCudaData<double, ENERGY_CFLOAT, x> (t_vector,6);
if(not cu_t_scalar) cu_t_scalar = new cCudaData<double, ENERGY_CFLOAT, x> (&t_scalar,1);
invoked_vector = update->ntimestep;
Cuda_ComputeTempPartialCuda_Vector(&cuda->shared_data,groupbit,(ENERGY_FLOAT*) cu_t_vector->dev_data(),xflag,yflag,zflag);
Cuda_ComputeTempPartialCuda_Vector(&cuda->shared_data,groupbit,(ENERGY_CFLOAT*) cu_t_vector->dev_data(),xflag,yflag,zflag);
cu_t_vector->download();
}
else
@ -269,7 +269,7 @@ void ComputeTempPartialCuda::remove_bias_all()
maxbias = atom->nmax;
memory->create(vbiasall,maxbias,3,"temp/partial:vbiasall");
delete cu_vbiasall;
cu_vbiasall = new cCudaData<double, V_FLOAT, yx> ((double*)vbiasall, atom->nmax, 3);
cu_vbiasall = new cCudaData<double, V_CFLOAT, yx> ((double*)vbiasall, atom->nmax, 3);
}
if(cuda->begin_setup)
{

6
src/USER-CUDA/compute_temp_partial_cuda.h
View File

@ -73,9 +73,9 @@ class ComputeTempPartialCuda : public Compute {
void dof_compute();
double t_vector[6];
double t_scalar;
cCudaData<double , ENERGY_FLOAT , x>* cu_t_scalar;
cCudaData<double , ENERGY_FLOAT , x>* cu_t_vector;
cCudaData<double, V_FLOAT, yx>* cu_vbiasall;
cCudaData<double , ENERGY_CFLOAT , x>* cu_t_scalar;
cCudaData<double , ENERGY_CFLOAT , x>* cu_t_vector;
cCudaData<double, V_CFLOAT, yx>* cu_vbiasall;
};
}

56
src/USER-CUDA/cuda.cpp
View File

@ -63,17 +63,17 @@ Cuda::Cuda(LAMMPS* lmp) : Pointers(lmp)
Cuda_Cuda_GetCompileSettings(&shared_data);
if(shared_data.compile_settings.prec_glob != sizeof(CUDA_FLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: Global Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_glob, sizeof(CUDA_FLOAT) / 4);
if(shared_data.compile_settings.prec_glob != sizeof(CUDA_CFLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: Global Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_glob, sizeof(CUDA_CFLOAT) / 4);
if(shared_data.compile_settings.prec_x != sizeof(X_FLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: X Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_x, sizeof(X_FLOAT) / 4);
if(shared_data.compile_settings.prec_x != sizeof(X_CFLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: X Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_x, sizeof(X_CFLOAT) / 4);
if(shared_data.compile_settings.prec_v != sizeof(V_FLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: V Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_v, sizeof(V_FLOAT) / 4);
if(shared_data.compile_settings.prec_v != sizeof(V_CFLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: V Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_v, sizeof(V_CFLOAT) / 4);
if(shared_data.compile_settings.prec_f != sizeof(F_FLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: F Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_f, sizeof(F_FLOAT) / 4);
if(shared_data.compile_settings.prec_f != sizeof(F_CFLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: F Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_f, sizeof(F_CFLOAT) / 4);
if(shared_data.compile_settings.prec_pppm != sizeof(PPPM_FLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: PPPM Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_pppm, sizeof(PPPM_FLOAT) / 4);
if(shared_data.compile_settings.prec_pppm != sizeof(PPPM_CFLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: PPPM Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_pppm, sizeof(PPPM_CFLOAT) / 4);
if(shared_data.compile_settings.prec_fft != sizeof(FFT_FLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: FFT Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_fft, sizeof(FFT_FLOAT) / 4);
if(shared_data.compile_settings.prec_fft != sizeof(FFT_CFLOAT) / 4) printf("\n\n # CUDA WARNING: Compile Settings of cuda and cpp code differ! \n # CUDA WARNING: FFT Precision: cuda %i cpp %i\n\n", shared_data.compile_settings.prec_fft, sizeof(FFT_CFLOAT) / 4);
#ifdef FFT_CUFFT
@ -382,9 +382,9 @@ void Cuda::allocate()
MYDBG(printf("# CUDA: Cuda::allocate ...\n");)
if(not cu_virial) {
cu_virial = new cCudaData<double, ENERGY_FLOAT, x > (NULL, & shared_data.pair.virial , 6);
cu_eng_vdwl = new cCudaData<double, ENERGY_FLOAT, x > (NULL, & shared_data.pair.eng_vdwl , 1);
cu_eng_coul = new cCudaData<double, ENERGY_FLOAT, x > (NULL, & shared_data.pair.eng_coul , 1);
cu_virial = new cCudaData<double, ENERGY_CFLOAT, x > (NULL, & shared_data.pair.virial , 6);
cu_eng_vdwl = new cCudaData<double, ENERGY_CFLOAT, x > (NULL, & shared_data.pair.eng_vdwl , 1);
cu_eng_coul = new cCudaData<double, ENERGY_CFLOAT, x > (NULL, & shared_data.pair.eng_coul , 1);
cu_extent = new cCudaData<double, double, x> (extent, 6);
shared_data.flag = CudaWrapper_AllocCudaData(sizeof(int));
int size = 2 * CUDA_MAX_DEBUG_SIZE;
@ -464,11 +464,11 @@ void Cuda::checkResize()
// do we have more atoms to upload than currently allocated memory on device? (also true if nothing yet allocated)
if(atom->nmax > cu_atom->nmax || cu_tag == NULL) {
delete cu_x;
cu_x = new cCudaData<double, X_FLOAT, yx> ((double*)atom->x , & cu_atom->x , atom->nmax, 3, 0, true); //cu_x->set_buffer(&(shared_data.buffer),&(shared_data.buffersize),true);
cu_x = new cCudaData<double, X_CFLOAT, yx> ((double*)atom->x , & cu_atom->x , atom->nmax, 3, 0, true); //cu_x->set_buffer(&(shared_data.buffer),&(shared_data.buffersize),true);
delete cu_v;
cu_v = new cCudaData<double, V_FLOAT, yx> ((double*)atom->v, & cu_atom->v , atom->nmax, 3);
cu_v = new cCudaData<double, V_CFLOAT, yx> ((double*)atom->v, & cu_atom->v , atom->nmax, 3);
delete cu_f;
cu_f = new cCudaData<double, F_FLOAT, yx> ((double*)atom->f, & cu_atom->f , atom->nmax, 3, 0, true);
cu_f = new cCudaData<double, F_CFLOAT, yx> ((double*)atom->f, & cu_atom->f , atom->nmax, 3, 0, true);
delete cu_tag;
cu_tag = new cCudaData<int , int , x > (atom->tag , & cu_atom->tag , atom->nmax, 0, true);
delete cu_type;
@ -480,31 +480,31 @@ void Cuda::checkResize()
if(atom->rmass) {
delete cu_rmass;
cu_rmass = new cCudaData<double, V_FLOAT, x > (atom->rmass , & cu_atom->rmass , atom->nmax);
cu_rmass = new cCudaData<double, V_CFLOAT, x > (atom->rmass , & cu_atom->rmass , atom->nmax);
}
if(cu_atom->q_flag) {
delete cu_q;
cu_q = new cCudaData<double, F_FLOAT, x > ((double*)atom->q, & cu_atom->q , atom->nmax, 0 , true);
cu_q = new cCudaData<double, F_CFLOAT, x > ((double*)atom->q, & cu_atom->q , atom->nmax, 0 , true);
}// cu_q->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
if(atom->radius) {
delete cu_radius;
cu_radius = new cCudaData<double, X_FLOAT, x > (atom->radius , & cu_atom->radius , atom->nmax);
cu_radius = new cCudaData<double, X_CFLOAT, x > (atom->radius , & cu_atom->radius , atom->nmax);
delete cu_v_radius;
cu_v_radius = new cCudaData<V_FLOAT, V_FLOAT, x> (v_radius , & cu_atom->v_radius , atom->nmax * 4);
cu_v_radius = new cCudaData<V_CFLOAT, V_CFLOAT, x> (v_radius , & cu_atom->v_radius , atom->nmax * 4);
delete cu_omega_rmass;
cu_omega_rmass = new cCudaData<V_FLOAT, V_FLOAT, x> (omega_rmass , & cu_atom->omega_rmass , atom->nmax * 4);
cu_omega_rmass = new cCudaData<V_CFLOAT, V_CFLOAT, x> (omega_rmass , & cu_atom->omega_rmass , atom->nmax * 4);
}
if(atom->omega) {
delete cu_omega;
cu_omega = new cCudaData<double, V_FLOAT, yx > (((double*) atom->omega) , & cu_atom->omega , atom->nmax, 3);
cu_omega = new cCudaData<double, V_CFLOAT, yx > (((double*) atom->omega) , & cu_atom->omega , atom->nmax, 3);
}
if(atom->torque) {
delete cu_torque;
cu_torque = new cCudaData<double, F_FLOAT, yx > (((double*) atom->torque) , & cu_atom->torque , atom->nmax, 3);
cu_torque = new cCudaData<double, F_CFLOAT, yx > (((double*) atom->torque) , & cu_atom->torque , atom->nmax, 3);
}
if(atom->special) {
@ -530,17 +530,17 @@ void Cuda::checkResize()
cu_atom->nmax = atom->nmax;
delete cu_x_type;
cu_x_type = new cCudaData<X_FLOAT, X_FLOAT, x> (x_type , & cu_atom->x_type , atom->nmax * 4);
cu_x_type = new cCudaData<X_CFLOAT, X_CFLOAT, x> (x_type , & cu_atom->x_type , atom->nmax * 4);
}
if(((cu_xhold == NULL) || (cu_xhold->get_dim()[0] < neighbor->maxhold)) && neighbor->xhold) {
delete cu_xhold;
cu_xhold = new cCudaData<double, X_FLOAT, yx> ((double*)neighbor->xhold, & cu_atom->xhold , neighbor->maxhold, 3);
cu_xhold = new cCudaData<double, X_CFLOAT, yx> ((double*)neighbor->xhold, & cu_atom->xhold , neighbor->maxhold, 3);
shared_data.atom.maxhold = neighbor->maxhold;
}
if(atom->mass && !cu_mass) {
cu_mass = new cCudaData<double, V_FLOAT, x > (atom->mass , & cu_atom->mass , atom->ntypes + 1);
cu_mass = new cCudaData<double, V_CFLOAT, x > (atom->mass , & cu_atom->mass , atom->ntypes + 1);
}
cu_atom->mass_host = atom->mass;
@ -618,11 +618,11 @@ void Cuda::evsetup_eatom_vatom(int eflag_atom, int vflag_atom)
{
if(eflag_atom) {
if(not cu_eatom)
cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > (force->pair->eatom, & (shared_data.atom.eatom) , atom->nmax); // cu_eatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > (force->pair->eatom, & (shared_data.atom.eatom) , atom->nmax); // cu_eatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
if(cu_eatom->get_dim()[0] != atom->nmax) {
//delete cu_eatom;
//cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > (force->pair->eatom, & (shared_data.atom.eatom) , atom->nmax );// cu_eatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
//cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > (force->pair->eatom, & (shared_data.atom.eatom) , atom->nmax );// cu_eatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
shared_data.atom.update_nmax = 2;
}
@ -632,11 +632,11 @@ void Cuda::evsetup_eatom_vatom(int eflag_atom, int vflag_atom)
if(vflag_atom) {
if(not cu_vatom)
cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)force->pair->vatom, & (shared_data.atom.vatom) , atom->nmax , 6);// cu_vatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)force->pair->vatom, & (shared_data.atom.vatom) , atom->nmax , 6);// cu_vatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
if(cu_vatom->get_dim()[0] != atom->nmax) {
//delete cu_vatom;
//cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)force->pair->vatom, & (shared_data.atom.vatom) , atom->nmax ,6 );// cu_vatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
//cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)force->pair->vatom, & (shared_data.atom.vatom) , atom->nmax ,6 );// cu_vatom->set_buffer(&(copy_buffer),&(copy_buffersize),true);}
shared_data.atom.update_nmax = 2;
}
@ -899,11 +899,11 @@ void Cuda::update_xhold(int &maxhold, double* xhold)
if(this->shared_data.atom.maxhold < atom->nmax) {
maxhold = atom->nmax;
delete this->cu_xhold;
this->cu_xhold = new cCudaData<double, X_FLOAT, yx> ((double*)xhold, & this->shared_data.atom.xhold , maxhold, 3);
this->cu_xhold = new cCudaData<double, X_CFLOAT, yx> ((double*)xhold, & this->shared_data.atom.xhold , maxhold, 3);
}
this->shared_data.atom.maxhold = maxhold;
CudaWrapper_CopyData(this->cu_xhold->dev_data(), this->cu_x->dev_data(), 3 * atom->nmax * sizeof(X_FLOAT));
CudaWrapper_CopyData(this->cu_xhold->dev_data(), this->cu_x->dev_data(), 3 * atom->nmax * sizeof(X_CFLOAT));
}
void Cuda::setTimingsZero()

6
src/USER-CUDA/fft3d_cuda.cpp
View File

@ -118,9 +118,9 @@ if(nprocs>1)
{
if(nprocs>1)
{
if(sizeof(FFT_FLOAT)==sizeof(double))cudaMemcpy((void*) (plan->cudata2), (void*) data, plan->cudatasize/2,cudaMemcpyHostToDevice);
if(sizeof(FFT_FLOAT)==sizeof(float)) cudaMemcpy((void*) (plan->cudata2), (void*) data, plan->cudatasize,cudaMemcpyHostToDevice);
initfftdata((double*)plan->cudata2,(FFT_FLOAT*)plan->cudata,plan->nfast,plan->nmid,plan->nslow);
if(sizeof(FFT_CFLOAT)==sizeof(double))cudaMemcpy((void*) (plan->cudata2), (void*) data, plan->cudatasize/2,cudaMemcpyHostToDevice);
if(sizeof(FFT_CFLOAT)==sizeof(float)) cudaMemcpy((void*) (plan->cudata2), (void*) data, plan->cudatasize,cudaMemcpyHostToDevice);
initfftdata((double*)plan->cudata2,(FFT_CFLOAT*)plan->cudata,plan->nfast,plan->nmid,plan->nslow);
}
}
if (flag == -1)

4
src/USER-CUDA/fix_addforce_cuda.cpp
View File

@ -101,7 +101,7 @@ int FixAddForceCuda::setmask()
void FixAddForceCuda::init()
{
if(not cu_foriginal)
cu_foriginal = new cCudaData<double, F_FLOAT, x> (foriginal,4);
cu_foriginal = new cCudaData<double, F_CFLOAT, x> (foriginal,4);
if (strstr(update->integrate_style,"respa"))
nlevels_respa = ((Respa *) update->integrate)->nlevels;
}
@ -144,7 +144,7 @@ void FixAddForceCuda::post_force(int vflag)
MYDBG( printf("# CUDA: FixAddForceCuda::postforce start\n"); )
force_flag = 0;
cu_foriginal->memset_device(0);
Cuda_FixAddForceCuda_PostForce(&cuda->shared_data, groupbit, xvalue, yvalue,zvalue,(F_FLOAT*) cu_foriginal->dev_data());
Cuda_FixAddForceCuda_PostForce(&cuda->shared_data, groupbit, xvalue, yvalue,zvalue,(F_CFLOAT*) cu_foriginal->dev_data());
cu_foriginal->download();
}

2
src/USER-CUDA/fix_addforce_cuda.h
View File

@ -53,7 +53,7 @@ class FixAddForceCuda : public Fix {
int iregion;
double xvalue,yvalue,zvalue;
double foriginal[4],foriginal_all[4];
cCudaData<double , F_FLOAT , x>* cu_foriginal;
cCudaData<double , F_CFLOAT , x>* cu_foriginal;
int force_flag;
int nlevels_respa;
};

4
src/USER-CUDA/fix_aveforce_cuda.cpp
View File

@ -143,7 +143,7 @@ int FixAveForceCuda::setmask()
void FixAveForceCuda::init()
{
if(not cu_foriginal)
cu_foriginal = new cCudaData<double, F_FLOAT, x> (foriginal,4);
cu_foriginal = new cCudaData<double, F_CFLOAT, x> (foriginal,4);
if (xstr) {
xvar = input->variable->find(xstr);
@ -207,7 +207,7 @@ void FixAveForceCuda::post_force(int vflag)
// sum forces on participating atoms
cu_foriginal->memset_device(0);
Cuda_FixAveForceCuda_PostForce_FOrg(&cuda->shared_data, groupbit,(F_FLOAT*) cu_foriginal->dev_data());
Cuda_FixAveForceCuda_PostForce_FOrg(&cuda->shared_data, groupbit,(F_CFLOAT*) cu_foriginal->dev_data());
cu_foriginal->download();
// average the force on participating atoms

2
src/USER-CUDA/fix_aveforce_cuda.h
View File

@ -56,7 +56,7 @@ class FixAveForceCuda : public Fix {
double xvalue,yvalue,zvalue;
double foriginal_all[4];
double foriginal[4];
cCudaData<double , F_FLOAT , x>* cu_foriginal;
cCudaData<double , F_CFLOAT , x>* cu_foriginal;
int nlevels_respa;
int varflag;
int xvar,yvar,zvar,xstyle,ystyle,zstyle;

4
src/USER-CUDA/fix_freeze_cuda.cpp
View File

@ -79,7 +79,7 @@ int FixFreezeCuda::setmask()
void FixFreezeCuda::init()
{
if(not cu_foriginal)
cu_foriginal = new cCudaData<double, F_FLOAT, x> (foriginal,3);
cu_foriginal = new cCudaData<double, F_CFLOAT, x> (foriginal,3);
int count = 0;
for (int i = 0; i < modify->nfix; i++)
if (strcmp(modify->fix[i]->style,"freeze") == 0) count++;
@ -113,7 +113,7 @@ void FixFreezeCuda::post_force(int vflag)
MYDBG( printf("# CUDA: FixFreezeCuda::postforce start\n"); )
force_flag = 0;
cu_foriginal->memset_device(0);
Cuda_FixFreezeCuda_PostForce(&cuda->shared_data, groupbit, (F_FLOAT*) cu_foriginal->dev_data());
Cuda_FixFreezeCuda_PostForce(&cuda->shared_data, groupbit, (F_CFLOAT*) cu_foriginal->dev_data());
cu_foriginal->download();
}

2
src/USER-CUDA/fix_freeze_cuda.h
View File

@ -47,7 +47,7 @@ class FixFreezeCuda : public Fix {
private:
class Cuda *cuda;
double foriginal[3],foriginal_all[3];
cCudaData<double , F_FLOAT , x>* cu_foriginal;
cCudaData<double , F_CFLOAT , x>* cu_foriginal;
int force_flag;
};

2
src/USER-CUDA/fix_nh_cuda.h
View File

@ -118,7 +118,7 @@ class FixNHCuda : public Fix {
void compute_press_target();
void nh_omega_dot();
X_FLOAT triggerneighsq;
X_CFLOAT triggerneighsq;
};
}

2
src/USER-CUDA/fix_nve_cuda.h
View File

@ -47,7 +47,7 @@ class FixNVECuda : public Fix
void final_integrate_respa(int, int);
void reset_dt();
X_FLOAT triggerneighsq;
X_CFLOAT triggerneighsq;
protected:
class Cuda *cuda;

4
src/USER-CUDA/fix_set_force_cuda.cpp
View File

@ -84,7 +84,7 @@ int FixSetForceCuda::setmask()
void FixSetForceCuda::init()
{
if(not cu_foriginal)
cu_foriginal = new cCudaData<double, F_FLOAT, x> (foriginal,3);
cu_foriginal = new cCudaData<double, F_CFLOAT, x> (foriginal,3);
if (strstr(update->integrate_style,"respa"))
nlevels_respa = ((Respa *) update->integrate)->nlevels;
}
@ -127,7 +127,7 @@ void FixSetForceCuda::post_force(int vflag)
MYDBG( printf("# CUDA: FixSetForceCuda::postforce start\n"); )
force_flag = 0;
cu_foriginal->memset_device(0);
Cuda_FixSetForceCuda_PostForce(&cuda->shared_data, groupbit, xvalue, yvalue,zvalue,(F_FLOAT*) cu_foriginal->dev_data(),flagx,flagy,flagz);
Cuda_FixSetForceCuda_PostForce(&cuda->shared_data, groupbit, xvalue, yvalue,zvalue,(F_CFLOAT*) cu_foriginal->dev_data(),flagx,flagy,flagz);
cu_foriginal->download();
}

2
src/USER-CUDA/fix_set_force_cuda.h
View File

@ -52,7 +52,7 @@ class FixSetForceCuda : public Fix {
int flagx,flagy,flagz;
double xvalue,yvalue,zvalue;
double foriginal[3],foriginal_all[3];
cCudaData<double , F_FLOAT , x>* cu_foriginal;
cCudaData<double , F_CFLOAT , x>* cu_foriginal;
int force_flag;
int nlevels_respa;
};

8
src/USER-CUDA/fix_shake_cuda.cpp
View File

@ -82,7 +82,7 @@ FixShakeCuda::FixShakeCuda(LAMMPS* lmp, int narg, char** arg) :
cu_list = NULL;
cu_bond_distance = NULL;
cu_angle_distance = NULL;
cu_virial = new cCudaData<double , ENERGY_FLOAT , xx >(virial, 6);
cu_virial = new cCudaData<double , ENERGY_CFLOAT , xx >(virial, 6);
grow_arrays(atom->nmax);
atom->add_callback(0);
@ -174,8 +174,8 @@ FixShakeCuda::FixShakeCuda(LAMMPS* lmp, int narg, char** arg) :
bond_distance = new double[atom->nbondtypes + 1];
angle_distance = new double[atom->nangletypes + 1];
cu_bond_distance = new cCudaData<double, X_FLOAT, xx> (bond_distance, atom->nbondtypes + 1);
cu_angle_distance = new cCudaData<double, X_FLOAT, xx> (angle_distance, atom->nangletypes + 1);
cu_bond_distance = new cCudaData<double, X_CFLOAT, xx> (bond_distance, atom->nbondtypes + 1);
cu_angle_distance = new cCudaData<double, X_CFLOAT, xx> (angle_distance, atom->nangletypes + 1);
// allocate statistics arrays
@ -2564,7 +2564,7 @@ void FixShakeCuda::grow_arrays(int nmax)
delete cu_shake_type;
cu_shake_type = new cCudaData<int, int, yx> ((int*)shake_type, nmax, 3);
delete cu_xshake;
cu_xshake = new cCudaData<double, X_FLOAT, xy> ((double*)xshake, nmax, 3);
cu_xshake = new cCudaData<double, X_CFLOAT, xy> ((double*)xshake, nmax, 3);
cu_shake_flag->upload();
cu_shake_atom->upload();
cu_shake_type->upload();

8
src/USER-CUDA/fix_shake_cuda.h
View File

@ -66,8 +66,8 @@ class FixShakeCuda : public Fix {
bool neighbor_step; // was neighboring done in this step -> need to run the Cuda_FixShake_Init
double *bond_distance,*angle_distance; // constraint distances
cCudaData<double , X_FLOAT , xx >* cu_bond_distance;
cCudaData<double , X_FLOAT , xx >* cu_angle_distance;
cCudaData<double , X_CFLOAT , xx >* cu_bond_distance;
cCudaData<double , X_CFLOAT , xx >* cu_angle_distance;
int ifix_respa; // rRESPA fix needed by SHAKE
int nlevels_respa; // copies of needed rRESPA variables
@ -93,9 +93,9 @@ class FixShakeCuda : public Fix {
cCudaData<int , int , xx >* cu_shake_flag;
cCudaData<int , int , yx >* cu_shake_atom;
cCudaData<int , int , yx >* cu_shake_type;
cCudaData<double , X_FLOAT , xy >* cu_xshake;
cCudaData<double , X_CFLOAT , xy >* cu_xshake;
cCudaData<int , int , xx >* cu_list;
cCudaData<double , ENERGY_FLOAT , xx >* cu_virial;
cCudaData<double , ENERGY_CFLOAT , xx >* cu_virial;
int* countoccur;
int vflag; // virial flag

2
src/USER-CUDA/fix_viscous_cuda.cpp
View File

@ -73,7 +73,7 @@ int FixViscousCuda::setmask()
void FixViscousCuda::setup(int vflag)
{
if(not cu_gamma)
cu_gamma = new cCudaData<double, F_FLOAT, x> (gamma,atom->ntypes+1);
cu_gamma = new cCudaData<double, F_CFLOAT, x> (gamma,atom->ntypes+1);
Cuda_FixViscousCuda_Init(&cuda->shared_data);
cu_gamma->upload();
// if (strcmp(update->integrate_style,"verlet/cuda") == 0)

2
src/USER-CUDA/fix_viscous_cuda.h
View File

@ -43,7 +43,7 @@ class FixViscousCuda : public FixViscous {
void setup(int);
void min_setup(int);
void post_force(int);
cCudaData<double, F_FLOAT, x>* cu_gamma;
cCudaData<double, F_CFLOAT, x>* cu_gamma;
private:
class Cuda *cuda;

6
src/USER-CUDA/pair_born_coul_long_cuda.cpp
View File

@ -117,7 +117,7 @@ void PairBornCoulLongCuda::compute(int eflag, int vflag)
void PairBornCoulLongCuda::settings(int narg, char **arg)
{
PairBornCoulLong::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
}
/* ---------------------------------------------------------------------- */
@ -175,9 +175,9 @@ void PairBornCoulLongCuda::ev_setup(int eflag, int vflag)
PairBornCoulLong::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

8
src/USER-CUDA/pair_buck_coul_cut_cuda.cpp
View File

@ -109,8 +109,8 @@ void PairBuckCoulCutCuda::compute(int eflag, int vflag)
void PairBuckCoulCutCuda::settings(int narg, char **arg)
{
PairBuckCoulCut::settings(narg, arg);
cuda->shared_data.pair.cut_coul_global = (F_FLOAT) cut_coul_global;
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_CFLOAT) cut_coul_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
}
/* ---------------------------------------------------------------------- */
@ -162,9 +162,9 @@ void PairBuckCoulCutCuda::ev_setup(int eflag, int vflag)
PairBuckCoulCut::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_buck_coul_long_cuda.cpp
View File

@ -115,7 +115,7 @@ void PairBuckCoulLongCuda::compute(int eflag, int vflag)
void PairBuckCoulLongCuda::settings(int narg, char **arg)
{
PairBuckCoulLong::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
}
/* ---------------------------------------------------------------------- */
@ -173,9 +173,9 @@ void PairBuckCoulLongCuda::ev_setup(int eflag, int vflag)
PairBuckCoulLong::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_buck_cuda.cpp
View File

@ -107,7 +107,7 @@ void PairBuckCuda::compute(int eflag, int vflag)
void PairBuckCuda::settings(int narg, char **arg)
{
PairBuck::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -158,9 +158,9 @@ void PairBuckCuda::ev_setup(int eflag, int vflag)
PairBuck::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

22
src/USER-CUDA/pair_eam_cuda.cpp
View File

@ -92,14 +92,14 @@ void PairEAMCuda::allocate()
if(! allocated) PairEAM::allocate();
cuda->shared_data.pair.cutsq = cutsq;
cuda->shared_data.pair.cut_global = (F_FLOAT) cutforcesq;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cutforcesq;
}
/* ---------------------------------------------------------------------- */
void PairEAMCuda::compute(int eflag, int vflag)
{
cuda->shared_data.pair.cut_global = (F_FLOAT) cutforcesq;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cutforcesq;
cuda->shared_data.pair.use_block_per_atom = 0;
cuda->shared_data.pair.collect_forces_later = 0;
@ -111,8 +111,8 @@ void PairEAMCuda::compute(int eflag, int vflag)
memory->create(fp, nmax, "pair:fp");
delete cu_rho;
delete cu_fp;
cu_rho = new cCudaData<double, F_FLOAT, x> (rho, atom->nmax);
cu_fp = new cCudaData<double, F_FLOAT, x> (fp, atom->nmax);
cu_rho = new cCudaData<double, F_CFLOAT, x> (rho, atom->nmax);
cu_fp = new cCudaData<double, F_CFLOAT, x> (fp, atom->nmax);
Cuda_PairEAMCuda_Init(&cuda->shared_data, rdr, rdrho, nfrho, nrhor, nr, nrho, nz2r,
cu_frho_spline->dev_data(), cu_rhor_spline->dev_data(), cu_z2r_spline->dev_data(),
cu_rho->dev_data(), cu_fp->dev_data(), type2frho, type2z2r, type2rhor);
@ -142,7 +142,7 @@ void PairEAMCuda::compute(int eflag, int vflag)
void PairEAMCuda::settings(int narg, char** arg)
{
PairEAM::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cutforcesq;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cutforcesq;
}
/* ---------------------------------------------------------------------- */
@ -171,9 +171,9 @@ void PairEAMCuda::init_style()
delete cu_z2r_spline;
delete cu_frho_spline;
cu_rhor_spline = new cCudaData<double, F_FLOAT, xyz>((double*)rhor_spline, nrhor, nr + 1, EAM_COEFF_LENGTH);
cu_z2r_spline = new cCudaData<double, F_FLOAT, xyz>((double*)z2r_spline, nz2r, nr + 1, EAM_COEFF_LENGTH);
cu_frho_spline = new cCudaData<double, F_FLOAT, xyz>((double*)frho_spline, nfrho, nrho + 1, EAM_COEFF_LENGTH);
cu_rhor_spline = new cCudaData<double, F_CFLOAT, xyz>((double*)rhor_spline, nrhor, nr + 1, EAM_COEFF_LENGTH);
cu_z2r_spline = new cCudaData<double, F_CFLOAT, xyz>((double*)z2r_spline, nz2r, nr + 1, EAM_COEFF_LENGTH);
cu_frho_spline = new cCudaData<double, F_CFLOAT, xyz>((double*)frho_spline, nfrho, nrho + 1, EAM_COEFF_LENGTH);
cu_rhor_spline->upload();
cu_z2r_spline->upload();
@ -236,7 +236,7 @@ int PairEAMCuda::pack_forward_comm(int n, int* iswap, double* buf,
{
Cuda_PairEAMCuda_PackComm(&cuda->shared_data, n, *iswap, buf);
if(sizeof(F_FLOAT) < sizeof(double)) return n;
if(sizeof(F_CFLOAT) < sizeof(double)) return n;
else return n;
}
@ -254,12 +254,12 @@ void PairEAMCuda::ev_setup(int eflag, int vflag)
if(eflag_atom && atom->nmax > maxeatomold) {
delete cuda->cu_eatom;
cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax);
cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax);
}
if(vflag_atom && atom->nmax > maxeatomold) {
delete cuda->cu_vatom;
cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6);
cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6);
}
}

10
src/USER-CUDA/pair_eam_cuda.h
View File

@ -66,11 +66,11 @@ class PairEAMCuda : public PairEAM
bool allocated2;
virtual void ev_setup(int eflag, int vflag);
class CudaNeighList* cuda_neigh_list;
cCudaData<double, F_FLOAT, x>* cu_rho;
cCudaData<double, F_FLOAT, x>* cu_fp;
cCudaData<double, F_FLOAT, xyz>* cu_rhor_spline;
cCudaData<double, F_FLOAT, xyz>* cu_z2r_spline;
cCudaData<double, F_FLOAT, xyz>* cu_frho_spline;
cCudaData<double, F_CFLOAT, x>* cu_rho;
cCudaData<double, F_CFLOAT, x>* cu_fp;
cCudaData<double, F_CFLOAT, xyz>* cu_rhor_spline;
cCudaData<double, F_CFLOAT, xyz>* cu_z2r_spline;
cCudaData<double, F_CFLOAT, xyz>* cu_frho_spline;
};

4
src/USER-CUDA/pair_gran_hooke_cuda.cpp
View File

@ -242,9 +242,9 @@ void PairGranHookeCuda::ev_setup(int eflag, int vflag)
PairGranHooke::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.eatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.eatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj96_cut_cuda.cpp
View File

@ -118,7 +118,7 @@ void PairLJ96CutCuda::compute(int eflag, int vflag)
void PairLJ96CutCuda::settings(int narg, char **arg)
{
PairLJ96Cut::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -171,9 +171,9 @@ void PairLJ96CutCuda::ev_setup(int eflag, int vflag)
PairLJ96Cut::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

18
src/USER-CUDA/pair_lj_charmm_coul_charmm_cuda.cpp
View File

@ -76,10 +76,10 @@ void PairLJCharmmCoulCharmmCuda::allocate()
cuda->shared_data.pair.coeff4 = lj4;
cuda->shared_data.pair.special_lj = force->special_lj;
cuda->shared_data.pair.special_coul = force->special_coul;
cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
}
}
@ -110,9 +110,9 @@ void PairLJCharmmCoulCharmmCuda::compute(int eflag, int vflag)
void PairLJCharmmCoulCharmmCuda::settings(int narg, char **arg)
{
PairLJCharmmCoulCharmm::settings(narg, arg);
cuda->shared_data.pair.cut_global = (X_FLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_FLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_FLOAT) cut_lj_inner;
cuda->shared_data.pair.cut_global = (X_CFLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_CFLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_CFLOAT) cut_lj_inner;
}
/* ---------------------------------------------------------------------- */
@ -180,9 +180,9 @@ void PairLJCharmmCoulCharmmCuda::ev_setup(int eflag, int vflag)
PairLJCharmmCoulCharmm::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

8
src/USER-CUDA/pair_lj_charmm_coul_charmm_cuda.h
View File

@ -50,10 +50,10 @@ class PairLJCharmmCoulCharmmCuda : public PairLJCharmmCoulCharmm
void allocate();
bool allocated2;
class CudaNeighList* cuda_neigh_list;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
};

18
src/USER-CUDA/pair_lj_charmm_coul_charmm_implicit_cuda.cpp
View File

@ -76,10 +76,10 @@ void PairLJCharmmCoulCharmmImplicitCuda::allocate()
cuda->shared_data.pair.coeff4 = lj4;
cuda->shared_data.pair.special_lj = force->special_lj;
cuda->shared_data.pair.special_coul = force->special_coul;
cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
}
}
@ -110,9 +110,9 @@ void PairLJCharmmCoulCharmmImplicitCuda::compute(int eflag, int vflag)
void PairLJCharmmCoulCharmmImplicitCuda::settings(int narg, char **arg)
{
PairLJCharmmCoulCharmmImplicit::settings(narg, arg);
cuda->shared_data.pair.cut_global = (X_FLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_FLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_FLOAT) cut_lj_inner;
cuda->shared_data.pair.cut_global = (X_CFLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_CFLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_CFLOAT) cut_lj_inner;
}
/* ---------------------------------------------------------------------- */
@ -175,9 +175,9 @@ void PairLJCharmmCoulCharmmImplicitCuda::ev_setup(int eflag, int vflag)
PairLJCharmmCoulCharmmImplicit::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

8
src/USER-CUDA/pair_lj_charmm_coul_charmm_implicit_cuda.h
View File

@ -50,10 +50,10 @@ class PairLJCharmmCoulCharmmImplicitCuda : public PairLJCharmmCoulCharmmImplicit
void allocate();
bool allocated2;
class CudaNeighList* cuda_neigh_list;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
};
}

18
src/USER-CUDA/pair_lj_charmm_coul_long_cuda.cpp
View File

@ -85,10 +85,10 @@ void PairLJCharmmCoulLongCuda::allocate()
cuda->shared_data.pair.offset = offset;
cuda->shared_data.pair.special_lj = force->special_lj;
cuda->shared_data.pair.special_coul = force->special_coul;
cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
}
}
@ -119,9 +119,9 @@ void PairLJCharmmCoulLongCuda::compute(int eflag, int vflag)
void PairLJCharmmCoulLongCuda::settings(int narg, char **arg)
{
PairLJCharmmCoulLong::settings(narg, arg);
cuda->shared_data.pair.cut_global = (X_FLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_FLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_FLOAT) cut_lj_inner;
cuda->shared_data.pair.cut_global = (X_CFLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_CFLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_CFLOAT) cut_lj_inner;
}
/* ---------------------------------------------------------------------- */
@ -188,9 +188,9 @@ void PairLJCharmmCoulLongCuda::ev_setup(int eflag, int vflag)
PairLJCharmmCoulLong::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

8
src/USER-CUDA/pair_lj_charmm_coul_long_cuda.h
View File

@ -50,10 +50,10 @@ class PairLJCharmmCoulLongCuda : public PairLJCharmmCoulLong
void allocate();
bool allocated2;
class CudaNeighList* cuda_neigh_list;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
};
}

8
src/USER-CUDA/pair_lj_class2_coul_cut_cuda.cpp
View File

@ -105,8 +105,8 @@ void PairLJClass2CoulCutCuda::compute(int eflag, int vflag)
void PairLJClass2CoulCutCuda::settings(int narg, char **arg)
{
PairLJClass2CoulCut::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_FLOAT) cut_coul_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_CFLOAT) cut_coul_global;
}
/* ---------------------------------------------------------------------- */
@ -154,9 +154,9 @@ void PairLJClass2CoulCutCuda::ev_setup(int eflag, int vflag)
PairLJClass2CoulCut::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj_class2_coul_long_cuda.cpp
View File

@ -111,7 +111,7 @@ void PairLJClass2CoulLongCuda::compute(int eflag, int vflag)
void PairLJClass2CoulLongCuda::settings(int narg, char **arg)
{
PairLJClass2CoulLong::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
}
/* ---------------------------------------------------------------------- */
@ -167,9 +167,9 @@ void PairLJClass2CoulLongCuda::ev_setup(int eflag, int vflag)
PairLJClass2CoulLong::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj_class2_cuda.cpp
View File

@ -118,7 +118,7 @@ void PairLJClass2Cuda::compute(int eflag, int vflag)
void PairLJClass2Cuda::settings(int narg, char **arg)
{
PairLJClass2::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -159,9 +159,9 @@ void PairLJClass2Cuda::ev_setup(int eflag, int vflag)
PairLJClass2::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

8
src/USER-CUDA/pair_lj_cut_coul_cut_cuda.cpp
View File

@ -105,8 +105,8 @@ void PairLJCutCoulCutCuda::compute(int eflag, int vflag)
void PairLJCutCoulCutCuda::settings(int narg, char **arg)
{
PairLJCutCoulCut::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_FLOAT) cut_coul_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_CFLOAT) cut_coul_global;
}
/* ---------------------------------------------------------------------- */
@ -154,9 +154,9 @@ void PairLJCutCoulCutCuda::ev_setup(int eflag, int vflag)
PairLJCutCoulCut::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

10
src/USER-CUDA/pair_lj_cut_coul_debye_cuda.cpp
View File

@ -105,9 +105,9 @@ void PairLJCutCoulDebyeCuda::compute(int eflag, int vflag)
void PairLJCutCoulDebyeCuda::settings(int narg, char **arg)
{
PairLJCutCoulDebye::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_FLOAT) cut_coul_global;
cuda->shared_data.pair.kappa = (F_FLOAT) kappa;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_CFLOAT) cut_coul_global;
cuda->shared_data.pair.kappa = (F_CFLOAT) kappa;
}
/* ---------------------------------------------------------------------- */
@ -155,9 +155,9 @@ void PairLJCutCoulDebyeCuda::ev_setup(int eflag, int vflag)
PairLJCutCoulDebye::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj_cut_coul_long_cuda.cpp
View File

@ -111,7 +111,7 @@ void PairLJCutCoulLongCuda::compute(int eflag, int vflag)
void PairLJCutCoulLongCuda::settings(int narg, char **arg)
{
PairLJCutCoulLong::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
}
/* ---------------------------------------------------------------------- */
@ -208,9 +208,9 @@ void PairLJCutCoulLongCuda::ev_setup(int eflag, int vflag)
PairLJCutCoulLong::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj_cut_cuda.cpp
View File

@ -118,7 +118,7 @@ void PairLJCutCuda::compute(int eflag, int vflag)
void PairLJCutCuda::settings(int narg, char **arg)
{
PairLJCut::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -171,9 +171,9 @@ void PairLJCutCuda::ev_setup(int eflag, int vflag)
PairLJCut::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj_cut_experimental_cuda.cpp
View File

@ -117,7 +117,7 @@ void PairLJCutExperimentalCuda::compute(int eflag, int vflag)
void PairLJCutExperimentalCuda::settings(int narg, char **arg)
{
PairLJCut::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -170,9 +170,9 @@ void PairLJCutExperimentalCuda::ev_setup(int eflag, int vflag)
PairLJCut::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

6
src/USER-CUDA/pair_lj_expand_cuda.cpp
View File

@ -120,7 +120,7 @@ void PairLJExpandCuda::compute(int eflag, int vflag)
void PairLJExpandCuda::settings(int narg, char **arg)
{
PairLJExpand::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -172,9 +172,9 @@ void PairLJExpandCuda::ev_setup(int eflag, int vflag)
PairLJExpand::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

28
src/USER-CUDA/pair_lj_gromacs_coul_gromacs_cuda.cpp
View File

@ -81,15 +81,15 @@ void PairLJGromacsCoulGromacsCuda::allocate()
cuda->shared_data.pair.coeff9 = ljsw5;
cuda->shared_data.pair.special_lj = force->special_lj;
cuda->shared_data.pair.special_coul = force->special_coul;
cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw1_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw1, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw2_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw2, &cuda->shared_data.pair.coeff6_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw3_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw3, &cuda->shared_data.pair.coeff7_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw4_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw4, &cuda->shared_data.pair.coeff8_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw5_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw5, &cuda->shared_data.pair.coeff9_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw1, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw2, &cuda->shared_data.pair.coeff6_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw3, &cuda->shared_data.pair.coeff7_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw4, &cuda->shared_data.pair.coeff8_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw5_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw5, &cuda->shared_data.pair.coeff9_gm, (atom->ntypes+1)*(atom->ntypes+1));
}
}
@ -120,9 +120,9 @@ void PairLJGromacsCoulGromacsCuda::compute(int eflag, int vflag)
void PairLJGromacsCoulGromacsCuda::settings(int narg, char **arg)
{
PairLJGromacsCoulGromacs::settings(narg, arg);
cuda->shared_data.pair.cut_global = (X_FLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_FLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_FLOAT) cut_lj_inner;
cuda->shared_data.pair.cut_global = (X_CFLOAT) cut_lj;
cuda->shared_data.pair.cut_coulsq_global = (X_CFLOAT) cut_coulsq;
cuda->shared_data.pair.cut_inner_global = (F_CFLOAT) cut_lj_inner;
}
/* ---------------------------------------------------------------------- */
@ -186,9 +186,9 @@ void PairLJGromacsCoulGromacsCuda::ev_setup(int eflag, int vflag)
PairLJGromacsCoulGromacs::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

18
src/USER-CUDA/pair_lj_gromacs_coul_gromacs_cuda.h
View File

@ -50,15 +50,15 @@ class PairLJGromacsCoulGromacsCuda : public PairLJGromacsCoulGromacs
void allocate();
bool allocated2;
class CudaNeighList* cuda_neigh_list;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw1_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw2_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw3_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw4_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw5_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw1_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw2_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw3_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw4_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw5_gm;
};

26
src/USER-CUDA/pair_lj_gromacs_cuda.cpp
View File

@ -82,15 +82,15 @@ void PairLJGromacsCuda::allocate()
cuda->shared_data.pair.coeff8 = ljsw4;
cuda->shared_data.pair.coeff9 = ljsw5;
cuda->shared_data.pair.special_lj = force->special_lj;
cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw1_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw1, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw2_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw2, &cuda->shared_data.pair.coeff6_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw3_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw3, &cuda->shared_data.pair.coeff7_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw4_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw4, &cuda->shared_data.pair.coeff8_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw5_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw5, &cuda->shared_data.pair.coeff9_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw1, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw2, &cuda->shared_data.pair.coeff6_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw3, &cuda->shared_data.pair.coeff7_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw4, &cuda->shared_data.pair.coeff8_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw5_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw5, &cuda->shared_data.pair.coeff9_gm, (atom->ntypes+1)*(atom->ntypes+1));
}
}
@ -119,8 +119,8 @@ void PairLJGromacsCuda::compute(int eflag, int vflag)
void PairLJGromacsCuda::settings(int narg, char **arg)
{
PairLJGromacs::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_inner_global = (F_FLOAT) cut_inner_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
cuda->shared_data.pair.cut_inner_global = (F_CFLOAT) cut_inner_global;
}
/* ---------------------------------------------------------------------- */
@ -169,9 +169,9 @@ void PairLJGromacsCuda::ev_setup(int eflag, int vflag)
PairLJGromacs::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

18
src/USER-CUDA/pair_lj_gromacs_cuda.h
View File

@ -50,15 +50,15 @@ class PairLJGromacsCuda : public PairLJGromacs
void allocate();
bool allocated2;
class CudaNeighList* cuda_neigh_list;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw1_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw2_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw3_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw4_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw5_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw1_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw2_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw3_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw4_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw5_gm;
};

8
src/USER-CUDA/pair_lj_sdk_coul_long_cuda.cpp
View File

@ -136,8 +136,8 @@ void PairLJSDKCoulLongCuda::compute(int eflag, int vflag)
void PairLJSDKCoulLongCuda::settings(int narg, char **arg)
{
PairLJSDKCoulLong::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_FLOAT) cut_coul;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_lj_global;
cuda->shared_data.pair.cut_coul_global = (F_CFLOAT) cut_coul;
}
/* ---------------------------------------------------------------------- */
@ -185,9 +185,9 @@ void PairLJSDKCoulLongCuda::ev_setup(int eflag, int vflag)
PairLJSDKCoulLong::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

16
src/USER-CUDA/pair_lj_sdk_cuda.cpp
View File

@ -96,11 +96,11 @@ void PairLJSDKCuda::allocate()
cuda->shared_data.pair.coeff3 = lj3;
cuda->shared_data.pair.coeff4 = lj4;
cuda->shared_data.pair.coeff5 = lj_type_double;
/*cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj_type_double_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj_type_double, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));*/
/*cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj_type_double_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj_type_double, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));*/
cuda->shared_data.pair.offset = offset;
cuda->shared_data.pair.special_lj = force->special_lj;
}
@ -135,7 +135,7 @@ void PairLJSDKCuda::compute(int eflag, int vflag)
void PairLJSDKCuda::settings(int narg, char **arg)
{
PairLJSDK::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -176,9 +176,9 @@ void PairLJSDKCuda::ev_setup(int eflag, int vflag)
PairLJSDK::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

10
src/USER-CUDA/pair_lj_sdk_cuda.h
View File

@ -52,11 +52,11 @@ class PairLJSDKCuda : public PairLJSDK
bool allocated2;
class CudaNeighList* cuda_neigh_list;
double** lj_type_double;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_FLOAT , x >* cu_lj_type_double_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj_type_double_gm;
};
}

26
src/USER-CUDA/pair_lj_smooth_cuda.cpp
View File

@ -82,15 +82,15 @@ void PairLJSmoothCuda::allocate()
cuda->shared_data.pair.coeff8 = ljsw4;
cuda->shared_data.pair.coeff9 = ljsw0;
cuda->shared_data.pair.special_lj = force->special_lj;
cu_lj1_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_FLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw0_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw0, &cuda->shared_data.pair.coeff9_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw1_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw1, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw2_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw2, &cuda->shared_data.pair.coeff6_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw3_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw3, &cuda->shared_data.pair.coeff7_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw4_gm = new cCudaData<double, F_FLOAT, x> ((double*)ljsw4, &cuda->shared_data.pair.coeff8_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj1, &cuda->shared_data.pair.coeff1_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj2, &cuda->shared_data.pair.coeff2_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj3, &cuda->shared_data.pair.coeff3_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_lj4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)lj4, &cuda->shared_data.pair.coeff4_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw0_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw0, &cuda->shared_data.pair.coeff9_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw1_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw1, &cuda->shared_data.pair.coeff5_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw2_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw2, &cuda->shared_data.pair.coeff6_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw3_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw3, &cuda->shared_data.pair.coeff7_gm, (atom->ntypes+1)*(atom->ntypes+1));
cu_ljsw4_gm = new cCudaData<double, F_CFLOAT, x> ((double*)ljsw4, &cuda->shared_data.pair.coeff8_gm, (atom->ntypes+1)*(atom->ntypes+1));
}
}
@ -119,8 +119,8 @@ void PairLJSmoothCuda::compute(int eflag, int vflag)
void PairLJSmoothCuda::settings(int narg, char **arg)
{
PairLJSmooth::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_inner_global = (F_FLOAT) cut_inner_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
cuda->shared_data.pair.cut_inner_global = (F_CFLOAT) cut_inner_global;
}
/* ---------------------------------------------------------------------- */
@ -169,9 +169,9 @@ void PairLJSmoothCuda::ev_setup(int eflag, int vflag)
PairLJSmooth::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

18
src/USER-CUDA/pair_lj_smooth_cuda.h
View File

@ -50,15 +50,15 @@ class PairLJSmoothCuda : public PairLJSmooth
void allocate();
bool allocated2;
class CudaNeighList* cuda_neigh_list;
cCudaData<double , F_FLOAT , x >* cu_lj1_gm;
cCudaData<double , F_FLOAT , x >* cu_lj2_gm;
cCudaData<double , F_FLOAT , x >* cu_lj3_gm;
cCudaData<double , F_FLOAT , x >* cu_lj4_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw0_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw1_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw2_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw3_gm;
cCudaData<double , F_FLOAT , x >* cu_ljsw4_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj1_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj2_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj3_gm;
cCudaData<double , F_CFLOAT , x >* cu_lj4_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw0_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw1_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw2_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw3_gm;
cCudaData<double , F_CFLOAT , x >* cu_ljsw4_gm;
};

6
src/USER-CUDA/pair_morse_cuda.cpp
View File

@ -117,7 +117,7 @@ void PairMorseCuda::compute(int eflag, int vflag)
void PairMorseCuda::settings(int narg, char **arg)
{
PairMorse::settings(narg, arg);
cuda->shared_data.pair.cut_global = (F_FLOAT) cut_global;
cuda->shared_data.pair.cut_global = (F_CFLOAT) cut_global;
}
/* ---------------------------------------------------------------------- */
@ -169,9 +169,9 @@ void PairMorseCuda::ev_setup(int eflag, int vflag)
PairMorse::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

4
src/USER-CUDA/pair_sw_cuda.cpp
View File

@ -200,8 +200,8 @@ void PairSWCuda::ev_setup(int eflag, int vflag)
PairSW::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

4
src/USER-CUDA/pair_tersoff_cuda.cpp
View File

@ -197,8 +197,8 @@ void PairTersoffCuda::ev_setup(int eflag, int vflag)
PairTersoff::ev_setup(eflag,vflag);
if (eflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_FLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
{delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData<double, ENERGY_CFLOAT, x > ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );}
if (vflag_atom && atom->nmax > maxeatomold)
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_FLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
{delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData<double, ENERGY_CFLOAT, yx > ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );}
}

54
src/USER-CUDA/pppm_cuda.cpp
View File

@ -747,7 +747,7 @@ void PPPMCuda::compute(int eflag, int vflag)
else
{
#ifdef FFT_CUFFT
pppm_initfftdata(&cuda->shared_data,(PPPM_FLOAT*)cu_density_brick->dev_data(),(FFT_FLOAT*)cu_work2->dev_data());
pppm_initfftdata(&cuda->shared_data,(PPPM_CFLOAT*)cu_density_brick->dev_data(),(FFT_CFLOAT*)cu_work2->dev_data());
#endif
}
@ -836,7 +836,7 @@ void PPPMCuda::allocate()
nxlo_out,nxhi_out,"pppm:density_brick_int");
cu_density_brick = new cCudaData<double, PPPM_FLOAT, x> ((double*) &(density_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
cu_density_brick = new cCudaData<double, PPPM_CFLOAT, x> ((double*) &(density_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
(nzhi_out-nzlo_out+1)*(nyhi_out-nylo_out+1)*(nxhi_out-nxlo_out+1));
cu_density_brick_int = new cCudaData<int, int, x> ((int*) &(density_brick_int[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
@ -847,48 +847,48 @@ void PPPMCuda::allocate()
memory->create3d_offset(vdx_brick_tmp,nzlo_out,nzhi_out,nylo_out,nyhi_out,
nxlo_out,nxhi_out,"pppm:vdx_brick_tmp");
cu_vdx_brick = new cCudaData<double, PPPM_FLOAT, x> ((double*) &(vdx_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
cu_vdx_brick = new cCudaData<double, PPPM_CFLOAT, x> ((double*) &(vdx_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
(nzhi_out-nzlo_out+1)*(nyhi_out-nylo_out+1)*(nxhi_out-nxlo_out+1));
memory->create3d_offset(vdy_brick,nzlo_out,nzhi_out,nylo_out,nyhi_out,
nxlo_out,nxhi_out,"pppm:vdy_brick");
cu_vdy_brick = new cCudaData<double, PPPM_FLOAT, x> ((double*) &(vdy_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
cu_vdy_brick = new cCudaData<double, PPPM_CFLOAT, x> ((double*) &(vdy_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
(nzhi_out-nzlo_out+1)*(nyhi_out-nylo_out+1)*(nxhi_out-nxlo_out+1));
memory->create3d_offset(vdz_brick,nzlo_out,nzhi_out,nylo_out,nyhi_out,
nxlo_out,nxhi_out,"pppm:vdz_brick");
cu_vdz_brick = new cCudaData<double, PPPM_FLOAT, x> ((double*) &(vdz_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
cu_vdz_brick = new cCudaData<double, PPPM_CFLOAT, x> ((double*) &(vdz_brick[nzlo_out][nylo_out][nxlo_out]), & (dev_tmp[n_cudata++]),
(nzhi_out-nzlo_out+1)*(nyhi_out-nylo_out+1)*(nxhi_out-nxlo_out+1));
memory->create(density_fft,nfft_both,"pppm:density_fft");
cu_density_fft = new cCudaData<double, PPPM_FLOAT, x> (density_fft, & (dev_tmp[n_cudata++]),nfft_both);
cu_density_fft = new cCudaData<double, PPPM_CFLOAT, x> (density_fft, & (dev_tmp[n_cudata++]),nfft_both);
cu_energy = new cCudaData<double, ENERGY_FLOAT, x> (NULL, &(dev_tmp[n_cudata++]),ny_pppm*nz_pppm);
cu_virial = new cCudaData<double, ENERGY_FLOAT, x> (NULL, &(dev_tmp[n_cudata++]),ny_pppm*nz_pppm*6);
cu_energy = new cCudaData<double, ENERGY_CFLOAT, x> (NULL, &(dev_tmp[n_cudata++]),ny_pppm*nz_pppm);
cu_virial = new cCudaData<double, ENERGY_CFLOAT, x> (NULL, &(dev_tmp[n_cudata++]),ny_pppm*nz_pppm*6);
memory->create(greensfn,nfft_both,"pppm:greensfn");
cu_greensfn = new cCudaData<double, PPPM_FLOAT, x> (greensfn, & (dev_tmp[n_cudata++]) , nx_pppm*ny_pppm*nz_pppm);
cu_greensfn = new cCudaData<double, PPPM_CFLOAT, x> (greensfn, & (dev_tmp[n_cudata++]) , nx_pppm*ny_pppm*nz_pppm);
memory->create(work1,2*nx_pppm*ny_pppm*nz_pppm,"pppm:work1");
memory->create(work2,2*nx_pppm*ny_pppm*nz_pppm,"pppm:work2");
memory->create(work3,2*nx_pppm*ny_pppm*nz_pppm,"pppm:work3");
cu_work1 = new cCudaData<double, FFT_FLOAT, x> (work1, & (dev_tmp[n_cudata++]) , 2*nx_pppm*ny_pppm*nz_pppm);
cu_work2 = new cCudaData<double, FFT_FLOAT, x> (work2, & (dev_tmp[n_cudata++]) , 2*nx_pppm*ny_pppm*nz_pppm);
cu_work3 = new cCudaData<double, FFT_FLOAT, x> (work3, & (dev_tmp[n_cudata++]) , 2*nx_pppm*ny_pppm*nz_pppm);
cu_work1 = new cCudaData<double, FFT_CFLOAT, x> (work1, & (dev_tmp[n_cudata++]) , 2*nx_pppm*ny_pppm*nz_pppm);
cu_work2 = new cCudaData<double, FFT_CFLOAT, x> (work2, & (dev_tmp[n_cudata++]) , 2*nx_pppm*ny_pppm*nz_pppm);
cu_work3 = new cCudaData<double, FFT_CFLOAT, x> (work3, & (dev_tmp[n_cudata++]) , 2*nx_pppm*ny_pppm*nz_pppm);
memory->create(fkx,nx_pppm,"pppmcuda:fkx");
cu_fkx = new cCudaData<double, PPPM_FLOAT, x> (fkx, & (dev_tmp[n_cudata++]) , nx_pppm);
cu_fkx = new cCudaData<double, PPPM_CFLOAT, x> (fkx, & (dev_tmp[n_cudata++]) , nx_pppm);
memory->create(fky,ny_pppm,"pppmcuda:fky");
cu_fky = new cCudaData<double, PPPM_FLOAT, x> (fky, & (dev_tmp[n_cudata++]) , ny_pppm);
cu_fky = new cCudaData<double, PPPM_CFLOAT, x> (fky, & (dev_tmp[n_cudata++]) , ny_pppm);
memory->create(fkz,nz_pppm,"pppmcuda:fkz");
cu_fkz = new cCudaData<double, PPPM_FLOAT, x> (fkz, & (dev_tmp[n_cudata++]) , nz_pppm);
cu_fkz = new cCudaData<double, PPPM_CFLOAT, x> (fkz, & (dev_tmp[n_cudata++]) , nz_pppm);
memory->create(vg,nfft_both,6,"pppm:vg");
cu_vg = new cCudaData<double, PPPM_FLOAT, xy> ((double*)vg, & (dev_tmp[n_cudata++]) , nfft_both,6);
cu_vg = new cCudaData<double, PPPM_CFLOAT, xy> ((double*)vg, & (dev_tmp[n_cudata++]) , nfft_both,6);
memory->create(buf1,nbuf,"pppm:buf1");
memory->create(buf2,nbuf,"pppm:buf2");
@ -898,14 +898,14 @@ void PPPMCuda::allocate()
gf_b = new double[order];
cu_gf_b = new cCudaData<double,PPPM_FLOAT,x> (gf_b, &(dev_tmp[n_cudata++]) , order);
cu_gf_b = new cCudaData<double,PPPM_CFLOAT,x> (gf_b, &(dev_tmp[n_cudata++]) , order);
memory->create2d_offset(rho1d,3,-order/2,order/2,"pppm:rho1d");
memory->create2d_offset(rho_coeff,order,(1-order)/2,order/2,"pppm:rho_coeff");
cu_rho_coeff = new cCudaData<double, PPPM_FLOAT, x> ((double*) &(rho_coeff[0][(1-order)/2]), & (dev_tmp[n_cudata++]) , order*(order/2-(1-order)/2+1));
cu_rho_coeff = new cCudaData<double, PPPM_CFLOAT, x> ((double*) &(rho_coeff[0][(1-order)/2]), & (dev_tmp[n_cudata++]) , order*(order/2-(1-order)/2+1));
debugdata=new PPPM_FLOAT[100];
cu_debugdata = new cCudaData<PPPM_FLOAT, PPPM_FLOAT, x> (debugdata,& (dev_tmp[n_cudata++]),100);
debugdata=new PPPM_CFLOAT[100];
cu_debugdata = new cCudaData<PPPM_CFLOAT, PPPM_CFLOAT, x> (debugdata,& (dev_tmp[n_cudata++]),100);
cu_flag = new cCudaData<int, int, x> (&global_flag,& (dev_tmp[n_cudata++]),3);
// create 2 FFTs and a Remap
@ -1308,7 +1308,7 @@ void PPPMCuda::poisson(int eflag, int vflag)
if (eflag || vflag) {
poisson_energy(nxlo_fft,nxhi_fft,nylo_fft,nyhi_fft,nzlo_fft,nzhi_fft,vflag);
ENERGY_FLOAT gpuvirial[6];
ENERGY_CFLOAT gpuvirial[6];
energy+=sum_energy(cu_virial->dev_data(),cu_energy->dev_data(),nx_pppm,ny_pppm,nz_pppm,vflag,gpuvirial);
if(vflag)
{
@ -1395,19 +1395,19 @@ void PPPMCuda::slabcorr(int eflag)
// compute local contribution to global dipole moment
if(slabbuf==NULL)
{
slabbuf=new ENERGY_FLOAT[(atom->nmax+31)/32];
cu_slabbuf = new cCudaData<ENERGY_FLOAT,ENERGY_FLOAT, x> (slabbuf, (atom->nmax+31)/32);
slabbuf=new ENERGY_CFLOAT[(atom->nmax+31)/32];
cu_slabbuf = new cCudaData<ENERGY_CFLOAT,ENERGY_CFLOAT, x> (slabbuf, (atom->nmax+31)/32);
}
if(unsigned((atom->nlocal+31)/32)*sizeof(ENERGY_FLOAT)>=unsigned(cu_slabbuf->dev_size()))
if(unsigned((atom->nlocal+31)/32)*sizeof(ENERGY_CFLOAT)>=unsigned(cu_slabbuf->dev_size()))
{
delete [] slabbuf;
delete cu_slabbuf;
slabbuf=new ENERGY_FLOAT[(atom->nmax+31)/32];
cu_slabbuf = new cCudaData<ENERGY_FLOAT,ENERGY_FLOAT, x> (slabbuf, (atom->nmax+31)/32);
slabbuf=new ENERGY_CFLOAT[(atom->nmax+31)/32];
cu_slabbuf = new cCudaData<ENERGY_CFLOAT,ENERGY_CFLOAT, x> (slabbuf, (atom->nmax+31)/32);
}
double dipole = cuda_slabcorr_energy(&cuda->shared_data,slabbuf,(ENERGY_FLOAT*) cu_slabbuf->dev_data());
double dipole = cuda_slabcorr_energy(&cuda->shared_data,slabbuf,(ENERGY_CFLOAT*) cu_slabbuf->dev_data());
double dipole_all;
MPI_Allreduce(&dipole,&dipole_all,1,MPI_DOUBLE,MPI_SUM,world);

52
src/USER-CUDA/pppm_cuda.h
View File

@ -53,43 +53,43 @@ class PPPMCuda : public PPPMOld {
class FFT3dCuda *fft1c,*fft2c;
double* work3;
cCudaData<double , FFT_FLOAT , x >* cu_work1;
cCudaData<double , FFT_FLOAT , x >* cu_work2;
cCudaData<double , FFT_FLOAT , x >* cu_work3;
cCudaData<double , PPPM_FLOAT , x >* cu_greensfn;
cCudaData<double , PPPM_FLOAT , x >* cu_gf_b;
cCudaData<double , PPPM_FLOAT , x >* cu_fkx;
cCudaData<double , PPPM_FLOAT , x >* cu_fky;
cCudaData<double , PPPM_FLOAT , x >* cu_fkz;
cCudaData<double , PPPM_FLOAT , xy>* cu_vg;
cCudaData<double , PPPM_FLOAT , x >* cu_density_brick;
cCudaData<double , FFT_CFLOAT , x >* cu_work1;
cCudaData<double , FFT_CFLOAT , x >* cu_work2;
cCudaData<double , FFT_CFLOAT , x >* cu_work3;
cCudaData<double , PPPM_CFLOAT , x >* cu_greensfn;
cCudaData<double , PPPM_CFLOAT , x >* cu_gf_b;
cCudaData<double , PPPM_CFLOAT , x >* cu_fkx;
cCudaData<double , PPPM_CFLOAT , x >* cu_fky;
cCudaData<double , PPPM_CFLOAT , x >* cu_fkz;
cCudaData<double , PPPM_CFLOAT , xy>* cu_vg;
cCudaData<double , PPPM_CFLOAT , x >* cu_density_brick;
cCudaData<int , int , x >* cu_density_brick_int;
cCudaData<double , PPPM_FLOAT , x >* cu_vdx_brick;
cCudaData<double , PPPM_FLOAT , x >* cu_vdy_brick;
cCudaData<double , PPPM_FLOAT , x >* cu_vdz_brick;
cCudaData<double , PPPM_FLOAT , x >* cu_density_fft;
cCudaData<double , ENERGY_FLOAT , x >* cu_energy;
cCudaData<double , ENERGY_FLOAT , x >* cu_virial;
cCudaData<double , X_FLOAT , yx>* cu_x;
cCudaData<double , V_FLOAT , yx>* cu_v;
cCudaData<double , F_FLOAT , yx>* cu_f;
cCudaData<double , F_FLOAT , yx>* cu_q;
cCudaData<double , PPPM_CFLOAT , x >* cu_vdx_brick;
cCudaData<double , PPPM_CFLOAT , x >* cu_vdy_brick;
cCudaData<double , PPPM_CFLOAT , x >* cu_vdz_brick;
cCudaData<double , PPPM_CFLOAT , x >* cu_density_fft;
cCudaData<double , ENERGY_CFLOAT , x >* cu_energy;
cCudaData<double , ENERGY_CFLOAT , x >* cu_virial;
cCudaData<double , X_CFLOAT , yx>* cu_x;
cCudaData<double , V_CFLOAT , yx>* cu_v;
cCudaData<double , F_CFLOAT , yx>* cu_f;
cCudaData<double , F_CFLOAT , yx>* cu_q;
cCudaData<int , int , yx>* cu_part2grid;
cCudaData<double , PPPM_FLOAT , x >* cu_rho_coeff;
cCudaData<PPPM_FLOAT , PPPM_FLOAT , x >* cu_debugdata;
cCudaData<double , PPPM_CFLOAT , x >* cu_rho_coeff;
cCudaData<PPPM_CFLOAT , PPPM_CFLOAT , x >* cu_debugdata;
cCudaData<int , int , x >* cu_flag;
cCudaData<int , int , x >* cu_pppm_grid_n;
cCudaData<int , int , x >* cu_pppm_grid_ids;
ENERGY_FLOAT* slabbuf;
cCudaData<ENERGY_FLOAT, ENERGY_FLOAT, x >* cu_slabbuf;
ENERGY_CFLOAT* slabbuf;
cCudaData<ENERGY_CFLOAT, ENERGY_CFLOAT, x >* cu_slabbuf;
int*** density_brick_int;
PPPM_FLOAT density_intScale;
PPPM_CFLOAT density_intScale;
int pppm_grid_nmax;
int* pppm2partgrid;
int* pppm_grid;
PPPM_FLOAT* debugdata;
PPPM_CFLOAT* debugdata;
bool firstpass;
void set_grid();

2
src/USER-CUDA/pppm_old.h
View File

@ -25,7 +25,7 @@ KSpaceStyle(pppm/old,PPPMOld)
#ifdef FFT_SINGLE
typedef float FFT_SCALAR;
#define MPI_FFT_SCALAR MPI_FLOAT
#define MPI_FFT_SCALAR MPI_CFLOAT
#else
typedef double FFT_SCALAR;
#define MPI_FFT_SCALAR MPI_DOUBLE

8
src/USER-CUDA/verlet_cuda.cpp
View File

@ -116,7 +116,7 @@ void VerletCuda::setup()
if(cuda->shared_data.me == 0)
printf("# CUDA: Using precision: Global: %u X: %u V: %u F: %u PPPM: %u \n", CUDA_PRECISION == 1 ? 4 : 8, (int) sizeof(X_FLOAT), (int) sizeof(V_FLOAT), (int) sizeof(F_FLOAT), (int) sizeof(PPPM_FLOAT));
printf("# CUDA: Using precision: Global: %u X: %u V: %u F: %u PPPM: %u \n", CUDA_PRECISION == 1 ? 4 : 8, (int) sizeof(X_CFLOAT), (int) sizeof(V_CFLOAT), (int) sizeof(F_CFLOAT), (int) sizeof(PPPM_CFLOAT));
cuda->allocate();
@ -638,9 +638,9 @@ void VerletCuda::run(int n)
if(cuda->neighbor_decide_by_integrator && cuda->cu_xhold) {
const int n = cuda->shared_data.atom.maxhold;
CudaWrapper_CopyData(cuda->cu_xhold->dev_data(), cuda->cu_x->dev_data(), n * sizeof(X_FLOAT));
CudaWrapper_CopyData((void*) & ((X_FLOAT*)cuda->cu_xhold->dev_data())[n], (void*) & ((X_FLOAT*)cuda->cu_x->dev_data())[atom->nmax], n * sizeof(X_FLOAT));
CudaWrapper_CopyData((void*) & ((X_FLOAT*)cuda->cu_xhold->dev_data())[2 * n], (void*) & ((X_FLOAT*)cuda->cu_x->dev_data())[2 * atom->nmax], n * sizeof(X_FLOAT));
CudaWrapper_CopyData(cuda->cu_xhold->dev_data(), cuda->cu_x->dev_data(), n * sizeof(X_CFLOAT));
CudaWrapper_CopyData((void*) & ((X_CFLOAT*)cuda->cu_xhold->dev_data())[n], (void*) & ((X_CFLOAT*)cuda->cu_x->dev_data())[atom->nmax], n * sizeof(X_CFLOAT));
CudaWrapper_CopyData((void*) & ((X_CFLOAT*)cuda->cu_xhold->dev_data())[2 * n], (void*) & ((X_CFLOAT*)cuda->cu_x->dev_data())[2 * atom->nmax], n * sizeof(X_CFLOAT));
}
cuda->shared_data.atom.reneigh_flag = 0;