Skip to content

GitLab

Unverified Commit b28d2e66 authored by Peter Eastman's avatar Peter Eastman Committed by GitHub
Browse files

Merged parallel code (#4649) 

* Unified lots of parallel computation code between platforms

* Unified test code between platforms

* Eliminated duplicated timing code
parent 78902bed
Hide whitespace changes
Inline Side-by-side
Showing with 34 additions and 1933 deletions
platforms/cuda/tests/TestCudaPeriodicTorsionForce.cpp
View file @ b28d2e66
...@@ -32,49 +32,6 @@ ...@@ -32,49 +32,6 @@
#include "CudaTests.h" #include "CudaTests.h"
#include "TestPeriodicTorsionForce.h" #include "TestPeriodicTorsionForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
PeriodicTorsionForce* force = new PeriodicTorsionForce();
for (int i = 3; i < numParticles; i++)
force->addTorsion(i-3, i-2, i-1, i, 2, 1.1, i);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, i%3);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, CudaPlatform::CudaDeviceIndex());
map<string, string> props;
props[CudaPlatform::CudaDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
// Try updating some parameters and see if they still match.
for (int i = 95; i < 102; i++) {
int p1, p2, p3, p4, periodicity;
double phase, k;
force->getTorsionParameters(i, p1, p2, p3, p4, periodicity, phase, k);
force->setTorsionParameters(i, p1, p2, p3, p4, periodicity, phase+0.1, 2*k);
}
force->updateParametersInContext(context1);
force->updateParametersInContext(context2);
State state3 = context1.getState(State::Energy);
State state4 = context2.getState(State::Energy);
ASSERT_EQUAL_TOL(state3.getPotentialEnergy(), state4.getPotentialEnergy(), 1e-5);
ASSERT(fabs(state1.getPotentialEnergy()-state3.getPotentialEnergy()) > 0.1);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/cuda/tests/TestCudaRBTorsionForce.cpp
View file @ b28d2e66
...@@ -32,34 +32,6 @@ ...@@ -32,34 +32,6 @@
#include "CudaTests.h" #include "CudaTests.h"
#include "TestRBTorsionForce.h" #include "TestRBTorsionForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
RBTorsionForce* force = new RBTorsionForce();
for (int i = 3; i < numParticles; i++)
force->addTorsion(i-3, i-2, i-1, i, 2, 0.1*i, 0.5*i, i, 1, 1);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, i%3);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, CudaPlatform::CudaDeviceIndex());
map<string, string> props;
props[CudaPlatform::CudaDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/include/HipContext.h
View file @ b28d2e66
...@@ -157,6 +157,11 @@ public: ...@@ -157,6 +157,11 @@ public:
* one ComputeContext is created for each device. * one ComputeContext is created for each device.
*/ */
std::vector<ComputeContext*> getAllContexts(); std::vector<ComputeContext*> getAllContexts();
/**
* Get a workspace used for accumulating energy when a simulation is parallelized across
* multiple devices.
*/
double& getEnergyWorkspace();
/** /**
* Get the stream currently being used for execution. * Get the stream currently being used for execution.
*/ */
......
platforms/hip/include/HipParallelKernels.h
View file @ b28d2e66
...@@ -83,7 +83,7 @@ private: ...@@ -83,7 +83,7 @@ private:
class FinishComputationTask; class FinishComputationTask;
HipPlatform::PlatformData& data; HipPlatform::PlatformData& data;
std::vector<Kernel> kernels; std::vector<Kernel> kernels;
std::vector<long long> completionTimes; std::vector<double> completionTimes;
std::vector<double> contextNonbondedFractions; std::vector<double> contextNonbondedFractions;
HipArray contextForces; HipArray contextForces;
void* pinnedPositionBuffer; void* pinnedPositionBuffer;
...@@ -95,322 +95,6 @@ private: ...@@ -95,322 +95,6 @@ private:
std::vector<hipStream_t> peerCopyStream; std::vector<hipStream_t> peerCopyStream;
}; };
/**
* This kernel is invoked by HarmonicBondForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcHarmonicBondForceKernel : public CalcHarmonicBondForceKernel {
public:
HipParallelCalcHarmonicBondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcHarmonicBondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcHarmonicBondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the HarmonicBondForce this kernel will be used for
*/
void initialize(const System& system, const HarmonicBondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the HarmonicBondForce to copy the parameters from
* @param firstBond the index of the first bond whose parameters might have changed
* @param lastBond the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force, int firstBond, int lastBond);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomBondForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcCustomBondForceKernel : public CalcCustomBondForceKernel {
public:
HipParallelCalcCustomBondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomBondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomBondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomBondForce this kernel will be used for
*/
void initialize(const System& system, const CustomBondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomBondForce to copy the parameters from
* @param firstBond the index of the first bond whose parameters might have changed
* @param lastBond the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomBondForce& force, int firstBond, int lastBond);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by HarmonicAngleForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcHarmonicAngleForceKernel : public CalcHarmonicAngleForceKernel {
public:
HipParallelCalcHarmonicAngleForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcHarmonicAngleForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcHarmonicAngleForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the HarmonicAngleForce this kernel will be used for
*/
void initialize(const System& system, const HarmonicAngleForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the HarmonicAngleForce to copy the parameters from
* @param firstAngle the index of the first bond whose parameters might have changed
* @param lastAngle the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force, int firstAngle, int lastAngle);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomAngleForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcCustomAngleForceKernel : public CalcCustomAngleForceKernel {
public:
HipParallelCalcCustomAngleForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomAngleForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomAngleForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomAngleForce this kernel will be used for
*/
void initialize(const System& system, const CustomAngleForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomAngleForce to copy the parameters from
* @param firstAngle the index of the first bond whose parameters might have changed
* @param lastAngle the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomAngleForce& force, int firstAngle, int lastAngle);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by PeriodicTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcPeriodicTorsionForceKernel : public CalcPeriodicTorsionForceKernel {
public:
HipParallelCalcPeriodicTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcPeriodicTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcPeriodicTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the PeriodicTorsionForce this kernel will be used for
*/
void initialize(const System& system, const PeriodicTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
class Task;
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the PeriodicTorsionForce to copy the parameters from
* @param firstTorsion the index of the first torsion whose parameters might have changed
* @param lastTorsion the index of the last torsion whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force, int firstTorsion, int lastTorsion);
private:
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by RBTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcRBTorsionForceKernel : public CalcRBTorsionForceKernel {
public:
HipParallelCalcRBTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcRBTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcRBTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the RBTorsionForce this kernel will be used for
*/
void initialize(const System& system, const RBTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the RBTorsionForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const RBTorsionForce& force);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CMAPTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcCMAPTorsionForceKernel : public CalcCMAPTorsionForceKernel {
public:
HipParallelCalcCMAPTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCMAPTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCMAPTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CMAPTorsionForce this kernel will be used for
*/
void initialize(const System& system, const CMAPTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CMAPTorsionForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcCustomTorsionForceKernel : public CalcCustomTorsionForceKernel {
public:
HipParallelCalcCustomTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomTorsionForce this kernel will be used for
*/
void initialize(const System& system, const CustomTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomTorsionForce to copy the parameters from
* @param firstTorsion the index of the first torsion whose parameters might have changed
* @param lastTorsion the index of the last torsion whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force, int firstTorsion, int lastTorsion);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/** /**
* This kernel is invoked by NonbondedForce to calculate the forces acting on the system. * This kernel is invoked by NonbondedForce to calculate the forces acting on the system.
*/ */
...@@ -473,162 +157,6 @@ private: ...@@ -473,162 +157,6 @@ private:
std::vector<Kernel> kernels; std::vector<Kernel> kernels;
}; };
/**
* This kernel is invoked by CustomNonbondedForce to calculate the forces acting on the system.
*/
class HipParallelCalcCustomNonbondedForceKernel : public CalcCustomNonbondedForceKernel {
public:
HipParallelCalcCustomNonbondedForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomNonbondedForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomNonbondedForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomNonbondedForce this kernel will be used for
*/
void initialize(const System& system, const CustomNonbondedForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomNonbondedForce to copy the parameters from
* @param firstParticle the index of the first particle whose parameters might have changed
* @param lastParticle the index of the last particle whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force, int firstParticle, int lastParticle);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomExternalForce to calculate the forces acting on the system and the energy of the system.
*/
class HipParallelCalcCustomExternalForceKernel : public CalcCustomExternalForceKernel {
public:
HipParallelCalcCustomExternalForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomExternalForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomExternalForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomExternalForce this kernel will be used for
*/
void initialize(const System& system, const CustomExternalForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomExternalForce to copy the parameters from
* @param firstParticle the index of the first particle whose parameters might have changed
* @param lastParticle the index of the last particle whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomExternalForce& force, int firstParticle, int lastParticle);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomHbondForce to calculate the forces acting on the system.
*/
class HipParallelCalcCustomHbondForceKernel : public CalcCustomHbondForceKernel {
public:
HipParallelCalcCustomHbondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomHbondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomHbondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomHbondForce this kernel will be used for
*/
void initialize(const System& system, const CustomHbondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomHbondForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const CustomHbondForce& force);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomCompoundBondForce to calculate the forces acting on the system.
*/
class HipParallelCalcCustomCompoundBondForceKernel : public CalcCustomCompoundBondForceKernel {
public:
HipParallelCalcCustomCompoundBondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system);
CommonCalcCustomCompoundBondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomCompoundBondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomCompoundBondForce this kernel will be used for
*/
void initialize(const System& system, const CustomCompoundBondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomCompoundBondForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force);
private:
class Task;
HipPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
} // namespace OpenMM } // namespace OpenMM
#endif /*OPENMM_HIPPARALLELKERNELS_H_*/ #endif /*OPENMM_HIPPARALLELKERNELS_H_*/
platforms/hip/src/HipContext.cpp
View file @ b28d2e66
...@@ -672,6 +672,10 @@ vector<ComputeContext*> HipContext::getAllContexts() { ...@@ -672,6 +672,10 @@ vector<ComputeContext*> HipContext::getAllContexts() {
return result; return result;
} }
double& HipContext::getEnergyWorkspace() {
return platformData.contextEnergy[contextIndex];
}
hipStream_t HipContext::getCurrentStream() { hipStream_t HipContext::getCurrentStream() {
return currentStream; return currentStream;
} }
......
platforms/hip/src/HipKernelFactory.cpp
View file @ b28d2e66
...@@ -30,6 +30,7 @@ ...@@ -30,6 +30,7 @@
#include "HipParallelKernels.h" #include "HipParallelKernels.h"
#include "HipPlatform.h" #include "HipPlatform.h"
#include "openmm/common/CommonKernels.h" #include "openmm/common/CommonKernels.h"
#include "openmm/common/CommonParallelKernels.h"
#include "openmm/internal/ContextImpl.h" #include "openmm/internal/ContextImpl.h"
#include "openmm/OpenMMException.h" #include "openmm/OpenMMException.h"
...@@ -37,39 +38,39 @@ using namespace OpenMM; ...@@ -37,39 +38,39 @@ using namespace OpenMM;
KernelImpl* HipKernelFactory::createKernelImpl(std::string name, const Platform& platform, ContextImpl& context) const { KernelImpl* HipKernelFactory::createKernelImpl(std::string name, const Platform& platform, ContextImpl& context) const {
HipPlatform::PlatformData& data = *static_cast<HipPlatform::PlatformData*>(context.getPlatformData()); HipPlatform::PlatformData& data = *static_cast<HipPlatform::PlatformData*>(context.getPlatformData());
HipContext& cu = *data.contexts[0];
if (data.contexts.size() > 1) { if (data.contexts.size() > 1) {
// We are running in parallel on multiple devices, so we may want to create a parallel kernel. // We are running in parallel on multiple devices, so we may want to create a parallel kernel.
if (name == CalcForcesAndEnergyKernel::Name()) if (name == CalcForcesAndEnergyKernel::Name())
return new HipParallelCalcForcesAndEnergyKernel(name, platform, data); return new HipParallelCalcForcesAndEnergyKernel(name, platform, data);
if (name == CalcHarmonicBondForceKernel::Name()) if (name == CalcHarmonicBondForceKernel::Name())
return new HipParallelCalcHarmonicBondForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcHarmonicBondForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCustomBondForceKernel::Name()) if (name == CalcCustomBondForceKernel::Name())
return new HipParallelCalcCustomBondForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomBondForceKernel(name, platform, cu, context.getSystem());
if (name == CalcHarmonicAngleForceKernel::Name()) if (name == CalcHarmonicAngleForceKernel::Name())
return new HipParallelCalcHarmonicAngleForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcHarmonicAngleForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCustomAngleForceKernel::Name()) if (name == CalcCustomAngleForceKernel::Name())
return new HipParallelCalcCustomAngleForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomAngleForceKernel(name, platform, cu, context.getSystem());
if (name == CalcPeriodicTorsionForceKernel::Name()) if (name == CalcPeriodicTorsionForceKernel::Name())
return new HipParallelCalcPeriodicTorsionForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcPeriodicTorsionForceKernel(name, platform, cu, context.getSystem());
if (name == CalcRBTorsionForceKernel::Name()) if (name == CalcRBTorsionForceKernel::Name())
return new HipParallelCalcRBTorsionForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcRBTorsionForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCMAPTorsionForceKernel::Name()) if (name == CalcCMAPTorsionForceKernel::Name())
return new HipParallelCalcCMAPTorsionForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCMAPTorsionForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCustomTorsionForceKernel::Name()) if (name == CalcCustomTorsionForceKernel::Name())
return new HipParallelCalcCustomTorsionForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomTorsionForceKernel(name, platform, cu, context.getSystem());
if (name == CalcNonbondedForceKernel::Name()) if (name == CalcNonbondedForceKernel::Name())
return new HipParallelCalcNonbondedForceKernel(name, platform, data, context.getSystem()); return new HipParallelCalcNonbondedForceKernel(name, platform, data, context.getSystem());
if (name == CalcCustomNonbondedForceKernel::Name()) if (name == CalcCustomNonbondedForceKernel::Name())
return new HipParallelCalcCustomNonbondedForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomNonbondedForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCustomExternalForceKernel::Name()) if (name == CalcCustomExternalForceKernel::Name())
return new HipParallelCalcCustomExternalForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomExternalForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCustomHbondForceKernel::Name()) if (name == CalcCustomHbondForceKernel::Name())
return new HipParallelCalcCustomHbondForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomHbondForceKernel(name, platform, cu, context.getSystem());
if (name == CalcCustomCompoundBondForceKernel::Name()) if (name == CalcCustomCompoundBondForceKernel::Name())
return new HipParallelCalcCustomCompoundBondForceKernel(name, platform, data, context.getSystem()); return new CommonParallelCalcCustomCompoundBondForceKernel(name, platform, cu, context.getSystem());
} }
HipContext& cu = *data.contexts[0];
if (name == CalcForcesAndEnergyKernel::Name()) if (name == CalcForcesAndEnergyKernel::Name())
return new HipCalcForcesAndEnergyKernel(name, platform, cu); return new HipCalcForcesAndEnergyKernel(name, platform, cu);
if (name == UpdateStateDataKernel::Name()) if (name == UpdateStateDataKernel::Name())
......
platforms/hip/src/HipParallelKernels.cpp
View file @ b28d2e66
...@@ -28,6 +28,7 @@ ...@@ -28,6 +28,7 @@
#include "HipParallelKernels.h" #include "HipParallelKernels.h"
#include "HipKernelSources.h" #include "HipKernelSources.h"
#include "openmm/common/ContextSelector.h" #include "openmm/common/ContextSelector.h"
#include "openmm/internal/timer.h"
using namespace OpenMM; using namespace OpenMM;
using namespace std; using namespace std;
...@@ -40,28 +41,6 @@ if (result != hipSuccess) { \ ...@@ -40,28 +41,6 @@ if (result != hipSuccess) { \
throw OpenMMException(m.str());\ throw OpenMMException(m.str());\
} }
/**
* Get the current clock time, measured in microseconds.
*/
#ifdef _MSC_VER
#include <Windows.h>
static long long getTime() {
FILETIME ft;
GetSystemTimeAsFileTime(&ft); // 100-nanoseconds since 1-1-1601
ULARGE_INTEGER result;
result.LowPart = ft.dwLowDateTime;
result.HighPart = ft.dwHighDateTime;
return result.QuadPart/10;
}
#else
#include <sys/time.h>
static long long getTime() {
struct timeval tod;
gettimeofday(&tod, 0);
return 1000000*tod.tv_sec+tod.tv_usec;
}
#endif
class HipParallelCalcForcesAndEnergyKernel::BeginComputationTask : public HipContext::WorkTask { class HipParallelCalcForcesAndEnergyKernel::BeginComputationTask : public HipContext::WorkTask {
public: public:
BeginComputationTask(ContextImpl& context, HipContext& cu, HipCalcForcesAndEnergyKernel& kernel, BeginComputationTask(ContextImpl& context, HipContext& cu, HipCalcForcesAndEnergyKernel& kernel,
...@@ -92,7 +71,7 @@ private: ...@@ -92,7 +71,7 @@ private:
class HipParallelCalcForcesAndEnergyKernel::FinishComputationTask : public HipContext::WorkTask { class HipParallelCalcForcesAndEnergyKernel::FinishComputationTask : public HipContext::WorkTask {
public: public:
FinishComputationTask(ContextImpl& context, HipContext& cu, HipCalcForcesAndEnergyKernel& kernel, FinishComputationTask(ContextImpl& context, HipContext& cu, HipCalcForcesAndEnergyKernel& kernel,
bool includeForce, bool includeEnergy, int groups, double& energy, long long& completionTime, long long* pinnedMemory, HipArray& contextForces, bool includeForce, bool includeEnergy, int groups, double& energy, double& completionTime, long long* pinnedMemory, HipArray& contextForces,
bool& valid, hipStream_t stream, hipEvent_t event, hipEvent_t localEvent) : bool& valid, hipStream_t stream, hipEvent_t event, hipEvent_t localEvent) :
context(context), cu(cu), kernel(kernel), includeForce(includeForce), includeEnergy(includeEnergy), groups(groups), energy(energy), context(context), cu(cu), kernel(kernel), includeForce(includeForce), includeEnergy(includeEnergy), groups(groups), energy(energy),
completionTime(completionTime), pinnedMemory(pinnedMemory), contextForces(contextForces), valid(valid), completionTime(completionTime), pinnedMemory(pinnedMemory), contextForces(contextForces), valid(valid),
...@@ -107,7 +86,7 @@ public: ...@@ -107,7 +86,7 @@ public:
// Record timing information for load balancing. Since this takes time, only do it at the start of the simulation. // Record timing information for load balancing. Since this takes time, only do it at the start of the simulation.
CHECK_RESULT(hipStreamSynchronize(cu.getCurrentStream()), "Error synchronizing HIP context"); CHECK_RESULT(hipStreamSynchronize(cu.getCurrentStream()), "Error synchronizing HIP context");
completionTime = getTime(); completionTime = getCurrentTime();
} }
if (includeForce) { if (includeForce) {
if (cu.getContextIndex() > 0) { if (cu.getContextIndex() > 0) {
...@@ -133,7 +112,7 @@ private: ...@@ -133,7 +112,7 @@ private:
bool includeForce, includeEnergy; bool includeForce, includeEnergy;
int groups; int groups;
double& energy; double& energy;
long long& completionTime; double& completionTime;
long long* pinnedMemory; long long* pinnedMemory;
HipArray& contextForces; HipArray& contextForces;
bool& valid; bool& valid;
...@@ -280,334 +259,6 @@ double HipParallelCalcForcesAndEnergyKernel::finishComputation(ContextImpl& cont ...@@ -280,334 +259,6 @@ double HipParallelCalcForcesAndEnergyKernel::finishComputation(ContextImpl& cont
return energy; return energy;
} }
class HipParallelCalcHarmonicBondForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcHarmonicBondForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcHarmonicBondForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcHarmonicBondForceKernel::HipParallelCalcHarmonicBondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcHarmonicBondForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcHarmonicBondForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcHarmonicBondForceKernel::initialize(const System& system, const HarmonicBondForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcHarmonicBondForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcHarmonicBondForceKernel::copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force, int firstBond, int lastBond) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstBond, lastBond);
}
class HipParallelCalcCustomBondForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomBondForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomBondForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomBondForceKernel::HipParallelCalcCustomBondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomBondForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomBondForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomBondForceKernel::initialize(const System& system, const CustomBondForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomBondForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomBondForceKernel::copyParametersToContext(ContextImpl& context, const CustomBondForce& force, int firstBond, int lastBond) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstBond, lastBond);
}
class HipParallelCalcHarmonicAngleForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcHarmonicAngleForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcHarmonicAngleForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcHarmonicAngleForceKernel::HipParallelCalcHarmonicAngleForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcHarmonicAngleForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcHarmonicAngleForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcHarmonicAngleForceKernel::initialize(const System& system, const HarmonicAngleForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcHarmonicAngleForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcHarmonicAngleForceKernel::copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force, int firstAngle, int lastAngle) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstAngle, lastAngle);
}
class HipParallelCalcCustomAngleForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomAngleForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomAngleForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomAngleForceKernel::HipParallelCalcCustomAngleForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomAngleForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomAngleForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomAngleForceKernel::initialize(const System& system, const CustomAngleForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomAngleForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomAngleForceKernel::copyParametersToContext(ContextImpl& context, const CustomAngleForce& force, int firstAngle, int lastAngle) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstAngle, lastAngle);
}
class HipParallelCalcPeriodicTorsionForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcPeriodicTorsionForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcPeriodicTorsionForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcPeriodicTorsionForceKernel::HipParallelCalcPeriodicTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcPeriodicTorsionForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcPeriodicTorsionForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcPeriodicTorsionForceKernel::initialize(const System& system, const PeriodicTorsionForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcPeriodicTorsionForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcPeriodicTorsionForceKernel::copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force, int firstTorsion, int lastTorsion) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstTorsion, lastTorsion);
}
class HipParallelCalcRBTorsionForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcRBTorsionForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcRBTorsionForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcRBTorsionForceKernel::HipParallelCalcRBTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcRBTorsionForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcRBTorsionForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcRBTorsionForceKernel::initialize(const System& system, const RBTorsionForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcRBTorsionForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcRBTorsionForceKernel::copyParametersToContext(ContextImpl& context, const RBTorsionForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force);
}
class HipParallelCalcCMAPTorsionForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCMAPTorsionForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCMAPTorsionForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCMAPTorsionForceKernel::HipParallelCalcCMAPTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCMAPTorsionForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCMAPTorsionForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCMAPTorsionForceKernel::initialize(const System& system, const CMAPTorsionForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCMAPTorsionForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCMAPTorsionForceKernel::copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force);
}
class HipParallelCalcCustomTorsionForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomTorsionForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomTorsionForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomTorsionForceKernel::HipParallelCalcCustomTorsionForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomTorsionForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomTorsionForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomTorsionForceKernel::initialize(const System& system, const CustomTorsionForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomTorsionForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomTorsionForceKernel::copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force, int firstTorsion, int lastTorsion) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstTorsion, lastTorsion);
}
class HipParallelCalcNonbondedForceKernel::Task : public HipContext::WorkTask { class HipParallelCalcNonbondedForceKernel::Task : public HipContext::WorkTask {
public: public:
Task(ContextImpl& context, HipCalcNonbondedForceKernel& kernel, bool includeForce, Task(ContextImpl& context, HipCalcNonbondedForceKernel& kernel, bool includeForce,
...@@ -656,167 +307,3 @@ void HipParallelCalcNonbondedForceKernel::getPMEParameters(double& alpha, int& n ...@@ -656,167 +307,3 @@ void HipParallelCalcNonbondedForceKernel::getPMEParameters(double& alpha, int& n
void HipParallelCalcNonbondedForceKernel::getLJPMEParameters(double& alpha, int& nx, int& ny, int& nz) const { void HipParallelCalcNonbondedForceKernel::getLJPMEParameters(double& alpha, int& nx, int& ny, int& nz) const {
dynamic_cast<const HipCalcNonbondedForceKernel&>(kernels[0].getImpl()).getLJPMEParameters(alpha, nx, ny, nz); dynamic_cast<const HipCalcNonbondedForceKernel&>(kernels[0].getImpl()).getLJPMEParameters(alpha, nx, ny, nz);
} }
class HipParallelCalcCustomNonbondedForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomNonbondedForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomNonbondedForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomNonbondedForceKernel::HipParallelCalcCustomNonbondedForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomNonbondedForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomNonbondedForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomNonbondedForceKernel::initialize(const System& system, const CustomNonbondedForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomNonbondedForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomNonbondedForceKernel::copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force, int firstParticle, int lastParticle) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstParticle, lastParticle);
}
class HipParallelCalcCustomExternalForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomExternalForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomExternalForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomExternalForceKernel::HipParallelCalcCustomExternalForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomExternalForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomExternalForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomExternalForceKernel::initialize(const System& system, const CustomExternalForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomExternalForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomExternalForceKernel::copyParametersToContext(ContextImpl& context, const CustomExternalForce& force, int firstParticle, int lastParticle) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force, firstParticle, lastParticle);
}
class HipParallelCalcCustomHbondForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomHbondForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomHbondForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomHbondForceKernel::HipParallelCalcCustomHbondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomHbondForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomHbondForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomHbondForceKernel::initialize(const System& system, const CustomHbondForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomHbondForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomHbondForceKernel::copyParametersToContext(ContextImpl& context, const CustomHbondForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force);
}
class HipParallelCalcCustomCompoundBondForceKernel::Task : public HipContext::WorkTask {
public:
Task(ContextImpl& context, CommonCalcCustomCompoundBondForceKernel& kernel, bool includeForce,
bool includeEnergy, double& energy) : context(context), kernel(kernel),
includeForce(includeForce), includeEnergy(includeEnergy), energy(energy) {
}
void execute() {
energy += kernel.execute(context, includeForce, includeEnergy);
}
private:
ContextImpl& context;
CommonCalcCustomCompoundBondForceKernel& kernel;
bool includeForce, includeEnergy;
double& energy;
};
HipParallelCalcCustomCompoundBondForceKernel::HipParallelCalcCustomCompoundBondForceKernel(std::string name, const Platform& platform, HipPlatform::PlatformData& data, const System& system) :
CalcCustomCompoundBondForceKernel(name, platform), data(data) {
for (int i = 0; i < (int) data.contexts.size(); i++)
kernels.push_back(Kernel(new CommonCalcCustomCompoundBondForceKernel(name, platform, *data.contexts[i], system)));
}
void HipParallelCalcCustomCompoundBondForceKernel::initialize(const System& system, const CustomCompoundBondForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).initialize(system, force);
}
double HipParallelCalcCustomCompoundBondForceKernel::execute(ContextImpl& context, bool includeForces, bool includeEnergy) {
for (int i = 0; i < (int) data.contexts.size(); i++) {
HipContext& cu = *data.contexts[i];
ComputeContext::WorkThread& thread = cu.getWorkThread();
thread.addTask(new Task(context, getKernel(i), includeForces, includeEnergy, data.contextEnergy[i]));
}
return 0.0;
}
void HipParallelCalcCustomCompoundBondForceKernel::copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force) {
for (int i = 0; i < (int) kernels.size(); i++)
getKernel(i).copyParametersToContext(context, force);
}
platforms/hip/tests/TestHipCustomAngleForce.cpp
View file @ b28d2e66
...@@ -33,35 +33,6 @@ ...@@ -33,35 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestCustomAngleForce.h" #include "TestCustomAngleForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
CustomAngleForce* force = new CustomAngleForce("(theta-1.1)^2");
vector<double> params;
for (int i = 2; i < numParticles; i++)
force->addAngle(i-2, i-1, i, params);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, 0);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipCustomBondForce.cpp
View file @ b28d2e66
...@@ -33,35 +33,6 @@ ...@@ -33,35 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestCustomBondForce.h" #include "TestCustomBondForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
CustomBondForce* force = new CustomBondForce(("(r-1.1)^2"));
vector<double> params;
for (int i = 1; i < numParticles; i++)
force->addBond(i-1, i, params);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, 0, 0);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipCustomCompoundBondForce.cpp
View file @ b28d2e66
...@@ -33,39 +33,6 @@ ...@@ -33,39 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestCustomCompoundBondForce.h" #include "TestCustomCompoundBondForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
CustomCompoundBondForce* force = new CustomCompoundBondForce(2, ("(distance(p1,p2)-1.1)^2"));
vector<int> particles(2);
vector<double> params;
for (int i = 1; i < numParticles; i++) {
particles[0] = i-1;
particles[1] = i;
force->addBond(particles, params);
}
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, 0, 0);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipCustomExternalForce.cpp
View file @ b28d2e66
...@@ -34,37 +34,6 @@ ...@@ -34,37 +34,6 @@
#include "TestCustomExternalForce.h" #include "TestCustomExternalForce.h"
#include "sfmt/SFMT.h" #include "sfmt/SFMT.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
CustomExternalForce* force = new CustomExternalForce("x^2+y^2+z^2");
vector<double> params;
for (int i = 0; i < numParticles; i++)
force->addParticle(i, params);
system.addForce(force);
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(5*genrand_real2(sfmt), 5*genrand_real2(sfmt), 5*genrand_real2(sfmt));
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipCustomNonbondedForce.cpp
View file @ b28d2e66
...@@ -33,43 +33,6 @@ ...@@ -33,43 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestCustomNonbondedForce.h" #include "TestCustomNonbondedForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
CustomNonbondedForce* force = new CustomNonbondedForce("4*eps*((sigma/r)^12-(sigma/r)^6); sigma=0.5; eps=1");
vector<double> params;
for (int i = 0; i < numParticles; i++)
force->addParticle(params);
system.addForce(force);
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(5*genrand_real2(sfmt), 5*genrand_real2(sfmt), 5*genrand_real2(sfmt));
for (int i = 0; i < numParticles; ++i)
for (int j = 0; j < i; ++j) {
Vec3 delta = positions[i]-positions[j];
if (delta.dot(delta) < 0.1)
force->addExclusion(i, j);
}
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipCustomTorsionForce.cpp
View file @ b28d2e66
...@@ -33,35 +33,6 @@ ...@@ -33,35 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestCustomTorsionForce.h" #include "TestCustomTorsionForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
CustomTorsionForce* force = new CustomTorsionForce("sin(theta-1.1)");
vector<double> params;
for (int i = 3; i < numParticles; i++)
force->addTorsion(i-3, i-2, i-1, i, params);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, i%3);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipHarmonicAngleForce.cpp
View file @ b28d2e66
...@@ -33,35 +33,6 @@ ...@@ -33,35 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestHarmonicAngleForce.h" #include "TestHarmonicAngleForce.h"
#include <map>
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
HarmonicAngleForce* force = new HarmonicAngleForce();
for (int i = 2; i < numParticles; i++)
force->addAngle(i-2, i-1, i, 1.1, i);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, 0);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
......
platforms/hip/tests/TestHipHarmonicBondForce.cpp
View file @ b28d2e66
...@@ -32,35 +32,6 @@ ...@@ -32,35 +32,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestHarmonicBondForce.h" #include "TestHarmonicBondForce.h"
#include <map>
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
HarmonicBondForce* force = new HarmonicBondForce();
for (int i = 1; i < numParticles; i++)
force->addBond(i-1, i, 1.1, i);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, 0, 0);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
......
platforms/hip/tests/TestHipNonbondedForce.cpp
View file @ b28d2e66
...@@ -34,98 +34,6 @@ ...@@ -34,98 +34,6 @@
#include "TestNonbondedForce.h" #include "TestNonbondedForce.h"
#include <hip/hip_runtime.h> #include <hip/hip_runtime.h>
void testParallelComputation(NonbondedForce::NonbondedMethod method) {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
NonbondedForce* force = new NonbondedForce();
for (int i = 0; i < numParticles; i++)
force->addParticle(i%2-0.5, 0.5, 1.0);
force->setNonbondedMethod(method);
system.addForce(force);
system.setDefaultPeriodicBoxVectors(Vec3(5,0,0), Vec3(0,5,0), Vec3(0,0,5));
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(5*genrand_real2(sfmt), 5*genrand_real2(sfmt), 5*genrand_real2(sfmt));
force->addGlobalParameter("scale", 0.5);
for (int i = 0; i < numParticles; ++i)
for (int j = 0; j < i; ++j) {
Vec3 delta = positions[i]-positions[j];
if (delta.dot(delta) < 0.1) {
force->addException(i, j, 0, 1, 0);
}
else if (delta.dot(delta) < 0.2) {
int index = force->addException(i, j, 0.5, 1, 1.0);
force->addExceptionParameterOffset("scale", index, 0.5, 0.4, 0.3);
}
}
// Create two contexts, one with a single device and one with two devices.
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
// See if they agree.
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
// Modify some particle parameters and see if they still agree.
for (int i = 0; i < numParticles; i += 5) {
double charge, sigma, epsilon;
force->getParticleParameters(i, charge, sigma, epsilon);
force->setParticleParameters(i, 0.9*charge, sigma, epsilon);
}
force->updateParametersInContext(context1);
force->updateParametersInContext(context2);
state1 = context1.getState(State::Forces | State::Energy);
state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void testReordering() {
// Check that reordering of atoms doesn't alter their positions.
const int numParticles = 200;
System system;
system.setDefaultPeriodicBoxVectors(Vec3(6, 0, 0), Vec3(2.1, 6, 0), Vec3(-1.5, -0.5, 6));
NonbondedForce *nonbonded = new NonbondedForce();
nonbonded->setNonbondedMethod(NonbondedForce::PME);
system.addForce(nonbonded);
vector<Vec3> positions;
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
for (int i = 0; i < numParticles; i++) {
system.addParticle(1.0);
nonbonded->addParticle(0.0, 0.0, 0.0);
positions.push_back(Vec3(genrand_real2(sfmt)-0.5, genrand_real2(sfmt)-0.5, genrand_real2(sfmt)-0.5)*20);
}
VerletIntegrator integrator(0.001);
Context context(system, integrator, platform);
context.setPositions(positions);
integrator.step(1);
State state = context.getState(State::Positions | State::Velocities);
for (int i = 0; i < numParticles; i++) {
ASSERT_EQUAL_VEC(positions[i], state.getPositions()[i], 1e-6);
}
}
void testDeterministicForces() { void testDeterministicForces() {
// Check that the HipDeterministicForces property works correctly. // Check that the HipDeterministicForces property works correctly.
......
platforms/hip/tests/TestHipPeriodicTorsionForce.cpp
View file @ b28d2e66
...@@ -33,34 +33,6 @@ ...@@ -33,34 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestPeriodicTorsionForce.h" #include "TestPeriodicTorsionForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
PeriodicTorsionForce* force = new PeriodicTorsionForce();
for (int i = 3; i < numParticles; i++)
force->addTorsion(i-3, i-2, i-1, i, 2, 1.1, i);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, i%3);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/hip/tests/TestHipRBTorsionForce.cpp
View file @ b28d2e66
...@@ -33,34 +33,6 @@ ...@@ -33,34 +33,6 @@
#include "HipTests.h" #include "HipTests.h"
#include "TestRBTorsionForce.h" #include "TestRBTorsionForce.h"
void testParallelComputation() {
System system;
const int numParticles = 200;
for (int i = 0; i < numParticles; i++)
system.addParticle(1.0);
RBTorsionForce* force = new RBTorsionForce();
for (int i = 3; i < numParticles; i++)
force->addTorsion(i-3, i-2, i-1, i, 2, 0.1*i, 0.5*i, i, 1, 1);
system.addForce(force);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numParticles; i++)
positions[i] = Vec3(i, i%2, i%3);
VerletIntegrator integrator1(0.01);
Context context1(system, integrator1, platform);
context1.setPositions(positions);
State state1 = context1.getState(State::Forces | State::Energy);
VerletIntegrator integrator2(0.01);
string deviceIndex = platform.getPropertyValue(context1, HipPlatform::HipDeviceIndex());
map<string, string> props;
props[HipPlatform::HipDeviceIndex()] = deviceIndex+","+deviceIndex;
Context context2(system, integrator2, platform, props);
context2.setPositions(positions);
State state2 = context2.getState(State::Forces | State::Energy);
ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);
for (int i = 0; i < numParticles; i++)
ASSERT_EQUAL_VEC(state1.getForces()[i], state2.getForces()[i], 1e-5);
}
void runPlatformTests() { void runPlatformTests() {
testParallelComputation(); testParallelComputation();
} }
platforms/opencl/include/OpenCLContext.h
View file @ b28d2e66
...@@ -197,6 +197,11 @@ public: ...@@ -197,6 +197,11 @@ public:
* one ComputeContext is created for each device. * one ComputeContext is created for each device.
*/ */
std::vector<ComputeContext*> getAllContexts(); std::vector<ComputeContext*> getAllContexts();
/**
* Get a workspace used for accumulating energy when a simulation is parallelized across
* multiple devices.
*/
double& getEnergyWorkspace();
/** /**
* Get the cl::CommandQueue currently being used for execution. * Get the cl::CommandQueue currently being used for execution.
*/ */
......
platforms/opencl/include/OpenCLParallelKernels.h
View file @ b28d2e66
...@@ -82,7 +82,7 @@ private: ...@@ -82,7 +82,7 @@ private:
class FinishComputationTask; class FinishComputationTask;
OpenCLPlatform::PlatformData& data; OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels; std::vector<Kernel> kernels;
std::vector<long long> completionTimes; std::vector<double> completionTimes;
std::vector<double> contextNonbondedFractions; std::vector<double> contextNonbondedFractions;
std::vector<int> tileCounts; std::vector<int> tileCounts;
OpenCLArray contextForces; OpenCLArray contextForces;
...@@ -92,314 +92,6 @@ private: ...@@ -92,314 +92,6 @@ private:
void* pinnedForceMemory; void* pinnedForceMemory;
}; };
/**
* This kernel is invoked by HarmonicBondForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcHarmonicBondForceKernel : public CalcHarmonicBondForceKernel {
public:
OpenCLParallelCalcHarmonicBondForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcHarmonicBondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcHarmonicBondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the HarmonicBondForce this kernel will be used for
*/
void initialize(const System& system, const HarmonicBondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param firstBond the index of the first bond whose parameters might have changed
* @param lastBond the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force, int firstBond, int lastBond);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomBondForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcCustomBondForceKernel : public CalcCustomBondForceKernel {
public:
OpenCLParallelCalcCustomBondForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomBondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomBondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomBondForce this kernel will be used for
*/
void initialize(const System& system, const CustomBondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param firstBond the index of the first bond whose parameters might have changed
* @param lastBond the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomBondForce& force, int firstBond, int lastBond);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by HarmonicAngleForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcHarmonicAngleForceKernel : public CalcHarmonicAngleForceKernel {
public:
OpenCLParallelCalcHarmonicAngleForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcHarmonicAngleForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcHarmonicAngleForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the HarmonicAngleForce this kernel will be used for
*/
void initialize(const System& system, const HarmonicAngleForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param firstAngle the index of the first bond whose parameters might have changed
* @param lastAngle the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force, int firstAngle, int lastAngle);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomAngleForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcCustomAngleForceKernel : public CalcCustomAngleForceKernel {
public:
OpenCLParallelCalcCustomAngleForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomAngleForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomAngleForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomAngleForce this kernel will be used for
*/
void initialize(const System& system, const CustomAngleForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param firstAngle the index of the first bond whose parameters might have changed
* @param lastAngle the index of the last bond whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomAngleForce& force, int firstAngle, int lastAngle);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by PeriodicTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcPeriodicTorsionForceKernel : public CalcPeriodicTorsionForceKernel {
public:
OpenCLParallelCalcPeriodicTorsionForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcPeriodicTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcPeriodicTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the PeriodicTorsionForce this kernel will be used for
*/
void initialize(const System& system, const PeriodicTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
class Task;
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the PeriodicTorsionForce to copy the parameters from
* @param firstTorsion the index of the first torsion whose parameters might have changed
* @param lastTorsion the index of the last torsion whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force, int firstTorsion, int lastTorsion);
private:
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by RBTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcRBTorsionForceKernel : public CalcRBTorsionForceKernel {
public:
OpenCLParallelCalcRBTorsionForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcRBTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcRBTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the RBTorsionForce this kernel will be used for
*/
void initialize(const System& system, const RBTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the RBTorsionForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const RBTorsionForce& force);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CMAPTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcCMAPTorsionForceKernel : public CalcCMAPTorsionForceKernel {
public:
OpenCLParallelCalcCMAPTorsionForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCMAPTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCMAPTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CMAPTorsionForce this kernel will be used for
*/
void initialize(const System& system, const CMAPTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CMAPTorsionForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomTorsionForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcCustomTorsionForceKernel : public CalcCustomTorsionForceKernel {
public:
OpenCLParallelCalcCustomTorsionForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomTorsionForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomTorsionForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomTorsionForce this kernel will be used for
*/
void initialize(const System& system, const CustomTorsionForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomTorsionForce to copy the parameters from
* @param firstTorsion the index of the first torsion whose parameters might have changed
* @param lastTorsion the index of the last torsion whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force, int firstTorsion, int lastTorsion);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/** /**
* This kernel is invoked by NonbondedForce to calculate the forces acting on the system. * This kernel is invoked by NonbondedForce to calculate the forces acting on the system.
*/ */
...@@ -462,162 +154,6 @@ private: ...@@ -462,162 +154,6 @@ private:
std::vector<Kernel> kernels; std::vector<Kernel> kernels;
}; };
/**
* This kernel is invoked by CustomNonbondedForce to calculate the forces acting on the system.
*/
class OpenCLParallelCalcCustomNonbondedForceKernel : public CalcCustomNonbondedForceKernel {
public:
OpenCLParallelCalcCustomNonbondedForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomNonbondedForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomNonbondedForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomNonbondedForce this kernel will be used for
*/
void initialize(const System& system, const CustomNonbondedForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomNonbondedForce to copy the parameters from
* @param firstParticle the index of the first particle whose parameters might have changed
* @param lastParticle the index of the last particle whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force, int firstParticle, int lastParticle);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomExternalForce to calculate the forces acting on the system and the energy of the system.
*/
class OpenCLParallelCalcCustomExternalForceKernel : public CalcCustomExternalForceKernel {
public:
OpenCLParallelCalcCustomExternalForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomExternalForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomExternalForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomExternalForce this kernel will be used for
*/
void initialize(const System& system, const CustomExternalForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomExternalForce to copy the parameters from
* @param firstParticle the index of the first particle whose parameters might have changed
* @param lastParticle the index of the last particle whose parameters might have changed
*/
void copyParametersToContext(ContextImpl& context, const CustomExternalForce& force, int firstParticle, int lastParticle);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomHbondForce to calculate the forces acting on the system.
*/
class OpenCLParallelCalcCustomHbondForceKernel : public CalcCustomHbondForceKernel {
public:
OpenCLParallelCalcCustomHbondForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomHbondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomHbondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomHbondForce this kernel will be used for
*/
void initialize(const System& system, const CustomHbondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomHbondForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const CustomHbondForce& force);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
/**
* This kernel is invoked by CustomCompoundBondForce to calculate the forces acting on the system.
*/
class OpenCLParallelCalcCustomCompoundBondForceKernel : public CalcCustomCompoundBondForceKernel {
public:
OpenCLParallelCalcCustomCompoundBondForceKernel(std::string name, const Platform& platform, OpenCLPlatform::PlatformData& data, const System& system);
CommonCalcCustomCompoundBondForceKernel& getKernel(int index) {
return dynamic_cast<CommonCalcCustomCompoundBondForceKernel&>(kernels[index].getImpl());
}
/**
* Initialize the kernel.
*
* @param system the System this kernel will be applied to
* @param force the CustomCompoundBondForce this kernel will be used for
*/
void initialize(const System& system, const CustomCompoundBondForce& force);
/**
* Execute the kernel to calculate the forces and/or energy.
*
* @param context the context in which to execute this kernel
* @param includeForces true if forces should be calculated
* @param includeEnergy true if the energy should be calculated
* @return the potential energy due to the force
*/
double execute(ContextImpl& context, bool includeForces, bool includeEnergy);
/**
* Copy changed parameters over to a context.
*
* @param context the context to copy parameters to
* @param force the CustomCompoundBondForce to copy the parameters from
*/
void copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force);
private:
class Task;
OpenCLPlatform::PlatformData& data;
std::vector<Kernel> kernels;
};
} // namespace OpenMM } // namespace OpenMM
#endif /*OPENMM_OPENCLPARALLELKERNELS_H_*/ #endif /*OPENMM_OPENCLPARALLELKERNELS_H_*/
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