From fa38047749bcb349e0a97a97efea8ec10f9013f1 Mon Sep 17 00:00:00 2001
From: Axel Kohlmeyer <akohlmey@gmail.com>
Date: Thu, 2 Mar 2023 18:29:05 -0500
Subject: [PATCH] use sin(x+pi/2) instead of cos(x) on Intel OpenCL with double
 precision

---
 lib/gpu/lal_soft.cu | 13 +++++++++++--
 1 file changed, 11 insertions(+), 2 deletions(-)

diff --git a/lib/gpu/lal_soft.cu b/lib/gpu/lal_soft.cu
index 74ac0e0c97..bc96a861df 100644
--- a/lib/gpu/lal_soft.cu
+++ b/lib/gpu/lal_soft.cu
@@ -22,6 +22,15 @@ _texture_2d( pos_tex,int4);
 #endif
 #else
 #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
 
 #define MY_PI (acctyp)3.14159265358979323846
@@ -95,7 +104,7 @@ __kernel void k_soft(const __global numtyp4 *restrict x_,
         f.z+=delz*force;
 
         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;
         }
         if (EVFLAG && vflag) {
@@ -186,7 +195,7 @@ __kernel void k_soft_fast(const __global numtyp4 *restrict x_,
         f.z+=delz*force;
 
         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;
         }
         if (EVFLAG && vflag) {