Removed CalcAmoebaForcesAndEnergyKernel and kClearBornSum

plugins/amoeba/platforms/cuda/src/AmoebaCudaKernelFactory.cpp

@@ -56,7 +56,7 @@ extern "C" OPENMMCUDA_EXPORT void registerKernelFactories() {
              platform.registerKernelFactory(CalcAmoebaGeneralizedKirkwoodForceKernel::Name(), factory);
              platform.registerKernelFactory(CalcAmoebaVdwForceKernel::Name(), factory);
              platform.registerKernelFactory(CalcAmoebaWcaDispersionForceKernel::Name(), factory);
-             platform.registerKernelFactory(CalcAmoebaForcesAndEnergyKernel::Name(), factory);
+             //platform.registerKernelFactory(CalcAmoebaForcesAndEnergyKernel::Name(), factory);
         }
     }
 }
@@ -109,7 +109,7 @@ KernelImpl* AmoebaCudaKernelFactory::createKernelImpl(std::string name, const Pl
     if( mapIterator == contextToAmoebaDataMap.end() ){
         amoebaCudaData                         = new AmoebaCudaData( cudaPlatformData );
         contextToAmoebaDataMap[&context]       = amoebaCudaData;
-        //amoebaCudaData->setLog( stderr );
+        amoebaCudaData->setLog( stderr );
         amoebaCudaData->setContextImpl( static_cast<void*>(&context) );
     } else {
         amoebaCudaData                         = mapIterator->second;
@@ -154,8 +154,5 @@ KernelImpl* AmoebaCudaKernelFactory::createKernelImpl(std::string name, const Pl
     if (name == CalcAmoebaUreyBradleyForceKernel::Name())
         return new CudaCalcAmoebaUreyBradleyForceKernel(name, platform, *amoebaCudaData, context.getSystem());
 
-    if (name == CalcAmoebaForcesAndEnergyKernel::Name())
-        return new CalcAmoebaForcesAndEnergyKernel(name, platform, *amoebaCudaData);
-
     throw OpenMMException((std::string("Tried to create kernel with illegal kernel name '")+name+"'").c_str());
 }

plugins/amoeba/platforms/cuda/src/AmoebaCudaKernels.cpp

@@ -44,63 +44,6 @@ extern "C" int gpuSetConstants( gpuContext gpu );
 using namespace OpenMM;
 using namespace std;
 
-void CalcAmoebaForcesAndEnergyKernel::initialize(const System& system) {
-}
-
-void CalcAmoebaForcesAndEnergyKernel::beginComputation(ContextImpl& context, bool includeForces, bool includeEnergy) {
-//fprintf( stderr, "In CalcAmoebaForcesAndEnergyKernel::beginComputation computeForceCount=%d inbMethod=%d GBSA=%d includeForces=%d includeEnergy=%d\n", 
-//         data.cudaPlatformData.computeForceCount, data.cudaPlatformData.nonbondedMethod,  data.getHasAmoebaGeneralizedKirkwood(), includeForces, includeEnergy ); fflush( stderr );
-
-    amoebaGpuContext amoebaGpu  = data.getAmoebaGpu();
-    _gpuContext* gpu            = data.getAmoebaGpu()->gpuContext;
-/*
-    if (data.cudaPlatformData.nonbondedMethod != NO_CUTOFF && data.cudaPlatformData.computeForceCount%100 == 0){
-        //fprintf( stderr, "In CalcAmoebaForcesAndEnergyKernel::beginComputation reordering atoms\n" ); fflush( stderr );
-        gpuReorderAtoms(gpu);
-    }
-
-    data.cudaPlatformData.computeForceCount++;
-
-    if(includeForces){
-        if( data.getHasAmoebaGeneralizedKirkwood() ){
-            kClearBornSumAndForces(gpu);
-        } else {
-            kClearForces(gpu);
-        }
-    }
-
-    if (includeEnergy)
-        kClearEnergy(gpu);
-*/
-    int originalIncludeGBSA = gpu->bIncludeGBSA;
-    if(includeForces && data.getHasAmoebaGeneralizedKirkwood() ){
-        gpu->bIncludeGBSA = 1;
-    }
-    cudaCalcForcesAndEnergyKernel->beginComputation( context, 1, includeEnergy);
-    gpu->bIncludeGBSA = originalIncludeGBSA;
-
-}
-
-double CalcAmoebaForcesAndEnergyKernel::finishComputation(ContextImpl& context, bool includeForces, bool includeEnergy) {
-
-    amoebaGpuContext amoebaGpu  = data.getAmoebaGpu();
-    _gpuContext* gpu            = data.getAmoebaGpu()->gpuContext;
-
-    return cudaCalcForcesAndEnergyKernel->finishComputation( context, includeForces, includeEnergy);
-/*
-    if( includeForces ){
-        kReduceForces(gpu);
-    }
-
-    double energy = 0.0; 
-    if( includeEnergy ){
-        energy = kReduceEnergy(gpu);
-    }    
-    return energy;
-*/
-
-}
-
 /* -------------------------------------------------------------------------- *
  *                           Calculates bonded forces                         *
  * -------------------------------------------------------------------------- */
@@ -848,6 +791,7 @@ static void computeAmoebaMultipoleForce( AmoebaCudaData& data ) {
     // calculate Born radii
 
     if( data.getHasAmoebaGeneralizedKirkwood() ){
+        kClearBornSum( gpu->gpuContext );
         kCalculateObcGbsaBornSum(gpu->gpuContext);
         kReduceObcGbsaBornSum(gpu->gpuContext);
     }   

plugins/amoeba/platforms/cuda/src/AmoebaCudaKernels.h

@@ -35,54 +35,6 @@
 
 namespace OpenMM {
 
-/**
- * This kernel is invoked at the beginning and end of force and energy computations.  It gives the
- * Platform a chance to clear buffers and do other initialization at the beginning, and to do any
- * necessary work at the end to determine the final results.
- */
-class CalcAmoebaForcesAndEnergyKernel : public CalcForcesAndEnergyKernel {
-public:
-    CalcAmoebaForcesAndEnergyKernel(std::string name, const Platform& platform, AmoebaCudaData& data) : CalcForcesAndEnergyKernel(name, platform), data(data) {
-        cudaCalcForcesAndEnergyKernel = new CudaCalcForcesAndEnergyKernel( name, platform, data.cudaPlatformData );
-    }   
-    ~CalcAmoebaForcesAndEnergyKernel() {
-        delete cudaCalcForcesAndEnergyKernel;
-    }   
-    /*** 
-     * Initialize the kernel.
-     * 
-     * @param system     the System this kernel will be applied to
-     */
-    void initialize(const System& system);
-
-    /** 
-     * This is called at the beginning of each force/energy computation, before calcForcesAndEnergy() has been called on
-     * any ForceImpl.
-     *
-     * @param context       the context in which to execute this kernel
-     * @param includeForce  true if forces should be computed
-     * @param includeEnergy true if potential energy should be computed
-     */
-    void beginComputation(ContextImpl& context, bool includeForce, bool includeEnergy);
-
-    /** 
-     * This is called at the end of each force/energy computation, after calcForcesAndEnergy() has been called on
-     * every ForceImpl.
-     *
-     * @param context       the context in which to execute this kernel
-     * @param includeForce  true if forces should be computed
-     * @param includeEnergy true if potential energy should be computed
-     * @return the potential energy of the system.  This value is added to all values returned by ForceImpls'
-     * calcForcesAndEnergy() methods.  That is, each force kernel may <i>either</i> return its contribution to the
-     * energy directly, <i>or</i> add it to an internal buffer so that it will be included here.
-     */
-    double finishComputation(ContextImpl& context, bool includeForce, bool includeEnergy);
-
-private:
-    AmoebaCudaData& data;
-    CudaCalcForcesAndEnergyKernel* cudaCalcForcesAndEnergyKernel;
-};
-
 /**
  * This kernel is invoked by AmoebaHarmonicBondForce to calculate the forces acting on the system and the energy of the system.
  */

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

@@ -160,6 +160,7 @@ extern void kClearFloat( amoebaGpuContext amoebaGpu, unsigned int entries, CUDAS
 extern void kClearFloat4( amoebaGpuContext amoebaGpu, unsigned int entries, CUDAStream<float4>* fieldToClear );
 extern void kClearFields_1( amoebaGpuContext amoebaGpu );
 extern void kClearFields_3( amoebaGpuContext amoebaGpu, unsigned int numberToClear );
+extern void kClearBornSum(gpuContext gpu);
 extern unsigned int getThreadsPerBlock( amoebaGpuContext amoebaGpu, unsigned int sharedMemoryPerThread, unsigned int sharedMemoryPerBlock );
 
 //extern int isNanOrInfinity( double number );

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

@@ -64,6 +64,26 @@ void kClearFloat4_kernel( unsigned int bufferLength, float4* fieldToClear )
     }   
 }
 
+__global__ 
+__launch_bounds__(384, 1)
+void kClearBornSum_kernel()
+{
+    unsigned int pos = blockIdx.x * blockDim.x + threadIdx.x;
+    while (pos < cSim.stride * cSim.nonbondOutputBuffers)
+    {   
+        cSim.pBornSum[pos]   = 0.0f;
+        cSim.pBornForce[pos] = 0.0f;
+        pos += gridDim.x * blockDim.x;
+    }   
+}
+
+void kClearBornSum(gpuContext gpu)
+{
+  //  printf("kClearBornSum\n");
+      kClearBornSum_kernel<<<gpu->sim.blocks, 384>>>();
+      LAUNCHERROR("kClearBornSum");
+}
+
 void kClearFloat4( amoebaGpuContext amoebaGpu, unsigned int entries, CUDAStream<float4>* fieldToClear )
 {
     kClearFloat4_kernel<<<amoebaGpu->gpuContext->blocksPerSM, 384>>>( entries, fieldToClear->_pDevData );