Commit e0741f7b authored by Mark Friedrichs's avatar Mark Friedrichs
Browse files

Removed logging output

parent a2dc3e23
...@@ -777,9 +777,11 @@ void cudaComputeAmoebaElectrostatic( amoebaGpuContext amoebaGpu ) ...@@ -777,9 +777,11 @@ void cudaComputeAmoebaElectrostatic( amoebaGpuContext amoebaGpu )
if (gpu->bOutputBufferPerWarp){ if (gpu->bOutputBufferPerWarp){
#ifdef AMOEBA_DEBUG
(void) fprintf( amoebaGpu->log, "kCalculateAmoebaCudaElectrostaticN2Forces warp: numBlocks=%u numThreads=%u bufferPerWarp=%u atm=%lu shrd=%lu ixnCt=%lu workUnits=%u\n", (void) fprintf( amoebaGpu->log, "kCalculateAmoebaCudaElectrostaticN2Forces warp: numBlocks=%u numThreads=%u bufferPerWarp=%u atm=%lu shrd=%lu ixnCt=%lu workUnits=%u\n",
amoebaGpu->nonbondBlocks, threadsPerBlock, amoebaGpu->bOutputBufferPerWarp, amoebaGpu->nonbondBlocks, threadsPerBlock, amoebaGpu->bOutputBufferPerWarp,
sizeof(ElectrostaticParticle), sizeof(ElectrostaticParticle)*threadsPerBlock, (*gpu->psInteractionCount)[0], gpu->sim.workUnits ); (void) fflush( amoebaGpu->log ); sizeof(ElectrostaticParticle), sizeof(ElectrostaticParticle)*threadsPerBlock, (*gpu->psInteractionCount)[0], gpu->sim.workUnits ); (void) fflush( amoebaGpu->log );
#endif
kCalculateAmoebaCudaElectrostaticN2ByWarpForces_kernel<<<amoebaGpu->nonbondBlocks, threadsPerBlock, sizeof(ElectrostaticParticle)*threadsPerBlock>>>( kCalculateAmoebaCudaElectrostaticN2ByWarpForces_kernel<<<amoebaGpu->nonbondBlocks, threadsPerBlock, sizeof(ElectrostaticParticle)*threadsPerBlock>>>(
......
...@@ -411,7 +411,6 @@ void cudaComputeAmoebaFixedEAndGkFields( amoebaGpuContext amoebaGpu ) ...@@ -411,7 +411,6 @@ void cudaComputeAmoebaFixedEAndGkFields( amoebaGpuContext amoebaGpu )
kClearFields_3( amoebaGpu, 3 ); kClearFields_3( amoebaGpu, 3 );
if (gpu->bOutputBufferPerWarp){ if (gpu->bOutputBufferPerWarp){
(void) fprintf( amoebaGpu->log, "N2 warp\n" );
kCalculateAmoebaFixedEAndGkFieldN2ByWarp_kernel<<<amoebaGpu->nonbondBlocks, threadsPerBlock, sizeof(FixedFieldParticle)*threadsPerBlock>>>( kCalculateAmoebaFixedEAndGkFieldN2ByWarp_kernel<<<amoebaGpu->nonbondBlocks, threadsPerBlock, sizeof(FixedFieldParticle)*threadsPerBlock>>>(
amoebaGpu->psWorkUnit->_pDevData, amoebaGpu->psWorkUnit->_pDevData,
gpu->psPosq4->_pDevData, gpu->psPosq4->_pDevData,
......
...@@ -1893,14 +1893,16 @@ void kCalculateAmoebaKirkwood( amoebaGpuContext amoebaGpu ) ...@@ -1893,14 +1893,16 @@ void kCalculateAmoebaKirkwood( amoebaGpuContext amoebaGpu )
//unsigned int eDiffhreadsPerBlock = getThreadsPerBlock( amoebaGpu, sizeof(KirkwoodEDiffParticle)); //unsigned int eDiffhreadsPerBlock = getThreadsPerBlock( amoebaGpu, sizeof(KirkwoodEDiffParticle));
//unsigned int maxThreadsPerBlock = threadsPerBlock> eDiffhreadsPerBlock ? threadsPerBlock : eDiffhreadsPerBlock; //unsigned int maxThreadsPerBlock = threadsPerBlock> eDiffhreadsPerBlock ? threadsPerBlock : eDiffhreadsPerBlock;
#ifdef AMOEBA_DEBUG
if( amoebaGpu->log ){ if( amoebaGpu->log ){
(void) fprintf( amoebaGpu->log, "kCalculateAmoebaCudaKirkwood: blcks=%u tds=%u %u bPrWrp=%u atm=%lu shrd=%lu ixnCt=%lu workUnits=%u\n", (void) fprintf( amoebaGpu->log, "kCalculateAmoebaCudaKirkwood: blcks=%u tds=%u %u bPrWrp=%u atm=%lu shrd=%lu ixnCt=%lu workUnits=%u\n",
amoebaGpu->nonbondBlocks, threadsPerBlock, maxThreads, amoebaGpu->bOutputBufferPerWarp, amoebaGpu->nonbondBlocks, threadsPerBlock, maxThreads, amoebaGpu->bOutputBufferPerWarp,
sizeof(KirkwoodParticle), sizeof(KirkwoodParticle)*threadsPerBlock, sizeof(KirkwoodParticle), sizeof(KirkwoodParticle)*threadsPerBlock,
(*gpu->psInteractionCount)[0], gpu->sim.workUnits ); (*gpu->psInteractionCount)[0], gpu->sim.workUnits );
(void) fflush( amoebaGpu->log ); (void) fflush( amoebaGpu->log );
} }
#endif
} }
kClearFields_1( amoebaGpu ); kClearFields_1( amoebaGpu );
......
...@@ -1626,6 +1626,7 @@ void kCalculateAmoebaLocalForces_kernel() ...@@ -1626,6 +1626,7 @@ void kCalculateAmoebaLocalForces_kernel()
void kCalculateAmoebaLocalForces(amoebaGpuContext gpu) void kCalculateAmoebaLocalForces(amoebaGpuContext gpu)
{ {
#ifdef AMOEBA_DEBUG
if( gpu->log ){ if( gpu->log ){
static int call = 1; static int call = 1;
if( call == 0 ){ if( call == 0 ){
...@@ -1634,6 +1635,7 @@ void kCalculateAmoebaLocalForces(amoebaGpuContext gpu) ...@@ -1634,6 +1635,7 @@ void kCalculateAmoebaLocalForces(amoebaGpuContext gpu)
call++; call++;
} }
} }
#endif
kCalculateAmoebaLocalForces_kernel<<<gpu->gpuContext->sim.blocks, gpu->gpuContext->sim.localForces_threads_per_block>>>(); kCalculateAmoebaLocalForces_kernel<<<gpu->gpuContext->sim.blocks, gpu->gpuContext->sim.localForces_threads_per_block>>>();
LAUNCHERROR("kCalculateAmoebaLocalForces"); LAUNCHERROR("kCalculateAmoebaLocalForces");
......
...@@ -701,7 +701,7 @@ void cudaComputeAmoebaMapTorques( amoebaGpuContext amoebaGpu, CUDAStream<float>* ...@@ -701,7 +701,7 @@ void cudaComputeAmoebaMapTorques( amoebaGpuContext amoebaGpu, CUDAStream<float>*
// check that BLOCK_SIZE >= amoebaGpu->maxMapTorqueDifference // check that BLOCK_SIZE >= amoebaGpu->maxMapTorqueDifference
if( amoebaGpu->maxMapTorqueDifference > BLOCK_SIZE ){ if( amoebaGpu->maxMapTorqueDifference > BLOCK_SIZE ){
(void) fprintf( amoebaGpu->log, "cudaComputeAmoebaMapTorques: block size (%d) in amoebaMapTorqueReduce_kernel is too small ( > %d)! -- aborting.\n" (void) fprintf( stderr, "cudaComputeAmoebaMapTorques: block size (%d) in amoebaMapTorqueReduce_kernel is too small ( > %d)! -- aborting.\n"
"This is likely due to the presence of disulfide bond or other bond between atoms whose indices is greater than %d; this limitation will be removed w/ the next software update.\n", "This is likely due to the presence of disulfide bond or other bond between atoms whose indices is greater than %d; this limitation will be removed w/ the next software update.\n",
BLOCK_SIZE, amoebaGpu->maxMapTorqueDifference, BLOCK_SIZE ); BLOCK_SIZE, amoebaGpu->maxMapTorqueDifference, BLOCK_SIZE );
exit(-1); exit(-1);
...@@ -872,7 +872,7 @@ void cudaComputeAmoebaMapTorquesAndAddTotalForce( amoebaGpuContext amoebaGpu, ...@@ -872,7 +872,7 @@ void cudaComputeAmoebaMapTorquesAndAddTotalForce( amoebaGpuContext amoebaGpu,
// check that BLOCK_SIZE >= amoebaGpu->maxMapTorqueDifference // check that BLOCK_SIZE >= amoebaGpu->maxMapTorqueDifference
if( amoebaGpu->maxMapTorqueDifference > BLOCK_SIZE ){ if( amoebaGpu->maxMapTorqueDifference > BLOCK_SIZE ){
(void) fprintf( amoebaGpu->log, "cudaComputeAmoebaMapTorquesAndAddTotalForce: block size (%d) in amoebaMapTorqueReduce_kernel is too small ( > %d)! -- aborting.\n" (void) fprintf( stderr, "cudaComputeAmoebaMapTorquesAndAddTotalForce: block size (%d) in amoebaMapTorqueReduce_kernel is too small ( > %d)! -- aborting.\n"
"This is likely due to the presence of disulfide bond or other bond between atoms whose indices is greater than %d; this limitation will be removed w/ the next software update.\n", "This is likely due to the presence of disulfide bond or other bond between atoms whose indices is greater than %d; this limitation will be removed w/ the next software update.\n",
BLOCK_SIZE, amoebaGpu->maxMapTorqueDifference, BLOCK_SIZE ); BLOCK_SIZE, amoebaGpu->maxMapTorqueDifference, BLOCK_SIZE );
exit(-1); exit(-1);
...@@ -1075,7 +1075,7 @@ void cudaComputeAmoebaMapTorquesAndAddTotalForce2( amoebaGpuContext amoebaGpu, ...@@ -1075,7 +1075,7 @@ void cudaComputeAmoebaMapTorquesAndAddTotalForce2( amoebaGpuContext amoebaGpu,
// check that BLOCK_SIZE >= amoebaGpu->maxMapTorqueDifference // check that BLOCK_SIZE >= amoebaGpu->maxMapTorqueDifference
if( amoebaGpu->maxMapTorqueDifference > BLOCK_SIZE ){ if( amoebaGpu->maxMapTorqueDifference > BLOCK_SIZE ){
(void) fprintf( amoebaGpu->log, "cudaComputeAmoebaMapTorquesAndAddTotalForce2: block size (%d) in amoebaMapTorqueReduce_kernel is too small ( > %d)! -- aborting.\n" (void) fprintf( stderr, "cudaComputeAmoebaMapTorquesAndAddTotalForce2: block size (%d) in amoebaMapTorqueReduce_kernel is too small ( > %d)! -- aborting.\n"
"This is likely due to the presence of disulfide bond or other bond between atoms whose indices is greater than %d; this limitation will be removed w/ the next software update.\n", "This is likely due to the presence of disulfide bond or other bond between atoms whose indices is greater than %d; this limitation will be removed w/ the next software update.\n",
BLOCK_SIZE, amoebaGpu->maxMapTorqueDifference, BLOCK_SIZE ); BLOCK_SIZE, amoebaGpu->maxMapTorqueDifference, BLOCK_SIZE );
exit(-1); exit(-1);
......
...@@ -1392,6 +1392,7 @@ extern double kReduceEnergy( gpuContext gpu); ...@@ -1392,6 +1392,7 @@ extern double kReduceEnergy( gpuContext gpu);
void cudaComputeAmoebaPmeElectrostatic( amoebaGpuContext amoebaGpu ) void cudaComputeAmoebaPmeElectrostatic( amoebaGpuContext amoebaGpu )
{ {
#ifdef AMOEBA_DEBUG
if( 0 ){ if( 0 ){
gpuContext gpu = amoebaGpu->gpuContext; gpuContext gpu = amoebaGpu->gpuContext;
std::vector<int> fileId; std::vector<int> fileId;
...@@ -1441,7 +1442,9 @@ void cudaComputeAmoebaPmeElectrostatic( amoebaGpuContext amoebaGpu ) ...@@ -1441,7 +1442,9 @@ void cudaComputeAmoebaPmeElectrostatic( amoebaGpuContext amoebaGpu )
//zeroForce( amoebaGpu ); //zeroForce( amoebaGpu );
exit(0); exit(0);
} }
#endif
#ifdef AMOEBA_DEBUG
if( 0 ){ if( 0 ){
gpuContext gpu = amoebaGpu->gpuContext; gpuContext gpu = amoebaGpu->gpuContext;
std::vector<int> fileId; std::vector<int> fileId;
...@@ -1476,6 +1479,7 @@ void cudaComputeAmoebaPmeElectrostatic( amoebaGpuContext amoebaGpu ) ...@@ -1476,6 +1479,7 @@ void cudaComputeAmoebaPmeElectrostatic( amoebaGpuContext amoebaGpu )
cudaLoadCudaFloatArray( gpu->natoms, 3, amoebaGpu->psForce, outputVector, gpu->psAtomIndex->_pSysData, 1.0f ); cudaLoadCudaFloatArray( gpu->natoms, 3, amoebaGpu->psForce, outputVector, gpu->psAtomIndex->_pSysData, 1.0f );
cudaWriteVectorOfDoubleVectorsToFile( "CudaPrePmeForce", fileId, outputVector ); cudaWriteVectorOfDoubleVectorsToFile( "CudaPrePmeForce", fileId, outputVector );
} }
#endif
cudaComputeAmoebaPmeDirectElectrostatic( amoebaGpu ); cudaComputeAmoebaPmeDirectElectrostatic( amoebaGpu );
kCalculateAmoebaPMEInducedDipoleForces( amoebaGpu ); kCalculateAmoebaPMEInducedDipoleForces( amoebaGpu );
......
...@@ -769,16 +769,16 @@ void) fflush( amoebaGpu->log ); ...@@ -769,16 +769,16 @@ void) fflush( amoebaGpu->log );
amoebaGpu->psCurrentEpsilon->_pSysData[1], amoebaGpu->psCurrentEpsilon->_pSysData[1],
amoebaGpu->psCurrentEpsilon->_pSysData[2], done ); amoebaGpu->psCurrentEpsilon->_pSysData[2], done );
(void) fflush( amoebaGpu->log ); (void) fflush( amoebaGpu->log );
#endif #endif
// exit if nan // exit if nan
if( 0 && amoebaGpu->mutualInducedCurrentEpsilon != amoebaGpu->mutualInducedCurrentEpsilon ){ if( 0 && amoebaGpu->mutualInducedCurrentEpsilon != amoebaGpu->mutualInducedCurrentEpsilon ){
(void) fprintf( amoebaGpu->log, "PME MI iteration=%3d eps is nan -- exiting.\n", iteration ); (void) fprintf( stderr, "PME MI iteration=%3d eps is nan -- exiting.\n", iteration );
exit(0); exit(0);
} }
iteration++; iteration++;
} }
......
...@@ -605,8 +605,8 @@ void kCalculateAmoebaMultipoleForces(amoebaGpuContext amoebaGpu, bool hasAmoebaG ...@@ -605,8 +605,8 @@ void kCalculateAmoebaMultipoleForces(amoebaGpuContext amoebaGpu, bool hasAmoebaG
// check if induce dipole calculation converged -- abort if it did not // check if induce dipole calculation converged -- abort if it did not
if( amoebaGpu->mutualInducedDone == 0 ){ if( amoebaGpu->mutualInducedDone == 0 ){
(void) fprintf( amoebaGpu->log, "%s induced dipole calculation did not converge -- aborting!\n", methodName.c_str() ); (void) fprintf( stderr, "%s induced dipole calculation did not converge -- aborting!\n", methodName.c_str() );
(void) fflush( amoebaGpu->log ); (void) fflush( stderr );
exit(-1); exit(-1);
} }
......
...@@ -583,6 +583,7 @@ void kCalculateAmoebaVdw14_7Forces( amoebaGpuContext amoebaGpu, int applyCutoff ...@@ -583,6 +583,7 @@ void kCalculateAmoebaVdw14_7Forces( amoebaGpuContext amoebaGpu, int applyCutoff
sizeof(unsigned int)*gpu->sim.nonbond_threads_per_block>>>(gpu->sim.pInteractingWorkUnit); sizeof(unsigned int)*gpu->sim.nonbond_threads_per_block>>>(gpu->sim.pInteractingWorkUnit);
LAUNCHERROR("kFindInteractionsWithinBlocksVdwPeriodic"); LAUNCHERROR("kFindInteractionsWithinBlocksVdwPeriodic");
#ifdef AMOEBA_DEBUG
if( 0 ){ if( 0 ){
gpu->psInteractionCount->Download(); gpu->psInteractionCount->Download();
gpu->psInteractingWorkUnit->Download(); gpu->psInteractingWorkUnit->Download();
...@@ -605,6 +606,7 @@ void kCalculateAmoebaVdw14_7Forces( amoebaGpuContext amoebaGpu, int applyCutoff ...@@ -605,6 +606,7 @@ void kCalculateAmoebaVdw14_7Forces( amoebaGpuContext amoebaGpu, int applyCutoff
} }
(void) fflush( amoebaGpu->log ); (void) fflush( amoebaGpu->log );
} }
#endif
if (gpu->bOutputBufferPerWarp){ if (gpu->bOutputBufferPerWarp){
kCalculateAmoebaVdw14_7CutoffByWarp_kernel<<<amoebaGpu->nonbondBlocks, threadsPerBlock, sizeof(Vdw14_7Particle)*threadsPerBlock>>>( kCalculateAmoebaVdw14_7CutoffByWarp_kernel<<<amoebaGpu->nonbondBlocks, threadsPerBlock, sizeof(Vdw14_7Particle)*threadsPerBlock>>>(
......
...@@ -30,10 +30,6 @@ void GetCalculateAmoebaCudaWcaDispersionSim(amoebaGpuContext amoebaGpu) ...@@ -30,10 +30,6 @@ void GetCalculateAmoebaCudaWcaDispersionSim(amoebaGpuContext amoebaGpu)
cudaError_t status; cudaError_t status;
gpuContext gpu = amoebaGpu->gpuContext; gpuContext gpu = amoebaGpu->gpuContext;
status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation)); status = cudaMemcpyFromSymbol(&gpu->sim, cSim, sizeof(cudaGmxSimulation));
fprintf( stderr, "In GetCalculateAmoebaCudaWcaDispersionSim: %p %lu %u\n",
gpu->psInteractionCount->_pSysData, gpu->psInteractionCount->_pSysData[0], gpu->sim.workUnits );
RTERROR(status, "GetCalculateAmoebaCudaWcaDispersionSim: cudaMemcpyFromSymbol: SetSim copy from cSim failed"); RTERROR(status, "GetCalculateAmoebaCudaWcaDispersionSim: cudaMemcpyFromSymbol: SetSim copy from cSim failed");
status = cudaMemcpyFromSymbol(&amoebaGpu->amoebaSim, cAmoebaSim, sizeof(cudaAmoebaGmxSimulation)); status = cudaMemcpyFromSymbol(&amoebaGpu->amoebaSim, cAmoebaSim, sizeof(cudaAmoebaGmxSimulation));
RTERROR(status, "GetCalculateAmoebaCudaWcaDispersionSim: cudaMemcpyFromSymbol: SetSim copy from cAmoebaSim failed"); RTERROR(status, "GetCalculateAmoebaCudaWcaDispersionSim: cudaMemcpyFromSymbol: SetSim copy from cAmoebaSim failed");
...@@ -612,6 +608,7 @@ void kCalculateAmoebaWcaDispersionForces( amoebaGpuContext amoebaGpu ) ...@@ -612,6 +608,7 @@ void kCalculateAmoebaWcaDispersionForces( amoebaGpuContext amoebaGpu )
kReduceWcaDispersionToFloat4( amoebaGpu, gpu->psForce4 ); kReduceWcaDispersionToFloat4( amoebaGpu, gpu->psForce4 );
#ifdef AMOEBA_DEBUG
if( 0 ){ if( 0 ){
gpu->psEnergy->Download(); gpu->psEnergy->Download();
double sum = 0.0; double sum = 0.0;
...@@ -622,6 +619,7 @@ void kCalculateAmoebaWcaDispersionForces( amoebaGpuContext amoebaGpu ) ...@@ -622,6 +619,7 @@ void kCalculateAmoebaWcaDispersionForces( amoebaGpuContext amoebaGpu )
} }
(void) fprintf( amoebaGpu->log,"%14.7e QQ SUM\n", sum ); (void) fprintf( amoebaGpu->log,"%14.7e QQ SUM\n", sum );
} }
#endif
if( 0 ){ if( 0 ){
int paddedNumberOfAtoms = amoebaGpu->gpuContext->sim.paddedNumberOfAtoms; int paddedNumberOfAtoms = amoebaGpu->gpuContext->sim.paddedNumberOfAtoms;
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment