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

use sin(x+pi/2) instead of cos(x) on Intel OpenCL with double precision

Browse Source
This commit is contained in:
Axel Kohlmeyer
2023-03-02 18:29:05 -05:00
parent da79aadc8d
commit fa38047749
1 changed files with 11 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
lib/gpu/lal_soft.cu
View File

@ -22,6 +22,15 @@ _texture_2d( pos_tex,int4);
#endif #endif
#else #else
#define pos_tex x_ #define pos_tex x_
// hack for Intel GPU with double precision
#if defined(_DOUBLE_DOUBLE) && (CONFIG_ID == 500)
#define MY_PI_HALF (acctyp)1.57079632679489661923
#define my_cos(x) sin(x+MY_PI_HALF)
#endif
#endif
#if !defined(my_cos)
#define my_cos(x) cos(x)
#endif #endif
#define MY_PI (acctyp)3.14159265358979323846 #define MY_PI (acctyp)3.14159265358979323846
@ -95,7 +104,7 @@ __kernel void k_soft(const __global numtyp4 *restrict x_,
f.z+=delz*force; f.z+=delz*force;
if (EVFLAG && eflag) { if (EVFLAG && eflag) {
numtyp e=coeff[mtype].x * ((numtyp)1.0+cos(arg)); numtyp e=coeff[mtype].x * ((numtyp)1.0+my_cos(arg));
energy+=factor_lj*e; energy+=factor_lj*e;
} }
if (EVFLAG && vflag) { if (EVFLAG && vflag) {
@ -186,7 +195,7 @@ __kernel void k_soft_fast(const __global numtyp4 *restrict x_,
f.z+=delz*force; f.z+=delz*force;
if (EVFLAG && eflag) { if (EVFLAG && eflag) {
numtyp e=coeff[mtype].x * ((numtyp)1.0+cos(arg)); numtyp e=coeff[mtype].x * ((numtyp)1.0+my_cos(arg));
energy+=factor_lj*e; energy+=factor_lj*e;
} }
if (EVFLAG && vflag) { if (EVFLAG && vflag) {