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

Edits to address potential thread-safety issues:

Removed static variables in CudaFreeEnergyKernels.cpp
Moved declaration of global feSim to within method where actually used  in kCalculateNonbondedSoftcore.cu
parent ab675f48
...@@ -386,8 +386,6 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::executeForces(ContextImpl& ...@@ -386,8 +386,6 @@ void CudaFreeEnergyCalcNonbondedSoftcoreForceKernel::executeForces(ContextImpl&
// --------------------------------------------------------------------------------------- // ---------------------------------------------------------------------------------------
_gpuContext* gpu = data.gpu; _gpuContext* gpu = data.gpu;
static int call = 0;
call++;
// write array, ... address's to board // write array, ... address's to board
...@@ -532,7 +530,6 @@ void CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::executeForces(ContextImpl& co ...@@ -532,7 +530,6 @@ void CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::executeForces(ContextImpl& co
_gpuContext* gpu = data.gpu; _gpuContext* gpu = data.gpu;
int debug = 1; int debug = 1;
static int call = 0;
// send address's of arrays, ... to device on first call // send address's of arrays, ... to device on first call
// required since force/energy buffers not set when CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::initialize() was called // required since force/energy buffers not set when CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::initialize() was called
...@@ -550,7 +547,6 @@ void CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::executeForces(ContextImpl& co ...@@ -550,7 +547,6 @@ void CudaFreeEnergyCalcGBSAOBCSoftcoreForceKernel::executeForces(ContextImpl& co
// calculate Born radii and first loop of Obc forces // calculate Born radii and first loop of Obc forces
if( debug && log ){ if( debug && log ){
call++;
if( log ){ if( log ){
(void) fprintf( log, "\n%s: calling kCalculateCDLJObcGbsaSoftcoreForces1\n", methodName.c_str() ); (void) fprintf( log, "\n%s: calling kCalculateCDLJObcGbsaSoftcoreForces1\n", methodName.c_str() );
(void) fflush( log ); (void) fflush( log );
...@@ -675,7 +671,6 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte ...@@ -675,7 +671,6 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
_gpuContext* gpu = data.gpu; _gpuContext* gpu = data.gpu;
int debug = 1; int debug = 1;
static int call = 0;
// send address's of arrays, ... to device on first call // send address's of arrays, ... to device on first call
// required since force/energy buffers not set when CudaFreeEnergyCalcGBVISoftcoreForceKernel::initialize() was called // required since force/energy buffers not set when CudaFreeEnergyCalcGBVISoftcoreForceKernel::initialize() was called
...@@ -691,7 +686,6 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte ...@@ -691,7 +686,6 @@ void CudaFreeEnergyCalcGBVISoftcoreForceKernel::executeForces(ContextImpl& conte
// calculate Born radii and first loop of GB/VI forces // calculate Born radii and first loop of GB/VI forces
if( debug && log ){ if( debug && log ){
call++;
if( log ){ if( log ){
(void) fprintf( log, "\n%s: calling kCalculateCDLJObcGbsaSoftcoreForces1 & %s\n", methodName.c_str(), (void) fprintf( log, "\n%s: calling kCalculateCDLJObcGbsaSoftcoreForces1 & %s\n", methodName.c_str(),
getQuinticScaling() ? "kReduceGBVIBornSumQuinticScaling" : "kReduceGBVIBornSum" ); getQuinticScaling() ? "kReduceGBVIBornSumQuinticScaling" : "kReduceGBVIBornSum" );
......
...@@ -34,7 +34,7 @@ ...@@ -34,7 +34,7 @@
struct cudaFreeEnergySimulationNonBonded { struct cudaFreeEnergySimulationNonBonded {
float* pParticleSoftCoreLJLambda; float* pParticleSoftCoreLJLambda;
}; };
struct cudaFreeEnergySimulationNonBonded feSim; //struct cudaFreeEnergySimulationNonBonded feSim;
// device handles // device handles
...@@ -55,6 +55,8 @@ void SetCalculateCDLJSoftcoreGpuSim( gpuContext gpu ) ...@@ -55,6 +55,8 @@ void SetCalculateCDLJSoftcoreGpuSim( gpuContext gpu )
void SetCalculateCDLJSoftcoreSupplementarySim( float* gpuParticleSoftCoreLJLambda) void SetCalculateCDLJSoftcoreSupplementarySim( float* gpuParticleSoftCoreLJLambda)
{ {
cudaError_t status; cudaError_t status;
struct cudaFreeEnergySimulationNonBonded feSim;
feSim.pParticleSoftCoreLJLambda = gpuParticleSoftCoreLJLambda; feSim.pParticleSoftCoreLJLambda = gpuParticleSoftCoreLJLambda;
status = cudaMemcpyToSymbol(feSimDev, &feSim, sizeof(cudaFreeEnergySimulationNonBonded)); status = cudaMemcpyToSymbol(feSimDev, &feSim, sizeof(cudaFreeEnergySimulationNonBonded));
RTERROR(status, "cudaMemcpyToSymbol: SetCalculateCDLJSoftcoreSupplementarySim"); RTERROR(status, "cudaMemcpyToSymbol: SetCalculateCDLJSoftcoreSupplementarySim");
......
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