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

update a few GPU kernels so they can be compiled on GPUs without double precisions support

Browse Source
This commit is contained in:
Axel Kohlmeyer
2021-10-21 07:33:00 -04:00
parent 2dce8923ee
commit 0325047c01
4 changed files with 11 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
lib/gpu/lal_atom.cu
View File

@ -18,7 +18,7 @@
#endif
__kernel void kernel_cast_x(__global numtyp4 *restrict x_type,
const __global double *restrict x,
const __global numtyp *restrict x,
const __global int *restrict type,
const int nall) {
int ii=GLOBAL_ID_X;

2
lib/gpu/lal_atom.h
View File

@ -475,7 +475,7 @@ class Atom {
UCL_Vector<numtyp,numtyp> v;
#ifdef GPU_CAST
UCL_Vector<double,double> x_cast;
UCL_Vector<numtyp,numtyp> x_cast;
UCL_Vector<int,int> type_cast;
#endif

12
lib/gpu/lal_zbl.cu
View File

@ -82,9 +82,9 @@ __kernel void k_zbl(const __global numtyp4 *restrict x_,
const __global numtyp4 *restrict coeff1,
const __global numtyp4 *restrict coeff2,
const __global numtyp4 *restrict coeff3,
const double cut_globalsq,
const double cut_innersq,
const double cut_inner,
const numtyp cut_globalsq,
const numtyp cut_innersq,
const numtyp cut_inner,
const int lj_types,
const __global int *dev_nbor,
const __global int *dev_packed,
@ -174,9 +174,9 @@ __kernel void k_zbl_fast(const __global numtyp4 *restrict x_,
const __global numtyp4 *restrict coeff1_in,
const __global numtyp4 *restrict coeff2_in,
const __global numtyp4 *restrict coeff3_in,
const double cut_globalsq,
const double cut_innersq,
const double cut_inner,
const numtyp cut_globalsq,
const numtyp cut_innersq,
const numtyp cut_inner,
const __global int *dev_nbor,
const __global int *dev_packed,
__global acctyp4 *restrict ans,

6
lib/gpu/lal_zbl.h
View File

@ -67,9 +67,9 @@ class ZBL : public BaseAtomic<numtyp, acctyp> {
/// If atom type constants fit in shared memory, use fast kernels
bool shared_types;
double _cut_globalsq;
double _cut_innersq;
double _cut_inner;
numtyp _cut_globalsq;
numtyp _cut_innersq;
numtyp _cut_inner;
/// Number of atom types
int _lj_types;