Cleanup of AmoebaTinkerParameterFile

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

@@ -657,10 +657,13 @@ CudaCalcAmoebaGeneralizedKirkwoodForceKernel::~CudaCalcAmoebaGeneralizedKirkwood
 void CudaCalcAmoebaGeneralizedKirkwoodForceKernel::initialize(const System& system, const AmoebaGeneralizedKirkwoodForce& force) {
 
     data.setHasAmoebaGeneralizedKirkwood( true );
+
     int numParticles = system.getNumParticles();
+
     std::vector<float> radius(numParticles);
     std::vector<float> scale(numParticles);
     std::vector<float> charge(numParticles);
+
     for( int ii = 0; ii < numParticles; ii++ ){
         double particleCharge, particleRadius, scalingFactor;
         force.getParticleParameters(ii, particleCharge, particleRadius, scalingFactor);

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

@@ -2925,6 +2925,10 @@ void amoebaGpuBuildScalingList( amoebaGpuContext amoebaGpu )
 
     // ---------------------------------------------------------------------------------------
 
+    if( amoebaGpu->psCovalentDegree == NULL ){
+        return;
+    }    
+
     const unsigned int paddedAtoms     = amoebaGpu->paddedNumberOfAtoms;
     const unsigned int actualAtoms     = amoebaGpu->gpuContext->natoms;
     const unsigned int grid            = amoebaGpu->gpuContext->grid;
@@ -4050,4 +4054,99 @@ void readFile( std::string fileName, StringVectorVector& fileContents ){
     return;
 }
 
+/**---------------------------------------------------------------------------------------
+
+   Report whether a number is a nan or infinity
+
+   @param number               number to test
+   @return 1 if number is  nan or infinity; else return 0
+
+   --------------------------------------------------------------------------------------- */
+
+int isNanOrInfinity( double number ){
+    return (number != number || number == std::numeric_limits<double>::infinity() || number == -std::numeric_limits<double>::infinity()) ? 1 : 0; 
+}
+
+/**---------------------------------------------------------------------------------------
+
+   Track iterations for MI dipoles
+
+   @param amoebaGpu            amoebaGpuContext reference
+   @param iteration            MI iteration
+
+   --------------------------------------------------------------------------------------- */
+
+void trackMutualInducedIterations( amoebaGpuContext amoebaGpu, int iteration){
+
+// ---------------------------------------------------------------------------------------
+
+    static const std::string methodName      = "trackMutualInducedIterations";
+    static int currentStep                   = 0; 
+    static double iterationStat[6]           = { 0.0, 0.0, 1000.0, 0.0, 0.0, 0.0 };
+
+// ---------------------------------------------------------------------------------------
+
+    if( amoebaGpu->log == NULL || currentStep > 20000 )return;
+
+    gpuContext gpu                       = amoebaGpu->gpuContext;
+    currentStep++;
+
+    double interationD = static_cast<double>(iteration);
+    iterationStat[0]  += interationD;
+    iterationStat[1]  += interationD*interationD;
+    iterationStat[2]   = interationD < iterationStat[2] ? interationD : iterationStat[2];
+    iterationStat[3]   = interationD > iterationStat[3] ? interationD : iterationStat[3];
+    iterationStat[4]  += 1.0; 
+    if( iterationStat[4] >= 1000.0 ){
+        double average = iterationStat[0]/iterationStat[4]; 
+        double stddev  = iterationStat[1] - average*average*iterationStat[4]; 
+               stddev  = sqrt( stddev )/(iterationStat[4]-1.0);
+        (void) fprintf( amoebaGpu->log, "%s %8d iteration=%10.3f stddev=%10.3f min/max[%10.3f %10.3f] %10.1f eps=%14.7e\n",
+                        methodName.c_str(), currentStep, average, stddev, iterationStat[2], iterationStat[3], iterationStat[4], amoebaGpu->mutualInducedCurrentEpsilon );
+        (void) fflush( amoebaGpu->log );
+        iterationStat[0] = iterationStat[1] = iterationStat[4] = 0.0; 
+    }
+    if( 0 ){ 
+        std::vector<int> fileId;
+        if( interationD < (amoebaGpu->mutualInducedMaxIterations-10) ) {
+            int id = (currentStep % 20); 
+            fileId.push_back( id );
+        } else {
+            fileId.push_back( currentStep );
+        }    
+        if( (currentStep % 20) == 0 || fileId[0] > 20 ){
+             (void) fprintf( amoebaGpu->log, "step=%d fileId=%d\n", currentStep, fileId[0] );
+        }
+        (void) fflush( amoebaGpu->log );
+        VectorOfDoubleVectors outputVector;
+        cudaLoadCudaFloat4Array( gpu->natoms, 3, gpu->psPosq4,                    outputVector );
+        cudaLoadCudaFloat4Array( gpu->natoms, 3, gpu->psVelm4,                    outputVector );
+/*
+            cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipole,      outputVector );
+            cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipolePolar, outputVector );
+            cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipoleS,     outputVector );
+            cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipolePolarS,outputVector );
+*/
+        cudaWriteVectorOfDoubleVectorsToFile( "CudaMIT", fileId, outputVector );
+        int nansPresent = isNanOrInfinity( amoebaGpu->mutualInducedCurrentEpsilon );
+        if( nansPresent == 0 ){ 
+            for( int ii = 0; ii < gpu->natoms && nansPresent == 0; ii++ ){
+                if( isNanOrInfinity( gpu->psPosq4->_pSysStream[0][ii].x ) || 
+                    isNanOrInfinity( gpu->psPosq4->_pSysStream[0][ii].y ) || 
+                    isNanOrInfinity( gpu->psPosq4->_pSysStream[0][ii].z ) || 
+                    isNanOrInfinity( gpu->psVelm4->_pSysStream[0][ii].x  ) || 
+                    isNanOrInfinity( gpu->psVelm4->_pSysStream[0][ii].y  ) || 
+                    isNanOrInfinity( gpu->psVelm4->_pSysStream[0][ii].z  ) ){ 
+                    nansPresent = 1; 
+                }
+            }
+        }
+        if( nansPresent ){
+            (void) fprintf( amoebaGpu->log, "epsilon nan - exiting\n" );
+            (void) fflush( amoebaGpu->log );
+            exit(-1);
+        }
+    }
+}
+
 #undef  AMOEBA_DEBUG

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

@@ -149,5 +149,8 @@ extern void kClearFields_1( amoebaGpuContext amoebaGpu );
 extern void kClearFields_3( amoebaGpuContext amoebaGpu, unsigned int numberToClear );
 extern unsigned int getThreadsPerBlock( amoebaGpuContext amoebaGpu, unsigned int sharedMemoryPerThread );
 
+extern int isNanOrInfinity( double number );
+extern void trackMutualInducedIterations( amoebaGpuContext amoebaGpu, int iteration);
+
 #endif //__AMOEBA_GPU_TYPES_H__
 

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

@@ -767,9 +767,8 @@ static void cudaComputeAmoebaMutualInducedAndGkFieldBySOR( amoebaGpuContext amoe
 
     static int timestep                  = 0;
     timestep++;
-    static const char* methodName        = "cudaComputeAmoebaMutualInducedAndGkFieldBySOR";
-    static double iterationStat[6]       = { 0.0, 0.0, 1000.0, 0.0, 0.0, 0.0 };
 #ifdef AMOEBA_DEBUG
+    static const char* methodName        = "cudaComputeAmoebaMutualInducedAndGkFieldBySOR";
     std::vector<int> fileId;
     fileId.resize( 2 );
     fileId[0] = timestep;
@@ -859,7 +858,6 @@ static void cudaComputeAmoebaMutualInducedAndGkFieldBySOR( amoebaGpuContext amoe
     done      = 0;
     iteration = 1;
 
-time_t start = clock();
     while( !done ){
 
         // matrix multiply
@@ -919,6 +917,7 @@ time_t start = clock();
 #endif
 
         // Debye=4.8033324f
+
         amoebaGpu->psCurrentEpsilon->Download();
         float currentEpsilon                     = amoebaGpu->psCurrentEpsilon->_pSysStream[0][0];
         amoebaGpu->mutualInducedCurrentEpsilon   = currentEpsilon;
@@ -983,43 +982,24 @@ time_t start = clock();
         }
 #endif
         iteration++;
-//done = 1;
-//if(  iteration > 1 )exit(0);
     }
 
     amoebaGpu->mutualInducedDone             = done;
     amoebaGpu->mutualInducedConverged        = ( !done || iteration > amoebaGpu->mutualInducedMaxIterations ) ? 0 : 1;
 
     if( amoebaGpu->log ){
-        static int count = 0;
-        count++;
-        double interationD = static_cast<double>(iteration);
-        iterationStat[0]  += interationD;
-        iterationStat[1]  += interationD*interationD;
-        iterationStat[2]   = interationD < iterationStat[2] ? interationD : iterationStat[2];
-        iterationStat[3]   = interationD > iterationStat[3] ? interationD : iterationStat[3];
-        iterationStat[4]  += 1.0;
-        if( count == 100 ){
-            double average = iterationStat[0]/iterationStat[4]; 
-            double stddev  = iterationStat[1] - average*average*iterationStat[4]; 
-                   stddev  = sqrt( stddev )/(iterationStat[4]-1.0);
-            (void) fprintf( amoebaGpu->log, "%s iteration=%10.3f stddev=%10.3f min/max[%10.3f %10.3f] %10.1f eps=%14.7e\n",
-                            methodName, average, stddev, iterationStat[2], iterationStat[3], iterationStat[4], amoebaGpu->mutualInducedCurrentEpsilon );
-            (void) fflush( amoebaGpu->log );
-            iterationStat[0] = iterationStat[1] = iterationStat[4] = 0.0;
-            count = 0;
-        }
+        trackMutualInducedIterations( amoebaGpu, iteration );
     }
 
 #ifdef AMOEBA_DEBUG
-    if( 1 ){
+    if( 0 ){
         std::vector<int> fileId;
         //fileId.push_back( 0 );
         VectorOfDoubleVectors outputVector;
         cudaLoadCudaFloat4Array( gpu->natoms, 3, gpu->psPosq4,                    outputVector );
         cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipole,      outputVector );
         cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipolePolar, outputVector );
-        cudaWriteVectorOfDoubleVectorsToFile( "CudaMI", fileId, outputVector );
+        cudaWriteVectorOfDoubleVectorsToFile( "CudaMI_GK", fileId, outputVector );
      }
 #endif
 

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

@@ -34,6 +34,7 @@ void GetCalculateAmoebaCudaMutualInducedFieldSim(amoebaGpuContext amoebaGpu)
 }
 
 //#define AMOEBA_DEBUG
+#undef AMOEBA_DEBUG
 
 __device__ void calculateMutualInducedFieldPairIxn_kernel( float4 atomCoordinatesI, float4 atomCoordinatesJ,
                                                            float dampingFactorI,    float dampingFactorJ,
@@ -602,9 +603,8 @@ static void cudaComputeAmoebaMutualInducedFieldBySOR( amoebaGpuContext amoebaGpu
   
    // ---------------------------------------------------------------------------------------
 
-    static const char* methodName = "cudaComputeAmoebaMutualInducedFieldBySOR";
-
 #ifdef AMOEBA_DEBUG
+    static const char* methodName = "cudaComputeAmoebaMutualInducedFieldBySOR";
     static int timestep = 0;
     std::vector<int> fileId;
     timestep++;
@@ -739,29 +739,16 @@ static void cudaComputeAmoebaMutualInducedFieldBySOR( amoebaGpuContext amoebaGpu
            amoebaGpu->psWorkVector[0]->_pDevStream[0],     amoebaGpu->psWorkVector[1]->_pDevStream[0] );
         LAUNCHERROR("kSorUpdateMutualInducedField");  
 
-        // // get total epsilon -- performing sums on gpu
+        // get total epsilon -- performing sums on gpu
 
         kReduceMutualInducedFieldDelta_kernel<<<1, amoebaGpu->epsilonThreadsPerBlock, 2*sizeof(float)*amoebaGpu->epsilonThreadsPerBlock>>>(
            3*gpu->natoms, amoebaGpu->psWorkVector[0]->_pDevStream[0], amoebaGpu->psWorkVector[1]->_pDevStream[0],
            amoebaGpu->psCurrentEpsilon->_pDevStream[0] );
         LAUNCHERROR("kReduceMutualInducedFieldDelta");
 
-#if 1
-        // get total epsilon -- performing sums on cpu
-{
-        float sum1                   = cudaGetSum( 3*gpu->natoms, amoebaGpu->psWorkVector[0]  );
-        float sum2                   = cudaGetSum( 3*gpu->natoms, amoebaGpu->psWorkVector[1] );
-        sum1                         = 4.8033324f*sqrtf( sum1/( (float) gpu->natoms) );
-        sum2                         = 4.8033324f*sqrtf( sum2/( (float) gpu->natoms) );
-
-        float currentEpsilon         = sum1 > sum2 ? sum1 : sum2;
-        amoebaGpu->mutualInducedCurrentEpsilon  = currentEpsilon;
-        (void) fprintf( amoebaGpu->log, "%s iteration=%3d eps %14.6e [%14.6e %14.6e] done=%d sums=%14.6e %14.6e\n",
-                        methodName, iteration, amoebaGpu->mutualInducedCurrentEpsilon,
-                           amoebaGpu->psCurrentEpsilon->_pSysStream[0][1], 
-                           amoebaGpu->psCurrentEpsilon->_pSysStream[0][2], done, sum1, sum2 );
-}
-#endif
+        if( amoebaGpu->log ){
+            trackMutualInducedIterations( amoebaGpu, iteration);
+        }
 
         // Debye=4.8033324f
         amoebaGpu->psCurrentEpsilon->Download();
@@ -810,14 +797,14 @@ static void cudaComputeAmoebaMutualInducedFieldBySOR( amoebaGpuContext amoebaGpu
         }
 #endif
         iteration++;
-//if(  iteration > 1 )exit(0);
     }
 
     amoebaGpu->mutualInducedDone             = done;
     amoebaGpu->mutualInducedConverged        = ( !done || iteration > amoebaGpu->mutualInducedMaxIterations ) ? 0 : 1;
 
 #ifdef AMOEBA_DEBUG
-    if( 1 ){
+/*
+    if( 0 ){
         std::vector<int> fileId;
         //fileId.push_back( 0 );
         VectorOfDoubleVectors outputVector;
@@ -826,6 +813,7 @@ static void cudaComputeAmoebaMutualInducedFieldBySOR( amoebaGpuContext amoebaGpu
         cudaLoadCudaFloatArray( gpu->natoms,  3, amoebaGpu->psInducedDipolePolar, outputVector );
         cudaWriteVectorOfDoubleVectorsToFile( "CudaMI", fileId, outputVector );
      }
+*/
 #endif
 
    // ---------------------------------------------------------------------------------------

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

@@ -48,8 +48,8 @@ int main( int numberOfArguments, char* argv[] ) {
         Platform::loadPluginsFromDirectory( Platform::getDefaultPluginsDirectory() );
         if( numberOfArguments > 1 ){
             MapStringString argumentMap;
-            appendInputArgumentsToArgumentMap( numberOfArguments, argv, argumentMap );
             argumentMap[INCLUDE_OBC_CAVITY_TERM] = "0";
+            appendInputArgumentsToArgumentMap( numberOfArguments, argv, argumentMap );
             runTestsUsingAmoebaTinkerParameterFile( argumentMap );
         }        
     }