Added static to __device__ declarations to allow compilation w/ CUDA 3.1

platforms/cuda/src/kernels/bbsort_kernel.cu

@@ -19,23 +19,23 @@ texture<unsigned int, 1, cudaReadModeElementType> tBucketOffsets;
 texture<unsigned int, 1, cudaReadModeElementType> tBucketOfSlices;
 texture<unsigned int, 1, cudaReadModeElementType> tSliceOffsetInBucket;
 
-__device__ int dGetValue(int2 v){
+static __device__ int dGetValue(int2 v){
 	return v.y;
 }
 
 template <typename T>
-__device__ T dGetValue(T v){
+static __device__ T dGetValue(T v){
 	return v;
 }
 
 
-__device__ void dPad(int2& v){
+static __device__ void dPad(int2& v){
 	v.x=0x3fffffff;
 	v.y=0x4fffffff;
 }
 
 template <typename T>
-__device__ void dPad(T & v){
+static __device__ void dPad(T & v){
 	v=0x7fffffff;
 }
 

platforms/cuda/src/kernels/kCalculateCDLJForces.cu

@@ -69,7 +69,7 @@ void GetCalculateCDLJForcesSim(gpuContext gpu)
 
 texture<float, 1, cudaReadModeElementType> tabulatedErfcRef;
 
-__device__ float fastErfc(float r)
+static __device__ float fastErfc(float r)
 {
     float normalized = cSim.tabulatedErfcScale*r;
     int index = (int) normalized;

platforms/cuda/src/kernels/kCalculateCDLJObcGbsaForces1.cu

@@ -66,7 +66,7 @@ void GetCalculateCDLJObcGbsaForces1Sim(gpuContext gpu)
 }
 texture<float, 1, cudaReadModeElementType> tabulatedErfcRef;
 
-__device__ float fastErfc(float r)
+static __device__ float fastErfc(float r)
 {
     float normalized = cSim.tabulatedErfcScale*r;
     int index = (int) normalized;

platforms/cuda/src/kernels/kCalculateGBVIAux.h

@@ -36,7 +36,7 @@
  * This file contains subroutines used in evaluating quantities associated w/ the GB/VI function
  */
 
-__device__ float getGBVI_L( float r, float x, float S )
+static __device__ float getGBVI_L( float r, float x, float S )
 {
 
    float rInv   = 1.0f/r;
@@ -48,7 +48,7 @@ __device__ float getGBVI_L( float r, float x, float S )
    return (1.5f*xInv2)*( (0.25f*rInv) - (xInv/3.0f) + (0.125f*diff2*xInv2*rInv) );
 }
 
-__device__ float getGBVI_Volume( float r_ij, float R, float S )
+static __device__ float getGBVI_Volume( float r_ij, float R, float S )
 {
 
      float upperBound        = r_ij + S; 
@@ -63,7 +63,7 @@ __device__ float getGBVI_Volume( float r_ij, float R, float S )
 
 }
 
-__device__ float getGBVI_dL_dr( float r, float x, float S )
+static __device__ float getGBVI_dL_dr( float r, float x, float S )
 {
 
    float rInv   = 1.0f/r;
@@ -80,7 +80,7 @@ __device__ float getGBVI_dL_dr( float r, float x, float S )
 
 }
 
-__device__ float getGBVI_dL_dx( float r, float x, float S )
+static __device__ float getGBVI_dL_dx( float r, float x, float S )
 {
 
    float rInv   = 1.0f/r;
@@ -96,7 +96,7 @@ __device__ float getGBVI_dL_dx( float r, float x, float S )
 
 }
 
-__device__ float getGBVI_dE2( float r, float R, float S, float bornForce )
+static __device__ float getGBVI_dE2( float r, float R, float S, float bornForce )
 {
 
     float diff              = S - R;
@@ -119,7 +119,7 @@ __device__ float getGBVI_dE2( float r, float R, float S, float bornForce )
 
 }
 
-__device__ float getGBVIBornForce2( float bornRadius, float R, float bornForce, float gamma )
+static __device__ float getGBVIBornForce2( float bornRadius, float R, float bornForce, float gamma )
 { 
     float ratio                     = (R/bornRadius);
     float returnBornForce           = bornForce + (3.0f*gamma*ratio*ratio*ratio)/bornRadius; // 'cavity' term

plugins/freeEnergy/platforms/cuda/src/kernels/kSoftcoreLJ.h

@@ -32,7 +32,7 @@
  */
 
 #ifdef USE_SOFTCORE_LJ
-__device__ float getSoftCoreLJ( float r2, float sig, float  eps, float lambdaI, float lambdaJ, float* energy)
+static __device__ float getSoftCoreLJ( float r2, float sig, float  eps, float lambdaI, float lambdaJ, float* energy)
 {
 
    float r                         = sqrt(r2);