Skip to content

GitLab

Commit d56c21b0 authored by peastman's avatar peastman
Browse files

Merged latest changes from main branch

parents 1049add2 8d062ebb
Show whitespace changes
Inline Side-by-side
Showing with 509 additions and 112 deletions
plugins/amoeba/platforms/reference/tests/TestReferenceWcaDispersionForce.cpp
View file @ d56c21b0
......@@ -50,6 +50,9 @@
using namespace OpenMM;
extern "C" OPENMM_EXPORT void registerAmoebaReferenceKernelFactories();
const double TOL = 1e-4;
void compareForcesEnergy( std::string& testName, double expectedEnergy, double energy,
......@@ -171,6 +174,7 @@ int main( int numberOfArguments, char* argv[] ) {
try {
std::cout << "TestCudaAmoebaWcaDispersionForce running test..." << std::endl;
registerAmoebaReferenceKernelFactories();
FILE* log = NULL;
......
plugins/drude/CMakeLists.txt
View file @ d56c21b0
......@@ -18,7 +18,7 @@
# The source is organized into subdirectories, but we handle them all from
# this CMakeLists file rather than letting CMake visit them as SUBDIRS.
SET(OPENMM_DRUDE_PLUGIN_SOURCE_SUBDIRS . openmmapi platforms/reference serialization)
SET(OPENMM_DRUDE_PLUGIN_SOURCE_SUBDIRS . openmmapi serialization)
SET(OPENMM_DRUDE_LIBRARY_NAME OpenMMDrude)
SET(SHARED_DRUDE_TARGET ${OPENMM_DRUDE_LIBRARY_NAME})
......@@ -70,20 +70,8 @@ FOREACH(subdir ${OPENMM_DRUDE_PLUGIN_SOURCE_SUBDIRS})
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include)
ENDFOREACH(subdir)
INCLUDE_DIRECTORIES(BEFORE ${OPENMM_DIR}/platforms/reference/src)
INCLUDE_DIRECTORIES(BEFORE ${OPENMM_DIR}/platforms/reference/src/SimTKReference)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/platforms/reference/src)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/platforms/reference/src/SimTKReference)
# ----------------------------------------------------------------------------
# If API_DRUDE wrappers are being generated, and add them to the build.
#IF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
# ADD_SUBDIRECTORY(wrappers)
# SET(SOURCE_DRUDE_FILES ${SOURCE_DRUDE_FILES} wrappers/DrudeOpenMMCWrapper.cpp wrappers/DrudeOpenMMFortranWrapper.cpp)
# SET_SOURCE_FILES_PROPERTIES(wrappers/DrudeOpenMMCWrapper.cpp wrappers/DrudeOpenMMFortranWrapper.cpp PROPERTIES GENERATED TRUE)
#ENDIF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/src)
ADD_LIBRARY(${SHARED_DRUDE_TARGET} SHARED ${SOURCE_DRUDE_FILES} ${SOURCE_DRUDE_INCLUDE_FILES} ${API_DRUDE_ABS_INCLUDE_FILES})
......@@ -97,16 +85,6 @@ IF(OPENMM_BUILD_STATIC_LIB)
SET_TARGET_PROPERTIES(${STATIC_DRUDE_TARGET} PROPERTIES LINK_FLAGS "${EXTRA_COMPILE_FLAGS}" COMPILE_FLAGS "${EXTRA_COMPILE_FLAGS} -DOPENMM_USE_STATIC_LIBRARIES -DOPENMM_BUILDING_STATIC_LIBRARY -DLEPTON_USE_STATIC_LIBRARIES -DLEPTON_BUILDING_STATIC_LIBRARY")
ENDIF(OPENMM_BUILD_STATIC_LIB)
#IF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
# ADD_DEPENDENCIES(${SHARED_DRUDE_TARGET} DrudeApiWrappers)
# IF( CREATE_SERIALIZABLE_OPENMM_DRUDE )
# ADD_DEPENDENCIES(${SHARED_DRUDE_SERIALIZABLE_TARGET} DrudeApiWrappers)
# ENDIF( CREATE_SERIALIZABLE_OPENMM_DRUDE )
# IF(OPENMM_BUILD_STATIC_LIB)
# ADD_DEPENDENCIES(${STATIC_DRUDE_TARGET} DrudeApiWrappers)
# ENDIF(OPENMM_BUILD_STATIC_LIB)
#ENDIF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
# ----------------------------------------------------------------------------
# On Linux need to link to libdl
......@@ -132,18 +110,9 @@ IF(OPENMM_BUILD_STATIC_LIB)
TARGET_LINK_LIBRARIES( ${STATIC_DRUDE_TARGET} ${STATIC_TARGET} )
ENDIF(OPENMM_BUILD_STATIC_LIB)
ADD_SUBDIRECTORY(platforms/reference/tests)
# Which hardware platforms to build
#SET(OPENMM_BUILD_DRUDE_PATH)
#SET(OPENMM_BUILD_DRUDE_CUDA_PATH)
#IF(OPENMM_BUILD_DRUDE_CUDA_LIB)
# ADD_SUBDIRECTORY(platforms/cuda)
# SET(OPENMM_BUILD_DRUDE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/platforms/cuda)
# SET(OPENMM_BUILD_DRUDE_CUDA_PATH ${CMAKE_CURRENT_SOURCE_DIR}/platforms/cuda)
# SET(OPENMM_DRUDE_CUDA_SOURCE_SUBDIRS . openmmapi olla platforms/cuda)
#ENDIF(OPENMM_BUILD_DRUDE_CUDA_LIB)
ADD_SUBDIRECTORY(platforms/reference)
IF(OPENCL_FOUND)
SET(OPENMM_BUILD_DRUDE_OPENCL_LIB ON CACHE BOOL "Build Drude implementation for OpenCL")
......
plugins/drude/platforms/reference/CMakeLists.txt 0 → 100644
View file @ d56c21b0
#---------------------------------------------------
# OpenMM Reference Drude Integrator
#
# Creates OpenMM library, base name=OpenMMDrudeReference.
# Default libraries are shared & optimized. Variants
# are created for debug (_d).
#
# Windows:
# OpenMMDrudeReference[_d].dll
# OpenMMDrudeReference[_d].lib
# Unix:
# libOpenMMDrudeReference[_d].so
#----------------------------------------------------
# The source is organized into subdirectories, but we handle them all from
# this CMakeLists file rather than letting CMake visit them as SUBDIRS.
SET(OPENMM_SOURCE_SUBDIRS .)
# Collect up information about the version of the OpenMM library we're building
# and make it available to the code so it can be built into the binaries.
SET(OPENMMDRUDEREFERENCE_LIBRARY_NAME OpenMMDrudeReference)
SET(SHARED_TARGET ${OPENMMDRUDEREFERENCE_LIBRARY_NAME})
# Ensure that debug libraries have "_d" appended to their names.
# CMake gets this right on Windows automatically with this definition.
IF (${CMAKE_GENERATOR} MATCHES "Visual Studio")
SET(CMAKE_DEBUG_POSTFIX "_d" CACHE INTERNAL "" FORCE)
ENDIF (${CMAKE_GENERATOR} MATCHES "Visual Studio")
# But on Unix or Cygwin we have to add the suffix manually
IF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
SET(SHARED_TARGET ${SHARED_TARGET}_d)
ENDIF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
# These are all the places to search for header files which are
# to be part of the API.
SET(API_INCLUDE_DIRS) # start empty
FOREACH(subdir ${OPENMM_SOURCE_SUBDIRS})
# append
SET(API_INCLUDE_DIRS ${API_INCLUDE_DIRS}
${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include
${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include/internal)
ENDFOREACH(subdir)
# We'll need both *relative* path names, starting with their API_INCLUDE_DIRS,
# and absolute pathnames.
SET(API_REL_INCLUDE_FILES) # start these out empty
SET(API_ABS_INCLUDE_FILES)
FOREACH(dir ${API_INCLUDE_DIRS})
FILE(GLOB fullpaths ${dir}/*.h) # returns full pathnames
SET(API_ABS_INCLUDE_FILES ${API_ABS_INCLUDE_FILES} ${fullpaths})
FOREACH(pathname ${fullpaths})
GET_FILENAME_COMPONENT(filename ${pathname} NAME)
SET(API_REL_INCLUDE_FILES ${API_REL_INCLUDE_FILES} ${dir}/${filename})
ENDFOREACH(pathname)
ENDFOREACH(dir)
# collect up source files
SET(SOURCE_FILES) # empty
SET(SOURCE_INCLUDE_FILES)
FOREACH(subdir ${OPENMM_SOURCE_SUBDIRS})
FILE(GLOB_RECURSE src_files ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/src/*.cpp ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/src/*.c)
FILE(GLOB incl_files ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/src/*.h)
SET(SOURCE_FILES ${SOURCE_FILES} ${src_files}) #append
SET(SOURCE_INCLUDE_FILES ${SOURCE_INCLUDE_FILES} ${incl_files})
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include)
ENDFOREACH(subdir)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/src)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_SOURCE_DIR}/platforms/reference/include)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_SOURCE_DIR}/platforms/reference/src)
# Create the library
INCLUDE_DIRECTORIES(${REFERENCE_INCLUDE_DIR})
ADD_LIBRARY(${SHARED_TARGET} SHARED ${SOURCE_FILES} ${SOURCE_INCLUDE_FILES} ${API_ABS_INCLUDE_FILES})
IF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
SET(MAIN_OPENMM_LIB ${OPENMM_LIBRARY_NAME}_d)
ELSE (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
SET(MAIN_OPENMM_LIB ${OPENMM_LIBRARY_NAME})
ENDIF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
TARGET_LINK_LIBRARIES(${SHARED_TARGET} ${MAIN_OPENMM_LIB})
TARGET_LINK_LIBRARIES(${SHARED_TARGET} debug ${SHARED_DRUDE_TARGET} optimized ${SHARED_DRUDE_TARGET})
SET_TARGET_PROPERTIES(${SHARED_TARGET} PROPERTIES LINK_FLAGS "${EXTRA_COMPILE_FLAGS}" COMPILE_FLAGS "${EXTRA_COMPILE_FLAGS} -msse2 -DOPENMM_BUILDING_SHARED_LIBRARY")
INSTALL(TARGETS ${SHARED_TARGET} DESTINATION ${CMAKE_INSTALL_PREFIX}/lib/plugins)
SUBDIRS (tests)
plugins/drude/platforms/reference/src/ReferenceDrudeKernelFactory.cpp
View file @ d56c21b0
......@@ -32,24 +32,23 @@
using namespace OpenMM;
#if defined(WIN32)
#include <windows.h>
extern "C" void initDrudeReferenceKernels();
BOOL WINAPI DllMain(HANDLE hModule, DWORD ul_reason_for_call, LPVOID lpReserved) {
if (ul_reason_for_call == DLL_PROCESS_ATTACH)
initDrudeReferenceKernels();
return TRUE;
}
#else
extern "C" void __attribute__((constructor)) initDrudeReferenceKernels();
#endif
extern "C" OPENMM_EXPORT void registerPlatforms() {
}
extern "C" void initDrudeReferenceKernels() {
Platform& platform = Platform::getPlatformByName("Reference");
extern "C" OPENMM_EXPORT void registerKernelFactories() {
for (int i = 0; i < Platform::getNumPlatforms(); i++) {
Platform& platform = Platform::getPlatform(i);
if (dynamic_cast<ReferencePlatform*>(&platform) != NULL) {
ReferenceDrudeKernelFactory* factory = new ReferenceDrudeKernelFactory();
platform.registerKernelFactory(CalcDrudeForceKernel::Name(), factory);
platform.registerKernelFactory(IntegrateDrudeLangevinStepKernel::Name(), factory);
platform.registerKernelFactory(IntegrateDrudeSCFStepKernel::Name(), factory);
}
}
}
extern "C" OPENMM_EXPORT void registerDrudeReferenceKernelFactories() {
registerKernelFactories();
}
KernelImpl* ReferenceDrudeKernelFactory::createKernelImpl(std::string name, const Platform& platform, ContextImpl& context) const {
......
plugins/drude/platforms/reference/tests/CMakeLists.txt
View file @ d56c21b0
......@@ -2,11 +2,9 @@
# Testing
#
ENABLE_TESTING()
#INCLUDE_DIRECTORIES(${CUDA_INCLUDE})
INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/reference/include)
INCLUDE_DIRECTORIES(${OPENMM_DIR}/openmmapi/include/openmm)
INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/reference/src)
#INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/reference/src/kernels)
SET(SHARED_OPENMM_DRUDE_TARGET OpenMMDrude)
......
plugins/drude/platforms/reference/tests/TestReferenceDrudeForce.cpp
View file @ d56c21b0
......@@ -47,6 +47,8 @@
using namespace OpenMM;
using namespace std;
extern "C" OPENMM_EXPORT void registerDrudeReferenceKernelFactories();
void validateForce(System& system, vector<Vec3>& positions, double expectedEnergy) {
// Given a System containing a Drude force, check that its energy has the expected value.
......@@ -193,6 +195,7 @@ void testChangingParameters() {
int main() {
try {
registerDrudeReferenceKernelFactories();
testSingleParticle();
testAnisotropicParticle();
testThole();
......
plugins/drude/platforms/reference/tests/TestReferenceDrudeLangevinIntegrator.cpp
View file @ d56c21b0
......@@ -48,6 +48,8 @@
using namespace OpenMM;
using namespace std;
extern "C" OPENMM_EXPORT void registerDrudeReferenceKernelFactories();
void testSinglePair() {
const double temperature = 300.0;
const double temperatureDrude = 10.0;
......@@ -174,6 +176,7 @@ void testWater() {
int main() {
try {
registerDrudeReferenceKernelFactories();
testSinglePair();
testWater();
}
......
plugins/drude/platforms/reference/tests/TestReferenceDrudeSCFIntegrator.cpp
View file @ d56c21b0
......@@ -48,6 +48,8 @@
using namespace OpenMM;
using namespace std;
extern "C" OPENMM_EXPORT void registerDrudeReferenceKernelFactories();
void testWater() {
// Create a box of SWM4-NDP water molecules. This involves constraints, virtual sites,
// and Drude particles.
......@@ -126,6 +128,7 @@ void testWater() {
int main() {
try {
registerDrudeReferenceKernelFactories();
testWater();
}
catch(const std::exception& e) {
......
plugins/rpmd/CMakeLists.txt
View file @ d56c21b0
......@@ -22,7 +22,7 @@ SET(CREATE_SERIALIZABLE_OPENMM_RPMD FALSE )
# The source is organized into subdirectories, but we handle them all from
# this CMakeLists file rather than letting CMake visit them as SUBDIRS.
SET(OPENMM_RPMD_PLUGIN_SOURCE_SUBDIRS . openmmapi platforms/reference)
SET(OPENMM_RPMD_PLUGIN_SOURCE_SUBDIRS . openmmapi)
SET(OPENMM_RPMD_LIBRARY_NAME OpenMMRPMD)
SET(SHARED_RPMD_TARGET ${OPENMM_RPMD_LIBRARY_NAME})
......@@ -80,46 +80,18 @@ FOREACH(subdir ${OPENMM_RPMD_PLUGIN_SOURCE_SUBDIRS})
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include)
ENDFOREACH(subdir)
INCLUDE_DIRECTORIES(BEFORE ${OPENMM_DIR}/platforms/reference/src)
INCLUDE_DIRECTORIES(BEFORE ${OPENMM_DIR}/platforms/reference/src/SimTKReference)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/platforms/reference/src)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/platforms/reference/src/SimTKReference)
# ----------------------------------------------------------------------------
# If API_RPMD wrappers are being generated, and add them to the build.
#IF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
# ADD_SUBDIRECTORY(wrappers)
# SET(SOURCE_RPMD_FILES ${SOURCE_RPMD_FILES} wrappers/RPMDOpenMMCWrapper.cpp wrappers/RPMDOpenMMFortranWrapper.cpp)
# SET_SOURCE_FILES_PROPERTIES(wrappers/RPMDOpenMMCWrapper.cpp wrappers/RPMDOpenMMFortranWrapper.cpp PROPERTIES GENERATED TRUE)
#ENDIF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/src)
ADD_LIBRARY(${SHARED_RPMD_TARGET} SHARED ${SOURCE_RPMD_FILES} ${SOURCE_RPMD_INCLUDE_FILES} ${API_RPMD_ABS_INCLUDE_FILES})
SET_TARGET_PROPERTIES(${SHARED_RPMD_TARGET} PROPERTIES LINK_FLAGS "${EXTRA_COMPILE_FLAGS}" COMPILE_FLAGS "${EXTRA_COMPILE_FLAGS} -DOPENMM_RPMD_BUILDING_SHARED_LIBRARY -DLEPTON_BUILDING_SHARED_LIBRARY")
#IF( CREATE_SERIALIZABLE_OPENMM_RPMD )
# ADD_LIBRARY(${SHARED_RPMD_SERIALIZABLE_TARGET} SHARED ${SOURCE_RPMD_FILES} ${SOURCE_RPMD_INCLUDE_FILES} ${API_RPMD_ABS_INCLUDE_FILES})
# SET_TARGET_PROPERTIES(${SHARED_RPMD_SERIALIZABLE_TARGET} PROPERTIES COMPILE_FLAGS "-DOPENMM_RPMD_BUILDING_SHARED_LIBRARY -DLEPTON_BUILDING_SHARED_LIBRARY -DOPENMM_SERIALIZE")
# INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/../../serialization/include)
#ENDIF( CREATE_SERIALIZABLE_OPENMM_RPMD )
IF(OPENMM_BUILD_STATIC_LIB)
ADD_LIBRARY(${STATIC_RPMD_TARGET} STATIC ${SOURCE_RPMD_FILES} ${SOURCE_RPMD_INCLUDE_FILES} ${API_RPMD_ABS_INCLUDE_FILES})
SET_TARGET_PROPERTIES(${STATIC_RPMD_TARGET} PROPERTIES LINK_FLAGS "${EXTRA_COMPILE_FLAGS}" COMPILE_FLAGS "${EXTRA_COMPILE_FLAGS} -DOPENMM_USE_STATIC_LIBRARIES -DOPENMM_BUILDING_STATIC_LIBRARY -DLEPTON_USE_STATIC_LIBRARIES -DLEPTON_BUILDING_STATIC_LIBRARY")
ENDIF(OPENMM_BUILD_STATIC_LIB)
#IF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
# ADD_DEPENDENCIES(${SHARED_RPMD_TARGET} RPMDApiWrappers)
# IF( CREATE_SERIALIZABLE_OPENMM_RPMD )
# ADD_DEPENDENCIES(${SHARED_RPMD_SERIALIZABLE_TARGET} RPMDApiWrappers)
# ENDIF( CREATE_SERIALIZABLE_OPENMM_RPMD )
# IF(OPENMM_BUILD_STATIC_LIB)
# ADD_DEPENDENCIES(${STATIC_RPMD_TARGET} RPMDApiWrappers)
# ENDIF(OPENMM_BUILD_STATIC_LIB)
#ENDIF(OPENMM_BUILD_C_AND_FORTRAN_WRAPPERS)
# ----------------------------------------------------------------------------
# On Linux need to link to libdl
......@@ -148,18 +120,9 @@ IF(OPENMM_BUILD_STATIC_LIB)
TARGET_LINK_LIBRARIES( ${STATIC_RPMD_TARGET} ${STATIC_TARGET} )
ENDIF(OPENMM_BUILD_STATIC_LIB)
ADD_SUBDIRECTORY(platforms/reference/tests)
# Which hardware platforms to build
#SET(OPENMM_BUILD_RPMD_PATH)
#SET(OPENMM_BUILD_RPMD_CUDA_PATH)
#IF(OPENMM_BUILD_RPMD_CUDA_LIB)
# ADD_SUBDIRECTORY(platforms/cuda)
# SET(OPENMM_BUILD_RPMD_PATH ${CMAKE_CURRENT_SOURCE_DIR}/platforms/cuda)
# SET(OPENMM_BUILD_RPMD_CUDA_PATH ${CMAKE_CURRENT_SOURCE_DIR}/platforms/cuda)
# SET(OPENMM_RPMD_CUDA_SOURCE_SUBDIRS . openmmapi olla platforms/cuda)
#ENDIF(OPENMM_BUILD_RPMD_CUDA_LIB)
ADD_SUBDIRECTORY(platforms/reference)
IF(OPENCL_FOUND)
SET(OPENMM_BUILD_RPMD_OPENCL_LIB ON CACHE BOOL "Build RPMD implementation for OpenCL")
......@@ -208,7 +171,3 @@ ENDIF (EXECUTABLE_OUTPUT_PATH)
#IF (OPENMM_BUILD_SERIALIZATION_SUPPORT)
# ADD_SUBDIRECTORY(serialization)
#ENDIF (OPENMM_BUILD_SERIALIZATION_SUPPORT)
#INCLUDE(ApiDoxygen.cmake)
#ADD_SUBDIRECTORY(tests)
plugins/rpmd/platforms/cuda/src/CudaRpmdKernels.cpp
View file @ d56c21b0
......@@ -252,7 +252,18 @@ void CudaIntegrateRPMDStepKernel::computeForces(ContextImpl& context) {
&cu.getPosq().getDevicePointer(), &cu.getAtomIndexArray().getDevicePointer(), &i};
cu.executeKernel(copyToContextKernel, copyToContextArgs, cu.getNumAtoms());
context.computeVirtualSites();
Vec3 initialBox[3];
context.getPeriodicBoxVectors(initialBox[0], initialBox[1], initialBox[2]);
context.updateContextState();
Vec3 finalBox[3];
context.getPeriodicBoxVectors(finalBox[0], finalBox[1], finalBox[2]);
if (initialBox[0] != finalBox[0] || initialBox[1] != finalBox[1] || initialBox[2] != finalBox[2]) {
// A barostat was applied during updateContextState(). Adjust the particle positions in all the
// other copies to match this one.
void* args[] = {&positions->getDevicePointer(), &cu.getPosq().getDevicePointer(), &cu.getAtomIndexArray().getDevicePointer(), &i};
cu.executeKernel(translateKernel, args, cu.getNumAtoms());
}
context.calcForcesAndEnergy(true, false, groupsNotContracted);
void* copyFromContextArgs[] = {&cu.getForce().getDevicePointer(), &forces->getDevicePointer(), &cu.getVelm().getDevicePointer(),
&velocities->getDevicePointer(), &cu.getPosq().getDevicePointer(), &positions->getDevicePointer(), &cu.getAtomIndexArray().getDevicePointer(), &i};
......
plugins/rpmd/platforms/cuda/src/kernels/rpmd.cu
View file @ d56c21b0
......@@ -217,7 +217,7 @@ extern "C" __global__ void copyDataFromContext(long long* srcForce, long long* d
}
/**
* Update atom positions so all copies are offset by the same number of periodic box widths.
* Atom positions in one copy have been modified. Apply the same offsets to all the other copies.
*/
extern "C" __global__ void applyCellTranslations(mixed4* posq, real4* movedPos, int* order, int movedCopy) {
for (int particle = blockIdx.x*blockDim.x+threadIdx.x; particle < NUM_ATOMS; particle += blockDim.x*gridDim.x) {
......
plugins/rpmd/platforms/cuda/tests/TestCudaRpmd.cpp
View file @ d56c21b0
......@@ -38,6 +38,7 @@
#include "openmm/Context.h"
#include "openmm/CustomNonbondedForce.h"
#include "openmm/HarmonicBondForce.h"
#include "openmm/MonteCarloBarostat.h"
#include "openmm/NonbondedForce.h"
#include "openmm/Platform.h"
#include "openmm/System.h"
......@@ -496,6 +497,80 @@ void testWithoutThermostat() {
}
}
void testWithBarostat() {
const int gridSize = 3;
const int numMolecules = gridSize*gridSize*gridSize;
const int numParticles = numMolecules*2;
const int numCopies = 10;
const double spacing = 2.0;
const double cutoff = 3.0;
const double boxSize = spacing*(gridSize+1);
const double temperature = 300.0;
System system;
system.setDefaultPeriodicBoxVectors(Vec3(boxSize, 0, 0), Vec3(0, boxSize, 0), Vec3(0, 0, boxSize));
HarmonicBondForce* bonds = new HarmonicBondForce();
system.addForce(bonds);
NonbondedForce* nonbonded = new NonbondedForce();
nonbonded->setCutoffDistance(cutoff);
nonbonded->setNonbondedMethod(NonbondedForce::PME);
nonbonded->setForceGroup(1);
nonbonded->setReciprocalSpaceForceGroup(2);
system.addForce(nonbonded);
system.addForce(new MonteCarloBarostat(0.5, temperature));
// Create a cloud of molecules.
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numMolecules; i++) {
system.addParticle(1.0);
system.addParticle(1.0);
nonbonded->addParticle(-0.2, 0.2, 0.2);
nonbonded->addParticle(0.2, 0.2, 0.2);
nonbonded->addException(2*i, 2*i+1, 0, 1, 0);
bonds->addBond(2*i, 2*i+1, 1.0, 10000.0);
}
RPMDIntegrator integ(numCopies, temperature, 50.0, 0.001);
Platform& platform = Platform::getPlatformByName("CUDA");
Context context(system, integ, platform);
for (int copy = 0; copy < numCopies; copy++) {
for (int i = 0; i < gridSize; i++)
for (int j = 0; j < gridSize; j++)
for (int k = 0; k < gridSize; k++) {
Vec3 pos = Vec3(spacing*(i+0.02*genrand_real2(sfmt)), spacing*(j+0.02*genrand_real2(sfmt)), spacing*(k+0.02*genrand_real2(sfmt)));
int index = k+gridSize*(j+gridSize*i);
positions[2*index] = pos;
positions[2*index+1] = Vec3(pos[0]+1.0, pos[1], pos[2]);
}
integ.setPositions(copy, positions);
}
// Check the temperature.
const int numSteps = 500;
integ.step(100);
vector<double> ke(numCopies, 0.0);
for (int i = 0; i < numSteps; i++) {
integ.step(1);
vector<State> state(numCopies);
for (int j = 0; j < numCopies; j++)
state[j] = integ.getState(j, State::Velocities, true);
for (int j = 0; j < numParticles; j++) {
for (int k = 0; k < numCopies; k++) {
Vec3 v = state[k].getVelocities()[j];
ke[k] += 0.5*system.getParticleMass(j)*v.dot(v);
}
}
}
double meanKE = 0.0;
for (int i = 0; i < numCopies; i++)
meanKE += ke[i];
meanKE /= numSteps*numCopies;
double expectedKE = 0.5*numCopies*numParticles*3*BOLTZ*temperature;
ASSERT_USUALLY_EQUAL_TOL(expectedKE, meanKE, 1e-2);
}
int main(int argc, char* argv[]) {
try {
registerRPMDCudaKernelFactories();
......@@ -507,6 +582,7 @@ int main(int argc, char* argv[]) {
testVirtualSites();
testContractions();
testWithoutThermostat();
testWithBarostat();
}
catch(const std::exception& e) {
std::cout << "exception: " << e.what() << std::endl;
......
plugins/rpmd/platforms/opencl/src/OpenCLRpmdKernels.cpp
View file @ d56c21b0
......@@ -268,7 +268,18 @@ void OpenCLIntegrateRPMDStepKernel::computeForces(ContextImpl& context) {
copyToContextKernel.setArg<cl_int>(5, i);
cl.executeKernel(copyToContextKernel, cl.getNumAtoms());
context.computeVirtualSites();
Vec3 initialBox[3];
context.getPeriodicBoxVectors(initialBox[0], initialBox[1], initialBox[2]);
context.updateContextState();
Vec3 finalBox[3];
context.getPeriodicBoxVectors(finalBox[0], finalBox[1], finalBox[2]);
if (initialBox[0] != finalBox[0] || initialBox[1] != finalBox[1] || initialBox[2] != finalBox[2]) {
// A barostat was applied during updateContextState(). Adjust the particle positions in all the
// other copies to match this one.
translateKernel.setArg<cl_int>(3, i);
cl.executeKernel(translateKernel, cl.getNumAtoms());
}
context.calcForcesAndEnergy(true, false, groupsNotContracted);
copyFromContextKernel.setArg<cl_int>(7, i);
cl.executeKernel(copyFromContextKernel, cl.getNumAtoms());
......
plugins/rpmd/platforms/opencl/src/kernels/rpmd.cl
View file @ d56c21b0
......@@ -204,7 +204,7 @@ __kernel void copyDataFromContext(__global real4* srcForce, __global real4* dstF
}
/**
* Update atom positions so all copies are offset by the same number of periodic box widths.
* Atom positions in one copy have been modified. Apply the same offsets to all the other copies.
*/
__kernel void applyCellTranslations(__global mixed4* posq, __global real4* movedPos, __global int* order, int movedCopy) {
for (int particle = get_global_id(0); particle < NUM_ATOMS; particle += get_global_size(0)) {
......
plugins/rpmd/platforms/opencl/tests/TestOpenCLRpmd.cpp
View file @ d56c21b0
......@@ -38,6 +38,7 @@
#include "openmm/Context.h"
#include "openmm/CustomNonbondedForce.h"
#include "openmm/HarmonicBondForce.h"
#include "openmm/MonteCarloBarostat.h"
#include "openmm/NonbondedForce.h"
#include "openmm/Platform.h"
#include "openmm/System.h"
......@@ -497,6 +498,80 @@ void testWithoutThermostat() {
}
}
void testWithBarostat() {
const int gridSize = 3;
const int numMolecules = gridSize*gridSize*gridSize;
const int numParticles = numMolecules*2;
const int numCopies = 10;
const double spacing = 2.0;
const double cutoff = 3.0;
const double boxSize = spacing*(gridSize+1);
const double temperature = 300.0;
System system;
system.setDefaultPeriodicBoxVectors(Vec3(boxSize, 0, 0), Vec3(0, boxSize, 0), Vec3(0, 0, boxSize));
HarmonicBondForce* bonds = new HarmonicBondForce();
system.addForce(bonds);
NonbondedForce* nonbonded = new NonbondedForce();
nonbonded->setCutoffDistance(cutoff);
nonbonded->setNonbondedMethod(NonbondedForce::PME);
nonbonded->setForceGroup(1);
nonbonded->setReciprocalSpaceForceGroup(2);
system.addForce(nonbonded);
system.addForce(new MonteCarloBarostat(0.5, temperature));
// Create a cloud of molecules.
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numMolecules; i++) {
system.addParticle(1.0);
system.addParticle(1.0);
nonbonded->addParticle(-0.2, 0.2, 0.2);
nonbonded->addParticle(0.2, 0.2, 0.2);
nonbonded->addException(2*i, 2*i+1, 0, 1, 0);
bonds->addBond(2*i, 2*i+1, 1.0, 10000.0);
}
RPMDIntegrator integ(numCopies, temperature, 50.0, 0.001);
Platform& platform = Platform::getPlatformByName("OpenCL");
Context context(system, integ, platform);
for (int copy = 0; copy < numCopies; copy++) {
for (int i = 0; i < gridSize; i++)
for (int j = 0; j < gridSize; j++)
for (int k = 0; k < gridSize; k++) {
Vec3 pos = Vec3(spacing*(i+0.02*genrand_real2(sfmt)), spacing*(j+0.02*genrand_real2(sfmt)), spacing*(k+0.02*genrand_real2(sfmt)));
int index = k+gridSize*(j+gridSize*i);
positions[2*index] = pos;
positions[2*index+1] = Vec3(pos[0]+1.0, pos[1], pos[2]);
}
integ.setPositions(copy, positions);
}
// Check the temperature.
const int numSteps = 500;
integ.step(100);
vector<double> ke(numCopies, 0.0);
for (int i = 0; i < numSteps; i++) {
integ.step(1);
vector<State> state(numCopies);
for (int j = 0; j < numCopies; j++)
state[j] = integ.getState(j, State::Velocities, true);
for (int j = 0; j < numParticles; j++) {
for (int k = 0; k < numCopies; k++) {
Vec3 v = state[k].getVelocities()[j];
ke[k] += 0.5*system.getParticleMass(j)*v.dot(v);
}
}
}
double meanKE = 0.0;
for (int i = 0; i < numCopies; i++)
meanKE += ke[i];
meanKE /= numSteps*numCopies;
double expectedKE = 0.5*numCopies*numParticles*3*BOLTZ*temperature;
ASSERT_USUALLY_EQUAL_TOL(expectedKE, meanKE, 1e-2);
}
int main(int argc, char* argv[]) {
try {
registerRPMDOpenCLKernelFactories();
......@@ -508,6 +583,7 @@ int main(int argc, char* argv[]) {
testVirtualSites();
testContractions();
testWithoutThermostat();
testWithBarostat();
}
catch(const std::exception& e) {
std::cout << "exception: " << e.what() << std::endl;
......
plugins/rpmd/platforms/reference/CMakeLists.txt 0 → 100644
View file @ d56c21b0
#---------------------------------------------------
# OpenMM Reference RPMD Integrator
#
# Creates OpenMM library, base name=OpenMMRPMDReference.
# Default libraries are shared & optimized. Variants
# are created for debug (_d).
#
# Windows:
# OpenMMRPMDReference[_d].dll
# OpenMMRPMDReference[_d].lib
# Unix:
# libOpenMMRPMDReference[_d].so
#----------------------------------------------------
# The source is organized into subdirectories, but we handle them all from
# this CMakeLists file rather than letting CMake visit them as SUBDIRS.
SET(OPENMM_SOURCE_SUBDIRS .)
# Collect up information about the version of the OpenMM library we're building
# and make it available to the code so it can be built into the binaries.
SET(OPENMMRPMDREFERENCE_LIBRARY_NAME OpenMMRPMDReference)
SET(SHARED_TARGET ${OPENMMRPMDREFERENCE_LIBRARY_NAME})
# Ensure that debug libraries have "_d" appended to their names.
# CMake gets this right on Windows automatically with this definition.
IF (${CMAKE_GENERATOR} MATCHES "Visual Studio")
SET(CMAKE_DEBUG_POSTFIX "_d" CACHE INTERNAL "" FORCE)
ENDIF (${CMAKE_GENERATOR} MATCHES "Visual Studio")
# But on Unix or Cygwin we have to add the suffix manually
IF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
SET(SHARED_TARGET ${SHARED_TARGET}_d)
ENDIF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
# These are all the places to search for header files which are
# to be part of the API.
SET(API_INCLUDE_DIRS) # start empty
FOREACH(subdir ${OPENMM_SOURCE_SUBDIRS})
# append
SET(API_INCLUDE_DIRS ${API_INCLUDE_DIRS}
${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include
${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include/internal)
ENDFOREACH(subdir)
# We'll need both *relative* path names, starting with their API_INCLUDE_DIRS,
# and absolute pathnames.
SET(API_REL_INCLUDE_FILES) # start these out empty
SET(API_ABS_INCLUDE_FILES)
FOREACH(dir ${API_INCLUDE_DIRS})
FILE(GLOB fullpaths ${dir}/*.h) # returns full pathnames
SET(API_ABS_INCLUDE_FILES ${API_ABS_INCLUDE_FILES} ${fullpaths})
FOREACH(pathname ${fullpaths})
GET_FILENAME_COMPONENT(filename ${pathname} NAME)
SET(API_REL_INCLUDE_FILES ${API_REL_INCLUDE_FILES} ${dir}/${filename})
ENDFOREACH(pathname)
ENDFOREACH(dir)
# collect up source files
SET(SOURCE_FILES) # empty
SET(SOURCE_INCLUDE_FILES)
FOREACH(subdir ${OPENMM_SOURCE_SUBDIRS})
FILE(GLOB_RECURSE src_files ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/src/*.cpp ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/src/*.c)
FILE(GLOB incl_files ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/src/*.h)
SET(SOURCE_FILES ${SOURCE_FILES} ${src_files}) #append
SET(SOURCE_INCLUDE_FILES ${SOURCE_INCLUDE_FILES} ${incl_files})
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/${subdir}/include)
ENDFOREACH(subdir)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_CURRENT_SOURCE_DIR}/src)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_SOURCE_DIR}/platforms/reference/include)
INCLUDE_DIRECTORIES(BEFORE ${CMAKE_SOURCE_DIR}/platforms/reference/src)
# Create the library
INCLUDE_DIRECTORIES(${REFERENCE_INCLUDE_DIR})
ADD_LIBRARY(${SHARED_TARGET} SHARED ${SOURCE_FILES} ${SOURCE_INCLUDE_FILES} ${API_ABS_INCLUDE_FILES})
IF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
SET(MAIN_OPENMM_LIB ${OPENMM_LIBRARY_NAME}_d)
ELSE (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
SET(MAIN_OPENMM_LIB ${OPENMM_LIBRARY_NAME})
ENDIF (UNIX AND CMAKE_BUILD_TYPE MATCHES Debug)
TARGET_LINK_LIBRARIES(${SHARED_TARGET} ${MAIN_OPENMM_LIB})
TARGET_LINK_LIBRARIES(${SHARED_TARGET} debug ${SHARED_RPMD_TARGET} optimized ${SHARED_RPMD_TARGET})
SET_TARGET_PROPERTIES(${SHARED_TARGET} PROPERTIES LINK_FLAGS "${EXTRA_COMPILE_FLAGS}" COMPILE_FLAGS "${EXTRA_COMPILE_FLAGS} -msse2 -DOPENMM_BUILDING_SHARED_LIBRARY")
INSTALL(TARGETS ${SHARED_TARGET} DESTINATION ${CMAKE_INSTALL_PREFIX}/lib/plugins)
SUBDIRS (tests)
plugins/rpmd/platforms/reference/src/ReferenceRpmdKernelFactory.cpp
View file @ d56c21b0
......@@ -32,22 +32,21 @@
using namespace OpenMM;
#if defined(WIN32)
#include <windows.h>
extern "C" void initRpmdReferenceKernels();
BOOL WINAPI DllMain(HANDLE hModule, DWORD ul_reason_for_call, LPVOID lpReserved) {
if (ul_reason_for_call == DLL_PROCESS_ATTACH)
initRpmdReferenceKernels();
return TRUE;
}
#else
extern "C" void __attribute__((constructor)) initRpmdReferenceKernels();
#endif
extern "C" OPENMM_EXPORT void registerPlatforms() {
}
extern "C" void initRpmdReferenceKernels() {
Platform& platform = Platform::getPlatformByName("Reference");
extern "C" OPENMM_EXPORT void registerKernelFactories() {
for (int i = 0; i < Platform::getNumPlatforms(); i++) {
Platform& platform = Platform::getPlatform(i);
if (dynamic_cast<ReferencePlatform*>(&platform) != NULL) {
ReferenceRpmdKernelFactory* factory = new ReferenceRpmdKernelFactory();
platform.registerKernelFactory(IntegrateRPMDStepKernel::Name(), factory);
}
}
}
extern "C" OPENMM_EXPORT void registerRpmdReferenceKernelFactories() {
registerKernelFactories();
}
KernelImpl* ReferenceRpmdKernelFactory::createKernelImpl(std::string name, const Platform& platform, ContextImpl& context) const {
......
plugins/rpmd/platforms/reference/src/ReferenceRpmdKernels.cpp
View file @ d56c21b0
......@@ -275,7 +275,22 @@ void ReferenceIntegrateRPMDStepKernel::computeForces(ContextImpl& context, const
pos = positions[i];
vel = velocities[i];
context.computeVirtualSites();
Vec3 initialBox[3];
context.getPeriodicBoxVectors(initialBox[0], initialBox[1], initialBox[2]);
context.updateContextState();
Vec3 finalBox[3];
context.getPeriodicBoxVectors(finalBox[0], finalBox[1], finalBox[2]);
if (initialBox[0] != finalBox[0] || initialBox[1] != finalBox[1] || initialBox[2] != finalBox[2]) {
// A barostat was applied during updateContextState(). Adjust the particle positions in all the
// other copies to match this one.
for (int j = 0; j < numParticles; j++) {
Vec3 delta = pos[j]-positions[i][j];
for (int k = 0; k < totalCopies; k++)
if (k != i)
positions[k][j] += delta;
}
}
positions[i] = pos;
velocities[i] = vel;
context.calcForcesAndEnergy(true, false, groupsNotContracted);
......
plugins/rpmd/platforms/reference/tests/CMakeLists.txt
View file @ d56c21b0
......@@ -2,11 +2,9 @@
# Testing
#
ENABLE_TESTING()
#INCLUDE_DIRECTORIES(${CUDA_INCLUDE})
INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/reference/include)
INCLUDE_DIRECTORIES(${OPENMM_DIR}/openmmapi/include/openmm)
INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/reference/src)
#INCLUDE_DIRECTORIES(${OPENMM_DIR}/platforms/reference/src/kernels)
SET(SHARED_OPENMM_RPMD_TARGET OpenMMRPMD)
......
plugins/rpmd/platforms/reference/tests/TestReferenceRpmd.cpp
View file @ d56c21b0
......@@ -37,6 +37,7 @@
#include "openmm/CMMotionRemover.h"
#include "openmm/Context.h"
#include "openmm/HarmonicBondForce.h"
#include "openmm/MonteCarloBarostat.h"
#include "openmm/NonbondedForce.h"
#include "openmm/Platform.h"
#include "openmm/System.h"
......@@ -50,6 +51,8 @@
using namespace OpenMM;
using namespace std;
extern "C" OPENMM_EXPORT void registerRpmdReferenceKernelFactories();
void testFreeParticles() {
const int numParticles = 100;
const int numCopies = 30;
......@@ -378,13 +381,89 @@ void testWithoutThermostat() {
}
}
void testWithBarostat() {
const int gridSize = 3;
const int numMolecules = gridSize*gridSize*gridSize;
const int numParticles = numMolecules*2;
const int numCopies = 5;
const double spacing = 2.0;
const double cutoff = 3.0;
const double boxSize = spacing*(gridSize+1);
const double temperature = 300.0;
System system;
system.setDefaultPeriodicBoxVectors(Vec3(boxSize, 0, 0), Vec3(0, boxSize, 0), Vec3(0, 0, boxSize));
HarmonicBondForce* bonds = new HarmonicBondForce();
system.addForce(bonds);
NonbondedForce* nonbonded = new NonbondedForce();
nonbonded->setCutoffDistance(cutoff);
nonbonded->setNonbondedMethod(NonbondedForce::PME);
nonbonded->setForceGroup(1);
nonbonded->setReciprocalSpaceForceGroup(2);
system.addForce(nonbonded);
system.addForce(new MonteCarloBarostat(0.5, temperature));
// Create a cloud of molecules.
OpenMM_SFMT::SFMT sfmt;
init_gen_rand(0, sfmt);
vector<Vec3> positions(numParticles);
for (int i = 0; i < numMolecules; i++) {
system.addParticle(1.0);
system.addParticle(1.0);
nonbonded->addParticle(-0.2, 0.2, 0.2);
nonbonded->addParticle(0.2, 0.2, 0.2);
nonbonded->addException(2*i, 2*i+1, 0, 1, 0);
bonds->addBond(2*i, 2*i+1, 1.0, 10000.0);
}
RPMDIntegrator integ(numCopies, temperature, 50.0, 0.001);
Platform& platform = Platform::getPlatformByName("Reference");
Context context(system, integ, platform);
for (int copy = 0; copy < numCopies; copy++) {
for (int i = 0; i < gridSize; i++)
for (int j = 0; j < gridSize; j++)
for (int k = 0; k < gridSize; k++) {
Vec3 pos = Vec3(spacing*(i+0.02*genrand_real2(sfmt)), spacing*(j+0.02*genrand_real2(sfmt)), spacing*(k+0.02*genrand_real2(sfmt)));
int index = k+gridSize*(j+gridSize*i);
positions[2*index] = pos;
positions[2*index+1] = Vec3(pos[0]+1.0, pos[1], pos[2]);
}
integ.setPositions(copy, positions);
}
// Check the temperature.
const int numSteps = 500;
integ.step(100);
vector<double> ke(numCopies, 0.0);
for (int i = 0; i < numSteps; i++) {
integ.step(1);
vector<State> state(numCopies);
for (int j = 0; j < numCopies; j++)
state[j] = integ.getState(j, State::Velocities, true);
for (int j = 0; j < numParticles; j++) {
for (int k = 0; k < numCopies; k++) {
Vec3 v = state[k].getVelocities()[j];
ke[k] += 0.5*system.getParticleMass(j)*v.dot(v);
}
}
}
double meanKE = 0.0;
for (int i = 0; i < numCopies; i++)
meanKE += ke[i];
meanKE /= numSteps*numCopies;
double expectedKE = 0.5*numCopies*numParticles*3*BOLTZ*temperature;
ASSERT_USUALLY_EQUAL_TOL(expectedKE, meanKE, 1e-2);
}
int main() {
try {
registerRpmdReferenceKernelFactories();
testFreeParticles();
testCMMotionRemoval();
testVirtualSites();
testContractions();
testWithoutThermostat();
testWithBarostat();
}
catch(const std::exception& e) {
std::cout << "exception: " << e.what() << std::endl;
......
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