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

Delete "float.h" from TIP4P GPU kernel for better OpenCL compatibility

Browse Source 
According to Khronos OpenCL docs, "The C99 standard headers <...>, float.h, <...> are not available and cannot be included by a program"
This commit is contained in:
Vsevak
2019-11-27 19:49:12 +03:00
parent 4febc7f794
commit 66a076b819
1 changed files with 5 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
lib/gpu/lal_lj_tip4p_long.cu
View File

@ -29,8 +29,6 @@ texture<int2> q_tex;
#define q_tex q_ #define q_tex q_
#endif #endif
#include <float.h>
ucl_inline int atom_mapping(const __global int *map, int glob){ ucl_inline int atom_mapping(const __global int *map, int glob){
return map[glob]; return map[glob];
} }
@ -184,6 +182,8 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
int tid, ii, offset; int tid, ii, offset;
atom_info(t_per_atom,ii,tid,offset); atom_info(t_per_atom,ii,tid,offset);
const numtyp eq_zero = 1e-6;
acctyp energy = (acctyp)0; acctyp energy = (acctyp)0;
acctyp e_coul = (acctyp)0; acctyp e_coul = (acctyp)0;
acctyp4 f, fO; acctyp4 f, fO;
@ -230,7 +230,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
iH1 = hneigh[i*4 ]; iH1 = hneigh[i*4 ];
iH2 = hneigh[i*4+1]; iH2 = hneigh[i*4+1];
} }
if(fabs(m[iO].w) <= FLT_EPSILON) { if(fabs(m[iO].w) <= eq_zero) {
compute_newsite(iO,iH1,iH2, &m[iO], alpha, x_); compute_newsite(iO,iH1,iH2, &m[iO], alpha, x_);
m[iO].w = qtmp; m[iO].w = qtmp;
} }
@ -270,7 +270,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
} }
if (iO >= inum) { if (iO >= inum) {
non_local_oxy = 1; non_local_oxy = 1;
if(fabs(m[iO].w) <= FLT_EPSILON) { if(fabs(m[iO].w) <= eq_zero) {
compute_newsite(iO,iH1,iH2, &m[iO], alpha, x_); compute_newsite(iO,iH1,iH2, &m[iO], alpha, x_);
numtyp qO; fetch(qO,iO,q_tex); numtyp qO; fetch(qO,iO,q_tex);
m[iO].w = qO; m[iO].w = qO;
@ -346,7 +346,7 @@ __kernel void k_lj_tip4p_long(const __global numtyp4 *restrict x_,
jH1 = hneigh[j*4 ]; jH1 = hneigh[j*4 ];
jH2 = hneigh[j*4+1]; jH2 = hneigh[j*4+1];
} }
if (fabs(m[j].w) <= FLT_EPSILON) { if (fabs(m[j].w) <= eq_zero) {
compute_newsite(j, jH1, jH2, &m[j], alpha, x_); compute_newsite(j, jH1, jH2, &m[j], alpha, x_);
m[j].w = qj; m[j].w = qj;
} }