Fixed compilation errors on Linux

plugins/amoeba/platforms/cuda/src/kernels/AmoebaGpu.cpp

@@ -38,6 +38,7 @@ extern void OPENMMCUDA_EXPORT SetForcesSim(gpuContext gpu);
 
 #include <sstream>
 #include <limits>
+#include <cstring>
 
 #ifdef WIN32
 //  #include <windows.h>

plugins/amoeba/platforms/cuda/src/kernels/amoebaScaleFactors.h

@@ -41,14 +41,14 @@ static __constant__ float directScale[5]  = { 0.0f, 1.0f, 1.0f, 1.0f, 1.0f };
            
 // subroutine to get masked scale factors
 
-__device__ void getMaskedDScaleFactor( unsigned int gridIndex, int scaleMask, float* dScale )
+__device__ static void getMaskedDScaleFactor( unsigned int gridIndex, int scaleMask, float* dScale )
 {
     unsigned int mask             = 1 << gridIndex;   
     *dScale                       = (scaleMask & mask) ? 0.0f : 1.0f;
 
 }
 
-__device__ void getMaskedPScaleFactor( unsigned int gridIndex, int2 scaleMask, float* pScale )
+__device__ static void getMaskedPScaleFactor( unsigned int gridIndex, int2 scaleMask, float* pScale )
 {
     unsigned int mask             = 1 << gridIndex;   
     *pScale                       = (scaleMask.x & mask) ? 0.5f : 1.0f;
@@ -56,7 +56,7 @@ __device__ void getMaskedPScaleFactor( unsigned int gridIndex, int2 scaleMask, f
 
 }
 
-__device__ void getMaskedMScaleFactor( unsigned int gridIndex, int2 scaleMask, float* mScale )
+__device__ static void getMaskedMScaleFactor( unsigned int gridIndex, int2 scaleMask, float* mScale )
 {
     unsigned int mask             = 1 << gridIndex;   
 
@@ -73,7 +73,7 @@ __device__ void getMaskedMScaleFactor( unsigned int gridIndex, int2 scaleMask, f
 
 // subroutine to get cell coordinates
 
-__device__ void decodeCell( unsigned int cellId, unsigned int* x, unsigned int* y, bool* exclusions )
+__device__ static void decodeCell( unsigned int cellId, unsigned int* x, unsigned int* y, bool* exclusions )
 {
     *x          = cellId;
     *y          = ((*x >> 2) & 0x7fff) << GRIDBITS;
@@ -83,7 +83,7 @@ __device__ void decodeCell( unsigned int cellId, unsigned int* x, unsigned int*
 
 }
 
-__device__ void load3dArrayBufferPerWarp( unsigned int offset, float* forceSum, float* outputForce )
+__device__ static void load3dArrayBufferPerWarp( unsigned int offset, float* forceSum, float* outputForce )
 {
 
     float of; 
@@ -101,7 +101,7 @@ __device__ void load3dArrayBufferPerWarp( unsigned int offset, float* forceSum,
 
 }
 
-__device__ void load3dArray( unsigned int offset, float* forceSum, float* outputForce )
+__device__ static void load3dArray( unsigned int offset, float* forceSum, float* outputForce )
 {
 
     outputForce[offset]                 = forceSum[0];  
@@ -110,7 +110,7 @@ __device__ void load3dArray( unsigned int offset, float* forceSum, float* output
 
 }
 
-__device__ void scale3dArray( float scaleFactor, float* force )
+__device__ static void scale3dArray( float scaleFactor, float* force )
 {
 
     force[0]  *= scaleFactor;  

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaFixedFieldParticle.h

@@ -124,7 +124,7 @@ __device__ void loadFixedFieldParticleData( struct FixedFieldParticle* sA,
 
 // zero fields
 
-__device__ void zeroFixedFieldParticleSharedField( struct FixedFieldParticle* sA ) 
+__device__ static void zeroFixedFieldParticleSharedField( struct FixedFieldParticle* sA )
 {
 
     sA->eField[0]    = 0.0f;
@@ -144,7 +144,7 @@ __device__ void zeroFixedFieldParticleSharedField( struct FixedFieldParticle* sA
 
 // body of fixed E-field calculation
 
-__device__ void calculateFixedEFieldPairIxn_kernel( float4 atomCoordinatesI, float4 atomCoordinatesJ,
+__device__ static void calculateFixedEFieldPairIxn_kernel( float4 atomCoordinatesI, float4 atomCoordinatesJ,
                                                     float dampingFactorI,    float dampingFactorJ,
                                                     float tholeI,            float tholeJ,
                                                     float* labDipoleI, float* labDipoleJ,

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwood.cu

@@ -1997,7 +1997,7 @@ int debugIndex               = atomJ;
  
 }
 
-__device__ int debugAccumulate( int index, float4* debugArray, float* field, unsigned int addMask, float idLabel )
+__device__ static int debugAccumulate( int index, float4* debugArray, float* field, unsigned int addMask, float idLabel )
 {
     index                             += cAmoebaSim.paddedNumberOfAtoms;
     debugArray[index].x                = addMask ? field[0] : 0.0f;

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaKirkwoodEDiff.cu

@@ -1031,7 +1031,7 @@ __device__ void calculateKirkwoodEDiffPairIxn_kernel( float4 atomCoordinatesI,
 
 }
 
-__device__ int debugAccumulate( unsigned int index, float4* debugArray, float* field, unsigned int addMask, float idLabel )
+__device__ static int debugAccumulate( unsigned int index, float4* debugArray, float* field, unsigned int addMask, float idLabel )
 {
     index                             += cAmoebaSim.paddedNumberOfAtoms;
     debugArray[index].x                = addMask ? field[0] : 0.0f;

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaMutualInducedAndGkFields.cu

@@ -267,7 +267,7 @@ __device__ void calculateMutualInducedAndGkFieldsGkPairIxn_kernel( float4 atomCo
 
 }
 
-__device__ int debugAccumulate( int index, float4* debugArray, float* field, unsigned int addMask, float idLabel )
+__device__ static int debugAccumulate( int index, float4* debugArray, float* field, unsigned int addMask, float idLabel )
 {
     index                             += cAmoebaSim.paddedNumberOfAtoms;
     debugArray[index].x                = addMask ? field[0] : 0.0f;

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaMutualInducedParticle.h

@@ -77,7 +77,7 @@ __device__ void loadMutualInducedShared( struct MutualInducedParticle* sA, unsig
 #endif
 }
 
-__device__ void zeroMutualInducedParticleSharedField( struct MutualInducedParticle* sA )
+__device__ static void zeroMutualInducedParticleSharedField( struct MutualInducedParticle* sA )
 
 {
     // zero shared fields

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaMapTorques.cu

@@ -28,7 +28,7 @@ void GetCalculateAmoebaCudaMapTorquesSim(amoebaGpuContext amoebaGpu)
     RTERROR(status, "GetCalculateAmoebaCudaMapTorquesSim: cudaMemcpyFromSymbol: SetSim copy from cAmoebaSim failed");
 }
 
-__device__ float normVector3( float* vector ) 
+__device__ static float normVector3( float* vector )
 {
 
     float norm                    = DOT3( vector, vector );

plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaRotateFrame.cu

@@ -34,7 +34,7 @@ void GetCalculateAmoebaMultipoleForcesSim(amoebaGpuContext amoebaGpu)
     RTERROR(status, "GetCalculateAmoebaMultipoleForcesSim: cudaMemcpyFromSymbol: SetSim copy from cAmoebaSim failed");
 }
 
-__device__ float normVector3( float* vector ) 
+__device__ static float normVector3( float* vector )
 {
 
     float norm                    = DOT3( vector, vector );

plugins/amoeba/platforms/cuda/tests/TestCudaAmoebaHarmonicBondForce.cpp

@@ -54,7 +54,7 @@ static void computeAmoebaHarmonicBondForce(int bondIndex,  std::vector<Vec3>& po
     double quadraticK;
     double cubicK;
     double quarticK;
-    amoebaHarmonicBondForce.getBondParameters(bondIndex, particle1, particle2,  bondLength,  quadraticK,  cubicK,  quarticK );
+    amoebaHarmonicBondForce.getBondParameters(bondIndex, particle1, particle2,  bondLength,  quadraticK );
 
     double deltaR[3];
     double r2 = 0.0;
@@ -152,7 +152,7 @@ void testOneBond( FILE* log ) {
     double quadraticK = 1.0;
     double cubicK     = 2.0;
     double quarticicK = 3.0;
-    amoebaHarmonicBondForce->addBond(0, 1, bondLength, quadraticK, cubicK,quarticicK);
+    amoebaHarmonicBondForce->addBond(0, 1, bondLength, quadraticK);
 
     system.addForce(amoebaHarmonicBondForce);
     Context context(system, integrator, Platform::getPlatformByName( "Cuda"));
@@ -181,8 +181,8 @@ void testTwoBond( FILE* log ) {
     double quadraticK = 1.0;
     double cubicK     = 2.0;
     double quarticicK = 3.0;
-    amoebaHarmonicBondForce->addBond(0, 1, bondLength, quadraticK, cubicK, quarticicK);
-    amoebaHarmonicBondForce->addBond(1, 2, bondLength, quadraticK, cubicK, quarticicK);
+    amoebaHarmonicBondForce->addBond(0, 1, bondLength, quadraticK);
+    amoebaHarmonicBondForce->addBond(1, 2, bondLength, quadraticK);
 
     system.addForce(amoebaHarmonicBondForce);
     Context context(system, integrator, Platform::getPlatformByName( "Cuda"));