Activated static tests

Fixed bug: was calling kReduceObcGbsaSoftcoreBornForces instead of kReduceGBVISoftcoreBornForces in GB/VI force calculation

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

@@ -398,8 +398,8 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::executeForces(ContextImpl&
                             methodName.c_str(), getIncludeGBSA(), getIncludeGBVI(), getNumExceptions() );
             (void) fflush( log );
         }
-        SetCalculateLocalSoftcoreGpuSim( gpu );
         SetCalculateCDLJSoftcoreGpuSim( gpu );
+        SetCalculateLocalSoftcoreGpuSim( gpu );
 
         // flip strides (unsure if this is needed)
 
@@ -701,6 +701,7 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
     }
 
     kClearSoftcoreBornForces(gpu);
+
     kCalculateGBVISoftcoreBornSum(gpu);
 
     if( getQuinticScaling() ){
@@ -708,9 +709,8 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
     } else {
         kReduceGBVISoftcoreBornSum(gpu);
     }
-    kCalculateCDLJObcGbsaSoftcoreForces1(gpu);
 
-//kPrintForces(gpu, "Post kCalculateCDLJObcGbsaSoftcoreForces1", call );
+    kCalculateCDLJObcGbsaSoftcoreForces1(gpu);
 
     if( debug && log ){
         (void) fprintf( log, "\n%s: calling %s\n", methodName.c_str(),
@@ -724,7 +724,7 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
         kReduceGBVIBornForcesQuinticScaling(gpu);
     } else {
         gpu->bIncludeGBVI = true;
-        kReduceObcGbsaSoftcoreBornForces(gpu);
+        kReduceGBVISoftcoreBornForces(gpu);
         gpu->bIncludeGBVI = false;
     }
 

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

@@ -34,6 +34,7 @@
 #include "GpuObcGbsaSoftcore.h"
 #include "GpuGBVISoftcore.h"
 #include <vector>
+#include <cuda.h>
 
 // Function prototypes
 

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

@@ -31,6 +31,7 @@
 
 #include "GpuGBVISoftcore.h"
 #include "GpuFreeEnergyCudaKernels.h"
+#include <cuda.h>
 
 struct cudaFreeEnergySimulationGBVI {
     float quinticLowerLimitFactor;

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

@@ -26,6 +26,7 @@
 
 #include "GpuLJ14Softcore.h"
 #include "GpuFreeEnergyCudaKernels.h"
+#include <cuda.h>
 
 static __constant__ cudaGmxSimulation cSim;
 static __constant__ cudaFreeEnergySimulationNonbonded14 feSim;

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

@@ -47,7 +47,7 @@ void SetCalculateCDLJSoftcoreGpuSim( gpuContext gpu )
 {
     cudaError_t status;
     status = cudaMemcpyToSymbol(cSim, &gpu->sim, sizeof(cudaGmxSimulation));
-    RTERROR(status, "cudaMemcpyToSymbol: SetSim copy to cSim failed");
+    RTERROR(status, "cudaMemcpyToSymbol: SetCalculateCDLJSoftcoreGpuSim copy to cSim failed");
 
     //(void) fprintf( stderr, "SetCalculateCDLJSoftcoreGpuSim gpu=%p cSim=%p sizeof=%u\n", gpu, &gpu->sim, sizeof(cudaGmxSimulation) ); fflush( stderr );
 }

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

@@ -187,6 +187,7 @@ void kReduceObcGbsaSoftcoreBornForces(gpuContext gpu)
 
 // Include versions of the kernels with cutoffs.
 
+#if 0
 #undef METHOD_NAME
 #undef USE_OUTPUT_BUFFER_PER_WARP
 #define USE_CUTOFF
@@ -208,6 +209,7 @@ void kReduceObcGbsaSoftcoreBornForces(gpuContext gpu)
 #undef METHOD_NAME
 #define METHOD_NAME(a, b) a##PeriodicByWarp##b
 #include "kCalculateObcGbsaSoftcoreBornSum.h"
+#endif
 
 __global__ void kReduceObcGbsaSoftcoreBornSum_kernel()
 {
@@ -345,12 +347,13 @@ fprintf( stderr, "kCalculateObcGbsaSoftcoreBornSum: bOutputBufferPerWarp=%u blks
 #undef GBSA
 
             if (gpu->bOutputBufferPerWarp)
-                kCalculateObcGbsaN2ByWarpBornSum_kernel<<<gpu->sim.nonbond_blocks, gpu->sim.nonbond_threads_per_block,
+                kCalculateObcGbsaSoftcoreN2ByWarpBornSum_kernel<<<gpu->sim.nonbond_blocks, gpu->sim.nonbond_threads_per_block,
                         sizeof(Atom)*gpu->sim.nonbond_threads_per_block>>>(gpu->sim.pWorkUnit);
             else
-                kCalculateObcGbsaN2BornSum_kernel<<<gpu->sim.nonbond_blocks, gpu->sim.nonbond_threads_per_block,
+                kCalculateObcGbsaSoftcoreN2BornSum_kernel<<<gpu->sim.nonbond_blocks, gpu->sim.nonbond_threads_per_block,
                         sizeof(Atom)*gpu->sim.nonbond_threads_per_block>>>(gpu->sim.pWorkUnit);
             break;
+#if 0
         case CUTOFF:
             if (gpu->bOutputBufferPerWarp)
                 kCalculateObcGbsaCutoffByWarpBornSum_kernel<<<gpu->sim.nonbond_blocks, gpu->sim.nonbond_threads_per_block,
@@ -367,6 +370,7 @@ fprintf( stderr, "kCalculateObcGbsaSoftcoreBornSum: bOutputBufferPerWarp=%u blks
                 kCalculateObcGbsaPeriodicBornSum_kernel<<<gpu->sim.nonbond_blocks, gpu->sim.nonbond_threads_per_block,
                         (sizeof(Atom)+sizeof(float))*gpu->sim.nonbond_threads_per_block>>>(gpu->sim.pInteractingWorkUnit);
             break;
+#endif
     }
     LAUNCHERROR("kCalculateObcGbsaSoftcoreBornSum");
 }

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

@@ -30,7 +30,7 @@
  * different versions of the kernels.
  */
 
-__global__ void METHOD_NAME(kCalculateObcGbsa, BornSum_kernel)(unsigned int* workUnit)
+__global__ void METHOD_NAME(kCalculateObcGbsaSoftcore, BornSum_kernel)(unsigned int* workUnit)
 {
     extern __shared__ Atom sA[];
     unsigned int totalWarps   = cSim.nonbond_blocks*cSim.nonbond_threads_per_block/GRID;

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

@@ -10,13 +10,14 @@ INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/cuda/include)
 INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/cuda/src)
 INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/cuda/src/kernels)
 Set( SHARED_OPENMM_TARGET OpenMMFreeEnergy)
+Set( STATIC_OPENMM_TARGET OpenMMFreeEnergy_static)
 Set( SHARED_CUDA_TARGET OpenMMCuda)
 Set( STATIC_CUDA_TARGET OpenMMCuda_static)
 IF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
     SET(SHARED_CUDA_TARGET   ${SHARED_CUDA_TARGET}_d)
     SET(SHARED_OPENMM_TARGET ${SHARED_OPENMM_TARGET}_d)
     SET(STATIC_CUDA_TARGET ${STATIC_CUDA_TARGET}_d)
-    SET(STATIC_CUDA_TARGET ${STATIC_CUDA_TARGET}_d)
+    Set(STATIC_OPENMM_TARGET ${STATIC_OPENMM_TARGET}_d)
 ENDIF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
 #LINK_DIRECTORIES
 
@@ -34,14 +35,14 @@ FOREACH(TEST_PROG ${TEST_PROGS})
 
 
     # Link with static library
-#     SET(TEST_STATIC ${TEST_ROOT}Static)
-#     CUDA_ADD_EXECUTABLE(${TEST_STATIC} ${TEST_PROG})
-#     SET_TARGET_PROPERTIES(${TEST_STATIC}
-#                 PROPERTIES
-#                 COMPILE_FLAGS "-DOPENMM_USE_STATIC_LIBRARIES"
-#                 )
-#     TARGET_LINK_LIBRARIES(${TEST_STATIC} ${STATIC_TARGET})
-#     ADD_TEST(${TEST_STATIC} ${EXECUTABLE_OUTPUT_PATH}/${TEST_STATIC})
+     SET(TEST_STATIC ${TEST_ROOT}Static)
+     CUDA_ADD_EXECUTABLE(${TEST_STATIC} ${TEST_PROG})
+     SET_TARGET_PROPERTIES(${TEST_STATIC}
+                 PROPERTIES
+                 COMPILE_FLAGS "-DOPENMM_USE_STATIC_LIBRARIES"
+                 )
+     TARGET_LINK_LIBRARIES(${TEST_STATIC} ${STATIC_TARGET} ${STATIC_OPENMM_TARGET} ${STATIC_CUDA_TARGET})
+     ADD_TEST(${TEST_STATIC} ${EXECUTABLE_OUTPUT_PATH}/${TEST_STATIC})
 
 ENDFOREACH(TEST_PROG ${TEST_PROGS})