Mods to get Free energy plugin libs to compile on Windows; still problem w/...

Mods to get Free energy plugin libs to compile on Windows; still problem w/ kCalculateLocalSoftcoreForces.cu
and hence OpenMMFreeEnergyCuda lib

platforms/cuda/src/kernels/gpu.cpp

@@ -57,6 +57,7 @@ using namespace std;
 // In case we're using some primitive version of Visual Studio this will
 // make sure that erf() and erfc() are defined.
 #include "openmm/internal/MSVC_erfc.h"
+#include "openmm/internal/windowsExport.h"
 
 using OpenMM::OpenMMException;
 using Lepton::Operation;
@@ -527,8 +528,7 @@ void gpuSetLJ14Parameters(gpuContext gpu, float epsfac, float fudge, const vecto
     psLJ14Parameter->Upload();
 }
 
-extern "C"
-void setExclusions(gpuContext gpu, const vector<vector<int> >& exclusions) {
+extern "C" void setExclusions(gpuContext gpu, const vector<vector<int> >& exclusions) {
     if (gpu->exclusions.size() > 0) {
         bool ok = (exclusions.size() == gpu->exclusions.size());
         for (int i = 0; i < exclusions.size() && ok; i++) {

platforms/cuda/src/kernels/kCalculateObcGbsaBornSum.cu

@@ -149,7 +149,7 @@ void kReduceObcGbsaBornSum(gpuContext gpu)
     LAUNCHERROR("kReduceObcGbsaBornSum");
 }
 
-void kClearObcGbsaBornSum(gpuContext gpu)
+extern void kClearObcGbsaBornSum(gpuContext gpu)
 {
   //  printf("kClearObcGbsaBornSum\n");
     kClearObcGbsaBornSum_kernel<<<gpu->sim.blocks, 384>>>();

platforms/cuda/src/kernels/windowsStdInt.h

+//
+// http://msinttypes.googlecode.com/svn/trunk/stdint.h
+//
 // ISO C9x  compliant stdint.h for Microsoft Visual Studio
 // Based on ISO/IEC 9899:TC2 Committee draft (May 6, 2005) WG14/N1124 
 // 

platforms/reference/src/SimTKReference/ReferenceBondForce.h

@@ -30,7 +30,7 @@
 
 // ---------------------------------------------------------------------------------------
 
-class ReferenceBondForce : public ReferenceForce {
+class OPENMM_EXPORT ReferenceBondForce : public ReferenceForce {
 
    private:
 

platforms/reference/src/SimTKReference/ReferenceBondIxn.h

@@ -29,7 +29,7 @@
 
 // ---------------------------------------------------------------------------------------
 
-class ReferenceBondIxn {
+class OPENMM_EXPORT ReferenceBondIxn {
 
    private:
 

platforms/reference/src/SimTKReference/ReferenceForce.h

@@ -25,9 +25,11 @@
 #ifndef __ReferenceForce_H__
 #define __ReferenceForce_H__
 
+#include "../../../../openmmapi/include/openmm/internal/windowsExport.h"
+
 // ---------------------------------------------------------------------------------------
 
-class ReferenceForce {
+class OPENMM_EXPORT  ReferenceForce {
 
    private:
        

platforms/reference/src/SimTKReference/ReferencePairIxn.h

@@ -25,11 +25,11 @@
 #ifndef __ReferencePairIxn_H__
 #define __ReferencePairIxn_H__
 
-// #include "ReferenceIxn.h"
+#include "../../../../openmmapi/include/openmm/internal/windowsExport.h"
 
 // ---------------------------------------------------------------------------------------
 
-class ReferencePairIxn {
+class OPENMM_EXPORT ReferencePairIxn {
 
    private:
 

platforms/reference/src/gbsa/CpuGBVI.cpp

@@ -501,7 +501,7 @@ RealOpenMM e3 = -partialChargeI2*partialCharges[atomJ]*Sgb( t )/deltaR[Reference
 
       energy += two*partialChargeI*atomIEnergy;
    }
-   energy *= CAL_TO_JOULE*preFactor;
+   energy *= 0.4184f*preFactor;
    energy -= cavityEnergy;
 
 #if( GBVIDebug == 1 )
@@ -717,7 +717,7 @@ if( atomI == 0 ){
 #endif
 
    const RealOpenMM* scaledRadii         = gbviParameters->getScaledRadii();
-   RealOpenMM stupidFactor               = three/CAL_TO_JOULE;
+   RealOpenMM stupidFactor               = three/0.4184f;
    for( int atomI = 0; atomI < numberOfAtoms; atomI++ ){
  
       RealOpenMM R        = atomicRadii[atomI];
@@ -830,7 +830,7 @@ if( atomI == 0 ){
 
    // convert from cal to Joule & apply prefactor tau = (1/diel_solute - 1/diel_solvent)
 
-   RealOpenMM conversion = (RealOpenMM)(CAL_TO_JOULE*gbviParameters->getTau());  
+   RealOpenMM conversion = (RealOpenMM)(0.4184f*gbviParameters->getTau());  
    for( int atomI = 0; atomI < numberOfAtoms; atomI++ ){
       inputForces[atomI][0] += conversion*forces[atomI][0];
       inputForces[atomI][1] += conversion*forces[atomI][1];

platforms/reference/src/gbsa/CpuImplicitSolvent.h

@@ -29,7 +29,7 @@
 
 // ---------------------------------------------------------------------------------------
 
-class CpuImplicitSolvent {
+class OPENMM_EXPORT CpuImplicitSolvent {
 
    public:
 

platforms/reference/src/gbsa/ImplicitSolventParameters.h

@@ -26,11 +26,12 @@
 #define __ImplicitSolventParameters_H__
 
 #include "../SimTKUtilities/SimTKOpenMMRealType.h"
+#include "../../../../openmmapi/include/openmm/internal/windowsExport.h"
 #include <string>
 
 // ---------------------------------------------------------------------------------------
 
-class ImplicitSolventParameters {
+class OPENMM_EXPORT ImplicitSolventParameters {
 
    protected:
 

plugins/freeEnergy/platforms/cuda/sharedTarget/CMakeLists.txt

@@ -72,6 +72,8 @@ CUDA_INCLUDE_DIRECTORIES(${OPENMM_BUILD_FREE_ENERGY_PATH}/platforms/cuda/../src
 CUDA_ADD_LIBRARY(${SHARED_TARGET} SHARED ${SOURCE_FILES} ${SOURCE_INCLUDE_FILES} ${API_ABS_INCLUDE_FILES})
 
 TARGET_LINK_LIBRARIES(${SHARED_TARGET} debug ${OPENMM_LIBRARY_NAME}_d optimized ${OPENMM_LIBRARY_NAME} )
+TARGET_LINK_LIBRARIES(${SHARED_TARGET} debug ${OPENMM_LIBRARY_NAME}Cuda_d optimized ${OPENMM_LIBRARY_NAME}Cuda )
+TARGET_LINK_LIBRARIES(${SHARED_TARGET} debug ${OPENMM_FREE_ENERGY_LIBRARY_NAME}_d optimized ${OPENMM_FREE_ENERGY_LIBRARY_NAME} )
 SET_TARGET_PROPERTIES(${SHARED_TARGET} PROPERTIES COMPILE_FLAGS "-DOPENMM_BUILDING_SHARED_LIBRARY")
 
 INSTALL_TARGETS(/lib/plugins RUNTIME_DIRECTORY /lib/plugins ${SHARED_TARGET})

plugins/freeEnergy/platforms/cuda/src/CudaFreeEnergyKernels.cpp

@@ -278,7 +278,6 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::initialize(const System& sy
         std::vector<float> softcoreLJLambdaArray(numParticles);
         std::vector<char> symbol;
         std::vector<std::vector<int> > exclusionList(numParticles);
-        float minSoftcoreLJLambda = 1.0e+20f;
         for (int i = 0; i < numParticles; i++) {
             double charge, radius, depth, softcoreLJLambda;
             force.getParticleParameters(i, charge, radius, depth, softcoreLJLambda);
@@ -287,9 +286,6 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::initialize(const System& sy
             c6[i]                    = static_cast<float>( (4*depth*pow(radius, 6.0)) );
             c12[i]                   = static_cast<float>( (4*depth*pow(radius, 12.0)) );
             softcoreLJLambdaArray[i] = static_cast<float>( softcoreLJLambda );
-            if( minSoftcoreLJLambda > softcoreLJLambda ){
-               minSoftcoreLJLambda = softcoreLJLambda;
-            }
             exclusionList[i].push_back(i);
         }
 
@@ -299,10 +295,11 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::initialize(const System& sy
         }
         Vec3 boxVectors[3];
         system.getPeriodicBoxVectors(boxVectors[0], boxVectors[1], boxVectors[2]);
-        gpuSetPeriodicBoxSize(gpu, static_cast<float>(boxVectors[0][0] ), static_cast<float>(boxVectors[1][1] ), static_cast<float>(boxVectors[2][2] ));
+        //gpuSetPeriodicBoxSize(gpu, static_cast<float>(boxVectors[0][0] ), static_cast<float>(boxVectors[1][1] ), static_cast<float>(boxVectors[2][2] ));
         CudaNonbondedMethod method = NO_CUTOFF;
         if (force.getNonbondedMethod() != NonbondedSoftcoreForce::NoCutoff) {
-            gpuSetNonbondedCutoff(gpu, static_cast<float>(force.getCutoffDistance() ), force.getReactionFieldDielectric());
+            throw OpenMMException( "NonbondedSoftcoreForce currently only handles NoCutoff option." );
+            //gpuSetNonbondedCutoff(gpu, static_cast<float>(force.getCutoffDistance() ), force.getReactionFieldDielectric());
             method = CUTOFF;
         }
         if (force.getNonbondedMethod() == NonbondedSoftcoreForce::CutoffPeriodic) {
@@ -325,7 +322,7 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::initialize(const System& sy
                     kmaxy++;
                 if (kmaxz%2 == 0)
                     kmaxz++;
-                gpuSetEwaldParameters(gpu, static_cast<float>( alpha ), kmaxx, kmaxy, kmaxz);
+                //gpuSetEwaldParameters(gpu, static_cast<float>( alpha ), kmaxx, kmaxy, kmaxz);
                 method = EWALD;
             }
             else {
@@ -335,7 +332,7 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::initialize(const System& sy
                 gridSizeX = ((gridSizeX+3)/4)*4;
                 gridSizeY = ((gridSizeY+3)/4)*4;
                 gridSizeZ = ((gridSizeZ+3)/4)*4;
-                gpuSetPMEParameters(gpu, static_cast<float>( alpha ), gridSizeX, gridSizeY, gridSizeZ);
+                //gpuSetPMEParameters(gpu, static_cast<float>( alpha ), gridSizeX, gridSizeY, gridSizeZ);
                 method = PARTICLE_MESH_EWALD;
             }
         }
@@ -454,16 +451,16 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::setIncludeGBVI( bool inputI
     bIncludeGBVI = inputIncludeGBVI;
 }
 
-bool CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::getIncludeSoftcore( void ) const {
-    return bIncludeSoftcore;
+int CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::getIncludeSoftcore( void ) const {
+    return includeSoftcore;
 }
 
 int CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::getNumExceptions( void ) const {
     return numExceptions;
 }
 
-void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::setIncludeSoftcore( bool inputIncludeSoftcore ){
-    bIncludeSoftcore = inputIncludeSoftcore;
+void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::setIncludeSoftcore( int inputIncludeSoftcore ){
+    includeSoftcore = inputIncludeSoftcore;
 }
 
 GpuLJ14Softcore* CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::getGpuLJ14Softcore( void ) const {
@@ -558,9 +555,9 @@ void CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::executeForces(ContextImpl& co
         }
     }
 
-    kClearBornForces(gpu);
+    kClearSoftcoreBornForces(gpu);
     kCalculateObcGbsaSoftcoreBornSum(gpu);
-    kReduceObcGbsaBornSum(gpu);
+    kReduceObcGbsaSoftcoreBornSum(gpu);
     kCalculateCDLJObcGbsaSoftcoreForces1(gpu);
 
 //kPrintForces(gpu, "Post kCalculateCDLJObcGbsaSoftcoreForces1", call );
@@ -703,13 +700,13 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
         }
     }
 
-    kClearBornForces(gpu);
+    kClearSoftcoreBornForces(gpu);
     kCalculateGBVISoftcoreBornSum(gpu);
 
     if( getQuinticScaling() ){
         kReduceGBVIBornSumQuinticScaling(gpu, gpuGBVISoftcore );
     } else {
-        kReduceGBVIBornSum(gpu);
+        kReduceGBVISoftcoreBornSum(gpu);
     }
     kCalculateCDLJObcGbsaSoftcoreForces1(gpu);
 
@@ -727,7 +724,7 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
         kReduceGBVIBornForcesQuinticScaling(gpu);
     } else {
         gpu->bIncludeGBVI = true;
-        kReduceObcGbsaBornForces(gpu);
+        kReduceObcGbsaSoftcoreBornForces(gpu);
         gpu->bIncludeGBVI = false;
     }
 

plugins/freeEnergy/platforms/cuda/src/CudaFreeEnergyKernels.h

@@ -61,7 +61,7 @@ public:
         numParticles         = 0;
         bIncludeGBSA         = false;
         bIncludeGBVI         = false;
-        bIncludeSoftcore     = false;
+        includeSoftcore      = false;
     }
 
     ~CudaFreeEnergyCalcNonbondedSoftcoreForceKernel();
@@ -114,13 +114,13 @@ public:
      *
      * @return flag
      */
-    bool getIncludeSoftcore( void ) const;
+    int getIncludeSoftcore( void ) const;
     /**
      * Set flag signalling whether GB/VI force is included
      *
      * @param inputIncludeGBVI input flag value
      */
-    void setIncludeSoftcore( bool inputSoftcore);
+    void setIncludeSoftcore( int inputSoftcore);
     /**
      * Get number of exceptions
      *
@@ -147,7 +147,7 @@ private:
     GpuLJ14Softcore* gpuLJ14Softcore;
     bool bIncludeGBSA;
     bool bIncludeGBVI;
-    bool bIncludeSoftcore;
+    int includeSoftcore;
     int numExceptions;
     FILE* log;
     int setSim;

plugins/freeEnergy/platforms/cuda/src/CudaFreeEnergyPlatform.cpp

@@ -27,7 +27,7 @@
 #include "CudaFreeEnergyPlatform.h"
 #include "CudaFreeEnergyKernelFactory.h"
 #include "openmm/freeEnergyKernels.h"
-#include "kernels/gputypes.h"
+#include "kernels/GpuFreeEnergyCudaKernels.h"
 
 using namespace OpenMM;
 
@@ -50,7 +50,7 @@ using namespace OpenMM;
 extern "C" void initOpenMMCudaFreeEnergyPlugin() {
 
 //(void) fprintf( stderr, "initOpenMMCudaFreeEnergyPlugin called\n");
-    if ( gpuIsAvailable() ){
+    if ( gpuIsAvailableSoftcore() ){
         Platform::registerPlatform(new CudaFreeEnergyPlatform());
     }   
 }

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

@@ -42,6 +42,9 @@
 // setup methods called from CudaFreeEnergyKernels
 // nonbonded and 1-4 ixns
 
+extern "C"
+bool gpuIsAvailableSoftcore();
+
 extern "C" 
 GpuNonbondedSoftcore* gpuSetNonbondedSoftcoreParameters(gpuContext gpu, float epsfac, const std::vector<int>& atom, const std::vector<float>& c6,
                                                         const std::vector<float>& c12, const std::vector<float>& q,
@@ -76,8 +79,6 @@ void kCalculateCDLJSoftcoreForces(gpuContext gpu );
 
 extern void kCalculateLocalSoftcoreForces( gpuContext gpu );
 
-
-
 // GB/VI softcore
 
 // setup method called from CudaFreeEnergyKernels
@@ -115,6 +116,10 @@ extern void kReduceGBVIBornSumQuinticScaling( gpuContext gpu, GpuGBVISoftcore* g
 extern void kCalculateGBVISoftcoreBornSum( gpuContext gpu );
 extern void kReduceGBVIBornForcesQuinticScaling( gpuContext gpu );
 extern void kCalculateGBVISoftcoreForces2( gpuContext gpu );
+extern void kReduceGBVISoftcoreBornForces(gpuContext gpu);
+extern void kReduceGBVISoftcoreBornSum(gpuContext gpu);
+
+extern void kClearSoftcoreBornForces(gpuContext gpu);
 
 // Obc softcore
 
@@ -158,8 +163,11 @@ void SetCalculateObcGbsaSoftcoreForces2Sim( gpuContext gpu );
 
 // kernel calls to device
 
+extern void kClearObcGbsaSoftcoreBornSum( gpuContext gpu );
 extern void kReduceObcGbsaSoftcoreBornForces( gpuContext gpu );
 extern void kCalculateObcGbsaSoftcoreBornSum( gpuContext gpu );
+extern void kReduceObcGbsaSoftcoreBornSum( gpuContext gpu );
+
 
 // this method is not needed; the OpenMM version can be used
 extern void kCalculateObcGbsaSoftcoreForces2( gpuContext gpu );

plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateGBVISoftcoreBornSum.cu

@@ -190,6 +190,123 @@ void kClearGBVISoftcoreBornSum(gpuContext gpu) {
     kClearGBVISoftcoreBornSum_kernel<<<gpu->sim.blocks, 384>>>();
 }
 
+__global__ void kReduceGBVISoftcoreBornForces_kernel()
+{
+    unsigned int pos = (blockIdx.x * blockDim.x + threadIdx.x);
+    float energy = 0.0f;
+    while (pos < cSim.atoms)
+    {
+        float bornRadius  = cSim.pBornRadii[pos];
+        float4 gbviData   = cSim.pGBVIData[pos];
+        float totalForce  = 0.0f;
+        float* pFt        = cSim.pBornForce + pos;
+
+        int i = cSim.nonbondOutputBuffers;
+        while (i >= 4)
+        {
+            float f1    = *pFt;
+            pFt        += cSim.stride;
+            float f2    = *pFt;
+            pFt        += cSim.stride;
+            float f3    = *pFt;
+            pFt        += cSim.stride;
+            float f4    = *pFt;
+            pFt        += cSim.stride;
+            totalForce += f1 + f2 + f3 + f4;
+            i -= 4;
+        }
+        if (i >= 2)
+        {
+            float f1    = *pFt;
+            pFt        += cSim.stride;
+            float f2    = *pFt;
+            pFt        += cSim.stride;
+            totalForce += f1 + f2;
+            i -= 2;
+        }
+        if (i > 0)
+        {
+            totalForce += *pFt;
+        }
+
+        float ratio         = (gbviData.x/bornRadius);
+        float ratio3        = ratio*ratio*ratio;
+        energy             -= gbviData.z*ratio3;
+        totalForce         += (3.0f*gbviData.z*ratio3)/bornRadius; // 'cavity' term
+        float br2           = bornRadius*bornRadius;
+        totalForce         *= (1.0f/3.0f)*br2*br2;
+
+        pFt = cSim.pBornForce + pos;
+        *pFt = totalForce;
+        pos += gridDim.x * blockDim.x;
+    }
+    cSim.pEnergy[blockIdx.x * blockDim.x + threadIdx.x] += energy;
+}
+
+void kReduceGBVISoftcoreBornForces(gpuContext gpu)
+{
+    kReduceGBVISoftcoreBornForces_kernel<<<gpu->sim.blocks, gpu->sim.bf_reduce_threads_per_block>>>();
+    LAUNCHERROR("kReduceGBVISoftcoreBornForces");
+
+}
+
+__global__ void kReduceGBVISoftcoreBornSum_kernel()
+{
+    unsigned int pos = (blockIdx.x * blockDim.x + threadIdx.x);
+    
+    while (pos < cSim.atoms)
+    {
+        float sum = 0.0f;
+        float* pSt = cSim.pBornSum + pos;
+        float4 atom = cSim.pGBVIData[pos];
+        
+        // Get summed Born data
+        for (int i = 0; i < cSim.nonbondOutputBuffers; i++)
+        {
+            sum += *pSt;
+       //     printf("%4d %4d A: %9.4f\n", pos, i, *pSt);
+            pSt += cSim.stride;
+        }
+        
+        // Now calculate Born radius
+
+        float Rinv           = 1.0f/atom.x;
+        sum                  = Rinv*Rinv*Rinv - sum; 
+        cSim.pBornRadii[pos] = pow( sum, (-1.0f/3.0f) ); 
+        pos += gridDim.x * blockDim.x;
+    }   
+}
+
+void kReduceGBVISoftcoreBornSum(gpuContext gpu)
+{
+    //printf("kReduceGBVISoftcoreBornSum\n");
+#define GBVISoftcore_DEBUG 0
+#if ( GBVISoftcore_DEBUG == 1 )
+               gpu->psGBVISoftcoreData->Download();
+               gpu->psBornSum->Download();
+               gpu->psPosq4->Download();
+                (void) fprintf( stderr, "\nkReduceGBVISoftcoreBornSum: Post BornSum %s Born radii & params\n", 
+                               (gpu->bIncludeGBVISoftcore ? "GBVI" : "Obc") );
+                for( int ii = 0; ii < gpu->natoms; ii++ ){
+                   (void) fprintf( stderr, "%d bSum=%14.6e param[%14.6e %14.6e %14.6e] x[%14.6f %14.6f %14.6f %14.6f]\n",
+                                   ii, 
+                                   gpu->psBornSum->_pSysStream[0][ii],
+                                   gpu->psGBVISoftcoreData->_pSysStream[0][ii].x,
+                                   gpu->psGBVISoftcoreData->_pSysStream[0][ii].y,
+                                   gpu->psGBVISoftcoreData->_pSysStream[0][ii].z,
+                                   gpu->psPosq4->_pSysStream[0][ii].x, gpu->psPosq4->_pSysStream[0][ii].y,
+                                   gpu->psPosq4->_pSysStream[0][ii].z, gpu->psPosq4->_pSysStream[0][ii].w
+                                 );  
+                }   
+#endif
+#undef GBVISoftcore_DEBUG
+
+
+    kReduceGBVISoftcoreBornSum_kernel<<<gpu->sim.blocks, 384>>>();
+    gpu->bRecalculateBornRadii = false;
+    LAUNCHERROR("kReduceGBVISoftcoreBornSum");
+}
+
 // Include versions of the kernels for N^2 calculations.
 
 #define METHOD_NAME(a, b) a##N2##b

plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateLocalSoftcoreForces.cu

@@ -24,12 +24,9 @@
  * along with this program.  If not, see <http://www.gnu.org/licenses/>.      *
  * -------------------------------------------------------------------------- */
 
-#include "gputypes.h"
 #include "GpuLJ14Softcore.h"
-
 #include <cuda.h>
 
-extern __shared__ Vectors sV[];
 static __constant__ cudaGmxSimulation cSim;
 static __constant__ cudaFreeEnergySimulationNonbonded14 feSim;
 
@@ -98,21 +95,18 @@ static __constant__ cudaFreeEnergySimulationNonbonded14 feSim;
     angle = (dp >= 0) ? angle : -angle; \
 }
 
-extern "C"
 void SetCalculateLocalSoftcoreGpuSim(gpuContext gpu)
 {
-    //(void) fprintf( stderr, "SetCalculateLocalSoftcoreForcesSim called\n" );
     cudaError_t status;
     status = cudaMemcpyToSymbol(cSim, &gpu->sim, sizeof(cudaGmxSimulation));     
     RTERROR(status, "cudaMemcpyToSymbol: SetCalculateLocalSoftcoreForcesSim copy to cSim failed");
 
 }
 
-static void SetCalculateLocalSoftcoreSim( GpuLJ14Softcore* gpuLJ14Softcore)
+void SetCalculateLocalSoftcoreSim( GpuLJ14Softcore* gpuLJ14Softcore)
 {
     cudaError_t status;
 
-    //(void) fprintf( stderr, "SetCalculateLocalSoftcoreSim called\n" );
     status = cudaMemcpyToSymbol(feSim, &gpuLJ14Softcore->feSim, sizeof(cudaFreeEnergySimulationNonbonded14));     
     RTERROR(status, "cudaMemcpyToSymbol: SetCalculateLocalSoftcoreSim copy to cSim failed");
 }
@@ -121,7 +115,7 @@ void GetCalculateLocalSoftcoreForcesSim(gpuContext gpu)
 {
     cudaError_t status;
     status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation));     
-    RTERROR(status, "cudaMemcpyFromSymbol: SetSim copy from cSim failed");
+    RTERROR(status, "cudaMemcpyFromSymbol: GetCalculateLocalSoftcoreForcesSim copy from cSim failed");
 }
     
 #define USE_SOFTCORE_LJ
@@ -129,7 +123,6 @@ void GetCalculateLocalSoftcoreForcesSim(gpuContext gpu)
 #include "kSoftcoreLJ.h"
 #endif
 
-//__global__ void METHOD_NAME(kCalculateLocalSoftcore, Forces_kernel)()
 __global__ void kCalculateLocalSoftcoreForces_kernel()
 {
     unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;

plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateNonbondedSoftcore.cu

@@ -26,6 +26,8 @@
 
 #include "GpuNonbondedSoftcore.h"
 #include "GpuFreeEnergyCudaKernels.h"
+#include "openmm/OpenMMException.h"
+#include <algorithm>
 
 // structure containing array of softcore lambdas
 
@@ -70,6 +72,24 @@ void GetCalculateCDLJSoftcoreForcesSim(float* gpuParticleSoftCoreLJLambda)
 // create, initialize and entrt SoftCoreLJLambda values
 // return handle to GpuNonbondedSoftcore object
 
+static void setSoftcoreExclusions(gpuContext gpu, const std::vector<std::vector<int> >& exclusions) {
+    if (gpu->exclusions.size() > 0) { 
+        bool ok = (exclusions.size() == gpu->exclusions.size());
+        for (unsigned int i = 0; i < exclusions.size() && ok; i++) {
+            if (exclusions[i].size() != gpu->exclusions[i].size())
+                ok = false;
+            else {
+                for (unsigned int j = 0; j < exclusions[i].size(); j++) 
+                    if (find(gpu->exclusions[i].begin(), gpu->exclusions[i].end(), exclusions[i][j]) == gpu->exclusions[i].end())
+                        ok = false;
+            }
+        }
+        if (!ok)
+            throw OpenMM::OpenMMException("All nonbonded forces must have identical sets of exceptions");
+    }    
+    gpu->exclusions = exclusions;
+}
+
 extern "C"
 GpuNonbondedSoftcore* gpuSetNonbondedSoftcoreParameters(gpuContext gpu, float epsfac, const std::vector<int>& atom, const std::vector<float>& c6,
                                                         const std::vector<float>& c12, const std::vector<float>& q,
@@ -80,7 +100,7 @@ GpuNonbondedSoftcore* gpuSetNonbondedSoftcoreParameters(gpuContext gpu, float ep
     gpu->sim.epsfac                    = epsfac;
     gpu->sim.nonbondedMethod           = method;
     if (numberOfParticles > 0)
-        setExclusions(gpu, exclusions);
+        setSoftcoreExclusions(gpu, exclusions);
     
     // create gpuNonbondedSoftcore
 
@@ -159,6 +179,14 @@ void gpuDeleteNonbondedSoftcoreParameters( void* gpuNonbondedSoftcore)
     delete internalGNonbondedSoftcore;
 }
 
+extern "C"
+bool gpuIsAvailableSoftcore()
+{
+    int deviceCount;
+    cudaGetDeviceCount(&deviceCount);
+    return (deviceCount > 0);
+}
+
 struct Atom {
     float x;
     float y;
@@ -399,7 +427,7 @@ fprintf( stderr, "kCalculateCDLJSoftcoreForces: bOutputBufferPerWarp=%u blks=%u
 void kPrintForces(gpuContext gpu, std::string idString, int call )
 {
  //   printf("kReduceForces\n");
-#define GBVI_DEBUG 4
+#define GBVI_DEBUG 0
 #if ( GBVI_DEBUG == 4 )
 
                 gpu->psBornRadii->Download();

plugins/freeEnergy/platforms/cuda/src/kernels/kCalculateObcGbsaSoftcoreBornSum.cu

@@ -75,6 +75,39 @@ void GetCalculateObcGbsaSoftcoreBornSumSim(gpuContext gpu)
     RTERROR(status, "GetCalculateObcGbsaSoftcoreBornSumSim: cudaMemcpyFromSymbol: SetSim copy from cSim failed");
 }
 
+__global__ void kClearObcGbsaSoftcoreBornSum_kernel()
+{
+    unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;
+    while (pos < cSim.stride * cSim.nonbondOutputBuffers)
+    {
+        ((float*)cSim.pBornSum)[pos] = 0.0f;
+        pos += gridDim.x * blockDim.x;
+    }
+}
+
+__global__ void kClearSoftcoreBornForces_kernel()
+{
+    unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;
+    while (pos < cSim.stride * cSim.nonbondOutputBuffers)
+    {   
+        ((float*)cSim.pBornForce)[pos] = 0.0f;
+        pos += gridDim.x * blockDim.x;
+    }   
+}
+
+void kClearSoftcoreBornForces(gpuContext gpu)
+{
+  //  printf("kClearBornForces\n");
+    kClearSoftcoreBornForces_kernel<<<gpu->sim.blocks, 384>>>();
+    LAUNCHERROR("kClearSoftcoreBornForces");
+}
+
+void kClearObcGbsaSoftcoreBornSum(gpuContext gpu)
+{
+  //  printf("kClearObcGbsaBornSum\n");
+    kClearObcGbsaSoftcoreBornSum_kernel<<<gpu->sim.blocks, 384>>>();
+}
+
 __global__ void kReduceObcGbsaSoftcoreBornForces_kernel()
 {
     unsigned int pos = (blockIdx.x * blockDim.x + threadIdx.x);
@@ -176,18 +209,7 @@ void kReduceObcGbsaSoftcoreBornForces(gpuContext gpu)
 #define METHOD_NAME(a, b) a##PeriodicByWarp##b
 #include "kCalculateObcGbsaSoftcoreBornSum.h"
 
-#if 0
-__global__ void kClearObcGbsaBornSum_kernel()
-{
-    unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;
-    while (pos < cSim.stride * cSim.nonbondOutputBuffers)
-    {
-        ((float*)cSim.pBornSum)[pos] = 0.0f;
-        pos += gridDim.x * blockDim.x;
-    }
-}
-
-__global__ void kReduceObcGbsaBornSum_kernel()
+__global__ void kReduceObcGbsaSoftcoreBornSum_kernel()
 {
     unsigned int pos = (blockIdx.x * blockDim.x + threadIdx.x);
     
@@ -221,14 +243,13 @@ __global__ void kReduceObcGbsaBornSum_kernel()
     }   
 }
 
-void kReduceObcGbsaBornSum(gpuContext gpu)
+void kReduceObcGbsaSoftcoreBornSum(gpuContext gpu)
 {
-//    printf("kReduceObcGbsaBornSum\n");
-    kReduceObcGbsaBornSum_kernel<<<gpu->sim.blocks, 384>>>();
+//    printf("kReduceObcGbsaSoftcoreBornSum\n");
+    kReduceObcGbsaSoftcoreBornSum_kernel<<<gpu->sim.blocks, 384>>>();
     gpu->bRecalculateBornRadii = false;
-    LAUNCHERROR("kReduceObcGbsaBornSum");
+    LAUNCHERROR("kReduceObcGbsaSoftcoreBornSum");
 }
-#endif
 
 /** 
  * Initialize parameters for Cuda Obc softcore
@@ -309,7 +330,7 @@ GpuObcGbsaSoftcore* gpuSetObcSoftcoreParameters(gpuContext gpu, float innerDiele
 void kCalculateObcGbsaSoftcoreBornSum(gpuContext gpu)
 {
   //  printf("kCalculateObcGbsaSoftcoreBornSum\n");
-    kClearObcGbsaBornSum( gpu );
+    kClearObcGbsaSoftcoreBornSum(gpu);
     LAUNCHERROR("kClearBornSum from kCalculateObcGbsaSoftcoreBornSum");
 
     switch (gpu->sim.nonbondedMethod)

plugins/freeEnergy/platforms/reference/src/SimTKReference/ReferenceFreeEnergyLJCoulomb14Softcore.cpp

@@ -281,7 +281,7 @@ int ReferenceFreeEnergyLJCoulomb14Softcore::calculateBondIxn( int* atomIndices,
          message << "]";
       }
 
-      SimTKOpenMMLog::printMessage( message );
+      //SimTKOpenMMLog::printMessage( message );
    }   
 
    return ReferenceForce::DefaultReturn;