Created #define for LOCAL_HACK_PI to sidestep bug on Windows cuda compiler...

Created #define for LOCAL_HACK_PI to sidestep bug on Windows cuda compiler which ignores "static const float ..." variables.
In files kCalculateCDLJEwaldFastReciprocal.h, kCalculateCDLJForces.h and kCalculateLocalForces.cu.

platforms/cuda/src/kernels/kCalculateCDLJEwaldFastReciprocal.h

@@ -29,6 +29,9 @@
  * Ewald summation method (Reciprocal space summation).
  */
 
+/* Cuda compiler on Windows does not recognized "static const float" values */
+#define LOCAL_HACK_PI 3.1415926535897932384626433832795
+
 /* Define multiply operations for floats */
 
 __device__ float2 MultofFloat2(float2 a, float2 b)
@@ -54,7 +57,7 @@ __device__ float2 ConjMultofFloat2(float2 a, float2 b)
 __global__ void kCalculateEwaldFastCosSinSums_kernel()
 {
     const float epsilon =  1.0;
-    const float recipCoeff = cSim.epsfac*(4*PI/cSim.cellVolume/epsilon);
+    const float recipCoeff = cSim.epsfac*(4*LOCAL_HACK_PI/cSim.cellVolume/epsilon);
     const unsigned int ksizex = 2*cSim.kmaxX-1;
     const unsigned int ksizey = 2*cSim.kmaxY-1;
     const unsigned int ksizez = 2*cSim.kmaxZ-1;
@@ -111,7 +114,7 @@ __global__ void kCalculateEwaldFastCosSinSums_kernel()
 __global__ void kCalculateEwaldFastForces_kernel()
 {
     const float epsilon =  1.0;
-    float recipCoeff = cSim.epsfac*(4*PI/cSim.cellVolume/epsilon);
+    float recipCoeff = cSim.epsfac*(4*LOCAL_HACK_PI/cSim.cellVolume/epsilon);
 
     unsigned int atom = threadIdx.x + blockIdx.x * blockDim.x;
 

platforms/cuda/src/kernels/kCalculateCDLJForces.h

@@ -30,6 +30,9 @@
  * different versions of the kernels.
  */
 
+/* Cuda compiler on Windows does not recognized "static const float" values */
+#define LOCAL_HACK_PI 3.1415926535897932384626433832795
+
 __global__ void METHOD_NAME(kCalculateCDLJ, Forces_kernel)(unsigned int* workUnit)
 {
     extern __shared__ Atom sA[];
@@ -45,7 +48,7 @@ __global__ void METHOD_NAME(kCalculateCDLJ, Forces_kernel)(unsigned int* workUni
 #endif
 
 #ifdef USE_EWALD
-    const float SQRT_PI = sqrt(PI);
+    const float SQRT_PI = sqrt(LOCAL_HACK_PI);
 #endif
 
     unsigned int lasty = 0xFFFFFFFF;

platforms/cuda/src/kernels/kCalculateLocalForces.cu

@@ -38,6 +38,8 @@ using namespace std;
 extern __shared__ Vectors sV[];
 static __constant__ cudaGmxSimulation cSim;
 
+/* Cuda compiler on Windows does not recognized "static const float" values */
+#define LOCAL_HACK_PI 3.1415926535897932384626433832795
 
 #define DOT3(v1, v2) (v1.x * v2.x + v1.y * v2.y + v1.z * v2.z)
 
@@ -59,14 +61,14 @@ static __constant__ cudaGmxSimulation cSim;
 #define GETPREFACTORSGIVENANGLECOSINE(cosine, param, dEdR) \
 { \
    float angle          = acos(cosine); \
-   float deltaIdeal     = angle - (param.x * (PI / 180.0f)); \
+   float deltaIdeal     = angle - (param.x * (LOCAL_HACK_PI / 180.0f)); \
    dEdR                 = param.y * deltaIdeal; \
 }
 
 #define GETENERGYGIVENANGLECOSINE(cosine, param, dEdR) \
 { \
    float angle          = acos(cosine); \
-   float deltaIdeal     = angle - (param.x * (PI / 180.0f)); \
+   float deltaIdeal     = angle - (param.x * (LOCAL_HACK_PI / 180.0f)); \
    dEdR                 = param.y * deltaIdeal * deltaIdeal; \
 }
 
@@ -250,7 +252,7 @@ __global__ void kCalculateLocalForces_kernel()
             float dihedralAngle;
             GETDIHEDRALANGLEBETWEENTHREEVECTORS(A->v0, A->v1, A->v2, A->v0, cp0, cp1, dihedralAngle);
             float4 dihedral         = cSim.pDihedralParameter[pos1];
-            float deltaAngle        = dihedral.z * dihedralAngle - (dihedral.y * PI / 180.0f);
+            float deltaAngle        = dihedral.z * dihedralAngle - (dihedral.y * LOCAL_HACK_PI / 180.0f);
 
 	    // ATTENTION: This section leads to a divergent deltaAngle values wrt
 	    // forces and energies. We separate the case dihedral.z = n = 0, which
@@ -259,8 +261,8 @@ __global__ void kCalculateLocalForces_kernel()
 /* E */     else
 	    {
 		float deltaAngle    = dihedralAngle - dihedral.y;
-		if (deltaAngle < -PI) deltaAngle += 2.0f * PI;
-		else if (deltaAngle > PI) deltaAngle -= 2.0f * PI;
+		if (deltaAngle < -LOCAL_HACK_PI) deltaAngle += 2.0f * LOCAL_HACK_PI;
+		else if (deltaAngle > LOCAL_HACK_PI) deltaAngle -= 2.0f * LOCAL_HACK_PI;
                 energy             += dihedral.x * deltaAngle * deltaAngle;
 	    }
 
@@ -353,11 +355,11 @@ __global__ void kCalculateLocalForces_kernel()
             GETDIHEDRALANGLECOSINEBETWEENTHREEVECTORS(A->v0, A->v1, A->v2, A->v0, cp0, cp1, dihedralAngle, cosPhi);
             if (dihedralAngle < 0.0f )
             {
-                dihedralAngle += PI;
+                dihedralAngle += LOCAL_HACK_PI;
             }
             else
             {
-                dihedralAngle -= PI;
+                dihedralAngle -= LOCAL_HACK_PI;
             }
             cosPhi                  = -cosPhi;
          //   printf("%4d: %9.4f %9.4f\n", pos1, dihedralAngle, cosPhi);