Skip to content

GitLab

Commit 2c559596 authored by peastman's avatar peastman Committed by GitHub
Browse files

Merge pull request #1740 from peastman/updateparams 

Optimizations to updateParametersInContext()
parents dd712841 2a1e9acf
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 381 additions and 193 deletions
platforms/cuda/include/CudaContext.h
View file @ 2c559596
...@@ -9,7 +9,7 @@ ...@@ -9,7 +9,7 @@
* Biological Structures at Stanford, funded under the NIH Roadmap for * * Biological Structures at Stanford, funded under the NIH Roadmap for *
* Medical Research, grant U54 GM072970. See https://simtk.org. * * Medical Research, grant U54 GM072970. See https://simtk.org. *
* * * *
* Portions copyright (c) 2009-2016 Stanford University and the Authors. * * Portions copyright (c) 2009-2017 Stanford University and the Authors. *
* Authors: Peter Eastman * * Authors: Peter Eastman *
* Contributors: * * Contributors: *
* * * *
...@@ -494,6 +494,10 @@ public: ...@@ -494,6 +494,10 @@ public:
CudaNonbondedUtilities& getNonbondedUtilities() { CudaNonbondedUtilities& getNonbondedUtilities() {
return *nonbonded; return *nonbonded;
} }
/**
* Set the particle charges. These are packed into the fourth element of the posq array.
*/
void setCharges(const std::vector<double>& charges);
/** /**
* Get the thread used by this context for executing parallel computations. * Get the thread used by this context for executing parallel computations.
*/ */
...@@ -577,6 +581,12 @@ public: ...@@ -577,6 +581,12 @@ public:
* and order to be revalidated. * and order to be revalidated.
*/ */
void invalidateMolecules(); void invalidateMolecules();
/**
* Mark that the current molecule definitions from one particular force (and hence the atom order)
* may be invalid. This should be called whenever force field parameters change. It will cause the
* definitions and order to be revalidated.
*/
bool invalidateMolecules(CudaForceInfo* force);
private: private:
/** /**
* Compute a sorted list of device indices in decreasing order of desirability * Compute a sorted list of device indices in decreasing order of desirability
...@@ -626,6 +636,7 @@ private: ...@@ -626,6 +636,7 @@ private:
CUfunction clearFourBuffersKernel; CUfunction clearFourBuffersKernel;
CUfunction clearFiveBuffersKernel; CUfunction clearFiveBuffersKernel;
CUfunction clearSixBuffersKernel; CUfunction clearSixBuffersKernel;
CUfunction setChargesKernel;
std::vector<CudaForceInfo*> forces; std::vector<CudaForceInfo*> forces;
std::vector<Molecule> molecules; std::vector<Molecule> molecules;
std::vector<MoleculeGroup> moleculeGroups; std::vector<MoleculeGroup> moleculeGroups;
...@@ -638,6 +649,7 @@ private: ...@@ -638,6 +649,7 @@ private:
CudaArray* energyBuffer; CudaArray* energyBuffer;
CudaArray* energyParamDerivBuffer; CudaArray* energyParamDerivBuffer;
CudaArray* atomIndexDevice; CudaArray* atomIndexDevice;
CudaArray* chargeBuffer;
std::vector<std::string> energyParamDerivNames; std::vector<std::string> energyParamDerivNames;
std::map<std::string, double> energyParamDerivWorkspace; std::map<std::string, double> energyParamDerivWorkspace;
std::vector<int> atomIndex; std::vector<int> atomIndex;
......
platforms/cuda/include/CudaKernels.h
View file @ 2c559596
...@@ -9,7 +9,7 @@ ...@@ -9,7 +9,7 @@
* Biological Structures at Stanford, funded under the NIH Roadmap for * * Biological Structures at Stanford, funded under the NIH Roadmap for *
* Medical Research, grant U54 GM072970. See https://simtk.org. * * Medical Research, grant U54 GM072970. See https://simtk.org. *
* * * *
* Portions copyright (c) 2008-2016 Stanford University and the Authors. * * Portions copyright (c) 2008-2017 Stanford University and the Authors. *
* Authors: Peter Eastman * * Authors: Peter Eastman *
* Contributors: * * Contributors: *
* * * *
...@@ -292,9 +292,11 @@ public: ...@@ -292,9 +292,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force); void copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force);
private: private:
class ForceInfo;
int numBonds; int numBonds;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaArray* params; CudaArray* params;
}; };
...@@ -332,9 +334,11 @@ public: ...@@ -332,9 +334,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomBondForce& force); void copyParametersToContext(ContextImpl& context, const CustomBondForce& force);
private: private:
class ForceInfo;
int numBonds; int numBonds;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
...@@ -375,9 +379,11 @@ public: ...@@ -375,9 +379,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force); void copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force);
private: private:
class ForceInfo;
int numAngles; int numAngles;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaArray* params; CudaArray* params;
}; };
...@@ -415,9 +421,11 @@ public: ...@@ -415,9 +421,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomAngleForce& force); void copyParametersToContext(ContextImpl& context, const CustomAngleForce& force);
private: private:
class ForceInfo;
int numAngles; int numAngles;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
...@@ -458,9 +466,11 @@ public: ...@@ -458,9 +466,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force); void copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaArray* params; CudaArray* params;
}; };
...@@ -498,9 +508,11 @@ public: ...@@ -498,9 +508,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const RBTorsionForce& force); void copyParametersToContext(ContextImpl& context, const RBTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaArray* params1; CudaArray* params1;
CudaArray* params2; CudaArray* params2;
...@@ -539,9 +551,11 @@ public: ...@@ -539,9 +551,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force); void copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
std::vector<int2> mapPositionsVec; std::vector<int2> mapPositionsVec;
CudaArray* coefficients; CudaArray* coefficients;
...@@ -582,9 +596,11 @@ public: ...@@ -582,9 +596,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force); void copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
...@@ -647,12 +663,14 @@ private: ...@@ -647,12 +663,14 @@ private:
const char* getMaxValue() const {return "make_int2(2147483647, 2147483647)";} const char* getMaxValue() const {return "make_int2(2147483647, 2147483647)";}
const char* getSortKey() const {return "value.y";} const char* getSortKey() const {return "value.y";}
}; };
class ForceInfo;
class PmeIO; class PmeIO;
class PmePreComputation; class PmePreComputation;
class PmePostComputation; class PmePostComputation;
class SyncStreamPreComputation; class SyncStreamPreComputation;
class SyncStreamPostComputation; class SyncStreamPostComputation;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
bool hasInitializedFFT; bool hasInitializedFFT;
CudaArray* sigmaEpsilon; CudaArray* sigmaEpsilon;
CudaArray* exceptionParams; CudaArray* exceptionParams;
...@@ -724,8 +742,10 @@ public: ...@@ -724,8 +742,10 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force); void copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force);
private: private:
class ForceInfo;
void initInteractionGroups(const CustomNonbondedForce& force, const std::string& interactionSource, const std::vector<std::string>& tableTypes); void initInteractionGroups(const CustomNonbondedForce& force, const std::string& interactionSource, const std::vector<std::string>& tableTypes);
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
CudaArray* interactionGroupData; CudaArray* interactionGroupData;
...@@ -775,10 +795,12 @@ public: ...@@ -775,10 +795,12 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const GBSAOBCForce& force); void copyParametersToContext(ContextImpl& context, const GBSAOBCForce& force);
private: private:
class ForceInfo;
double prefactor, surfaceAreaFactor, cutoff; double prefactor, surfaceAreaFactor, cutoff;
bool hasCreatedKernels; bool hasCreatedKernels;
int maxTiles; int maxTiles;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
CudaArray* params; CudaArray* params;
CudaArray* bornSum; CudaArray* bornSum;
CudaArray* bornRadii; CudaArray* bornRadii;
...@@ -825,10 +847,12 @@ public: ...@@ -825,10 +847,12 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomGBForce& force); void copyParametersToContext(ContextImpl& context, const CustomGBForce& force);
private: private:
class ForceInfo;
double cutoff; double cutoff;
bool hasInitializedKernels, needParameterGradient, needEnergyParamDerivs; bool hasInitializedKernels, needParameterGradient, needEnergyParamDerivs;
int maxTiles, numComputedValues; int maxTiles, numComputedValues;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
CudaParameterSet* params; CudaParameterSet* params;
CudaParameterSet* computedValues; CudaParameterSet* computedValues;
CudaParameterSet* energyDerivs; CudaParameterSet* energyDerivs;
...@@ -882,9 +906,11 @@ public: ...@@ -882,9 +906,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomExternalForce& force); void copyParametersToContext(ContextImpl& context, const CustomExternalForce& force);
private: private:
class ForceInfo;
int numParticles; int numParticles;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
const System& system; const System& system;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
...@@ -926,9 +952,11 @@ public: ...@@ -926,9 +952,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomHbondForce& force); void copyParametersToContext(ContextImpl& context, const CustomHbondForce& force);
private: private:
class ForceInfo;
int numDonors, numAcceptors; int numDonors, numAcceptors;
bool hasInitializedKernel; bool hasInitializedKernel;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
CudaParameterSet* donorParams; CudaParameterSet* donorParams;
CudaParameterSet* acceptorParams; CudaParameterSet* acceptorParams;
CudaArray* globals; CudaArray* globals;
...@@ -978,9 +1006,11 @@ public: ...@@ -978,9 +1006,11 @@ public:
void copyParametersToContext(ContextImpl& context, const CustomCentroidBondForce& force); void copyParametersToContext(ContextImpl& context, const CustomCentroidBondForce& force);
private: private:
class ForceInfo;
int numGroups, numBonds; int numGroups, numBonds;
bool needEnergyParamDerivs; bool needEnergyParamDerivs;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
CudaArray* groupParticles; CudaArray* groupParticles;
...@@ -1031,8 +1061,10 @@ public: ...@@ -1031,8 +1061,10 @@ public:
void copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force); void copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force);
private: private:
class ForceInfo;
int numBonds; int numBonds;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
CudaParameterSet* params; CudaParameterSet* params;
CudaArray* globals; CudaArray* globals;
std::vector<std::string> globalParamNames; std::vector<std::string> globalParamNames;
...@@ -1077,7 +1109,9 @@ public: ...@@ -1077,7 +1109,9 @@ public:
void copyParametersToContext(ContextImpl& context, const CustomManyParticleForce& force); void copyParametersToContext(ContextImpl& context, const CustomManyParticleForce& force);
private: private:
class ForceInfo;
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
bool hasInitializedKernel; bool hasInitializedKernel;
NonbondedMethod nonbondedMethod; NonbondedMethod nonbondedMethod;
int maxNeighborPairs, forceWorkgroupSize, findNeighborsWorkgroupSize; int maxNeighborPairs, forceWorkgroupSize, findNeighborsWorkgroupSize;
...@@ -1139,9 +1173,11 @@ public: ...@@ -1139,9 +1173,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const GayBerneForce& force); void copyParametersToContext(ContextImpl& context, const GayBerneForce& force);
private: private:
class ForceInfo;
class ReorderListener; class ReorderListener;
void sortAtoms(); void sortAtoms();
CudaContext& cu; CudaContext& cu;
ForceInfo* info;
bool hasInitializedKernels; bool hasInitializedKernels;
int numRealParticles, numExceptions, maxNeighborBlocks; int numRealParticles, numExceptions, maxNeighborBlocks;
GayBerneForce::NonbondedMethod nonbondedMethod; GayBerneForce::NonbondedMethod nonbondedMethod;
......
platforms/cuda/src/CudaContext.cpp
View file @ 2c559596
...@@ -6,7 +6,7 @@ ...@@ -6,7 +6,7 @@
* Biological Structures at Stanford, funded under the NIH Roadmap for * * Biological Structures at Stanford, funded under the NIH Roadmap for *
* Medical Research, grant U54 GM072970. See https://simtk.org. * * Medical Research, grant U54 GM072970. See https://simtk.org. *
* * * *
* Portions copyright (c) 2009-2016 Stanford University and the Authors. * * Portions copyright (c) 2009-2017 Stanford University and the Authors. *
* Authors: Peter Eastman * * Authors: Peter Eastman *
* Contributors: * * Contributors: *
* * * *
...@@ -108,7 +108,8 @@ static int executeInWindows(const string &command) { ...@@ -108,7 +108,8 @@ static int executeInWindows(const string &command) {
CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlockingSync, const string& precision, const string& compiler, CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlockingSync, const string& precision, const string& compiler,
const string& tempDir, const std::string& hostCompiler, CudaPlatform::PlatformData& platformData) : system(system), currentStream(0), const string& tempDir, const std::string& hostCompiler, CudaPlatform::PlatformData& platformData) : system(system), currentStream(0),
time(0.0), platformData(platformData), stepCount(0), computeForceCount(0), stepsSinceReorder(99999), contextIsValid(false), atomsWereReordered(false), hasCompilerKernel(false), isNvccAvailable(false), time(0.0), platformData(platformData), stepCount(0), computeForceCount(0), stepsSinceReorder(99999), contextIsValid(false), atomsWereReordered(false), hasCompilerKernel(false), isNvccAvailable(false),
pinnedBuffer(NULL), posq(NULL), posqCorrection(NULL), velm(NULL), force(NULL), energyBuffer(NULL), energyParamDerivBuffer(NULL), atomIndexDevice(NULL), integration(NULL), expression(NULL), bonded(NULL), nonbonded(NULL), thread(NULL) { pinnedBuffer(NULL), posq(NULL), posqCorrection(NULL), velm(NULL), force(NULL), energyBuffer(NULL), energyParamDerivBuffer(NULL), atomIndexDevice(NULL), chargeBuffer(NULL),
integration(NULL), expression(NULL), bonded(NULL), nonbonded(NULL), thread(NULL) {
// Determine what compiler to use. // Determine what compiler to use.
this->compiler = "\""+compiler+"\""; this->compiler = "\""+compiler+"\"";
...@@ -291,6 +292,7 @@ CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlocking ...@@ -291,6 +292,7 @@ CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlocking
clearFourBuffersKernel = getKernel(utilities, "clearFourBuffers"); clearFourBuffersKernel = getKernel(utilities, "clearFourBuffers");
clearFiveBuffersKernel = getKernel(utilities, "clearFiveBuffers"); clearFiveBuffersKernel = getKernel(utilities, "clearFiveBuffers");
clearSixBuffersKernel = getKernel(utilities, "clearSixBuffers"); clearSixBuffersKernel = getKernel(utilities, "clearSixBuffers");
setChargesKernel = getKernel(utilities, "setCharges");
// Set defines based on the requested precision. // Set defines based on the requested precision.
...@@ -407,6 +409,8 @@ CudaContext::~CudaContext() { ...@@ -407,6 +409,8 @@ CudaContext::~CudaContext() {
delete energyParamDerivBuffer; delete energyParamDerivBuffer;
if (atomIndexDevice != NULL) if (atomIndexDevice != NULL)
delete atomIndexDevice; delete atomIndexDevice;
if (chargeBuffer != NULL)
delete chargeBuffer;
if (integration != NULL) if (integration != NULL)
delete integration; delete integration;
if (expression != NULL) if (expression != NULL)
...@@ -860,6 +864,25 @@ void CudaContext::clearAutoclearBuffers() { ...@@ -860,6 +864,25 @@ void CudaContext::clearAutoclearBuffers() {
} }
} }
void CudaContext::setCharges(const vector<double>& charges) {
if (chargeBuffer == NULL)
chargeBuffer = new CudaArray(*this, numAtoms, useDoublePrecision ? sizeof(double) : sizeof(float), "chargeBuffer");
if (getUseDoublePrecision()) {
double* c = (double*) getPinnedBuffer();
for (int i = 0; i < charges.size(); i++)
c[i] = charges[i];
chargeBuffer->upload(c);
}
else {
float* c = (float*) getPinnedBuffer();
for (int i = 0; i < charges.size(); i++)
c[i] = (float) charges[i];
chargeBuffer->upload(c);
}
void* args[] = {&chargeBuffer->getDevicePointer(), &posq->getDevicePointer(), &atomIndexDevice->getDevicePointer(), &numAtoms};
executeKernel(setChargesKernel, args, numAtoms);
}
/** /**
* This class ensures that atom reordering doesn't break virtual sites. * This class ensures that atom reordering doesn't break virtual sites.
*/ */
...@@ -1058,9 +1081,19 @@ void CudaContext::findMoleculeGroups() { ...@@ -1058,9 +1081,19 @@ void CudaContext::findMoleculeGroups() {
} }
void CudaContext::invalidateMolecules() { void CudaContext::invalidateMolecules() {
for (int i = 0; i < forces.size(); i++)
if (invalidateMolecules(forces[i]))
return;
}
bool CudaContext::invalidateMolecules(CudaForceInfo* force) {
if (numAtoms == 0 || nonbonded == NULL || !nonbonded->getUseCutoff()) if (numAtoms == 0 || nonbonded == NULL || !nonbonded->getUseCutoff())
return; return false;
bool valid = true; bool valid = true;
int forceIndex = -1;
for (int i = 0; i < forces.size(); i++)
if (forces[i] == force)
forceIndex = i;
for (int group = 0; valid && group < (int) moleculeGroups.size(); group++) { for (int group = 0; valid && group < (int) moleculeGroups.size(); group++) {
MoleculeGroup& mol = moleculeGroups[group]; MoleculeGroup& mol = moleculeGroups[group];
vector<int>& instances = mol.instances; vector<int>& instances = mol.instances;
...@@ -1075,22 +1108,21 @@ void CudaContext::invalidateMolecules() { ...@@ -1075,22 +1108,21 @@ void CudaContext::invalidateMolecules() {
Molecule& m2 = molecules[instances[j]]; Molecule& m2 = molecules[instances[j]];
int offset2 = offsets[j]; int offset2 = offsets[j];
for (int i = 0; i < (int) atoms.size() && valid; i++) { for (int i = 0; i < (int) atoms.size() && valid; i++) {
for (int k = 0; k < (int) forces.size(); k++) if (!force->areParticlesIdentical(atoms[i]+offset1, atoms[i]+offset2))
if (!forces[k]->areParticlesIdentical(atoms[i]+offset1, atoms[i]+offset2)) valid = false;
valid = false;
} }
// See if the force groups are identical. // See if the force groups are identical.
for (int i = 0; i < (int) forces.size() && valid; i++) { if (valid && forceIndex > -1) {
for (int k = 0; k < (int) m1.groups[i].size() && valid; k++) for (int k = 0; k < (int) m1.groups[forceIndex].size() && valid; k++)
if (!forces[i]->areGroupsIdentical(m1.groups[i][k], m2.groups[i][k])) if (!force->areGroupsIdentical(m1.groups[forceIndex][k], m2.groups[forceIndex][k]))
valid = false; valid = false;
} }
} }
} }
if (valid) if (valid)
return; return false;
// The list of which molecules are identical is no longer valid. We need to restore the // The list of which molecules are identical is no longer valid. We need to restore the
// atoms to their original order, rebuild the list of identical molecules, and sort them // atoms to their original order, rebuild the list of identical molecules, and sort them
...@@ -1158,6 +1190,7 @@ void CudaContext::invalidateMolecules() { ...@@ -1158,6 +1190,7 @@ void CudaContext::invalidateMolecules() {
for (int i = 0; i < (int) reorderListeners.size(); i++) for (int i = 0; i < (int) reorderListeners.size(); i++)
reorderListeners[i]->execute(); reorderListeners[i]->execute();
reorderAtoms(); reorderAtoms();
return true;
} }
void CudaContext::reorderAtoms() { void CudaContext::reorderAtoms() {
......
platforms/cuda/src/CudaKernels.cpp
View file @ 2c559596
This diff is collapsed.
platforms/cuda/src/kernels/utilities.cu
View file @ 2c559596
...@@ -73,4 +73,11 @@ __global__ void clearSixBuffers(int* __restrict__ buffer1, int size1, int* __res ...@@ -73,4 +73,11 @@ __global__ void clearSixBuffers(int* __restrict__ buffer1, int size1, int* __res
clearSingleBuffer(buffer6, size6); clearSingleBuffer(buffer6, size6);
} }
/**
* Record the atomic charges into the posq array.
*/
__global__ void setCharges(real* __restrict__ charges, real4* __restrict__ posq, int* __restrict__ atomOrder, int numAtoms) {
for (int i = blockDim.x*blockIdx.x+threadIdx.x; i < numAtoms; i += blockDim.x*gridDim.x)
posq[i].w = charges[atomOrder[i]];
}
} }
\ No newline at end of file
platforms/opencl/include/OpenCLContext.h
View file @ 2c559596
...@@ -609,6 +609,10 @@ public: ...@@ -609,6 +609,10 @@ public:
OpenCLNonbondedUtilities& getNonbondedUtilities() { OpenCLNonbondedUtilities& getNonbondedUtilities() {
return *nonbonded; return *nonbonded;
} }
/**
* Set the particle charges. These are packed into the fourth element of the posq array.
*/
void setCharges(const std::vector<double>& charges);
/** /**
* Get the thread used by this context for executing parallel computations. * Get the thread used by this context for executing parallel computations.
*/ */
...@@ -692,6 +696,12 @@ public: ...@@ -692,6 +696,12 @@ public:
* and order to be revalidated. * and order to be revalidated.
*/ */
void invalidateMolecules(); void invalidateMolecules();
/**
* Mark that the current molecule definitions from one particular force (and hence the atom order)
* may be invalid. This should be called whenever force field parameters change. It will cause the
* definitions and order to be revalidated.
*/
bool invalidateMolecules(OpenCLForceInfo* force);
private: private:
struct Molecule; struct Molecule;
struct MoleculeGroup; struct MoleculeGroup;
...@@ -739,6 +749,7 @@ private: ...@@ -739,6 +749,7 @@ private:
cl::Kernel clearSixBuffersKernel; cl::Kernel clearSixBuffersKernel;
cl::Kernel reduceReal4Kernel; cl::Kernel reduceReal4Kernel;
cl::Kernel reduceForcesKernel; cl::Kernel reduceForcesKernel;
cl::Kernel setChargesKernel;
std::vector<OpenCLForceInfo*> forces; std::vector<OpenCLForceInfo*> forces;
std::vector<Molecule> molecules; std::vector<Molecule> molecules;
std::vector<MoleculeGroup> moleculeGroups; std::vector<MoleculeGroup> moleculeGroups;
...@@ -754,6 +765,7 @@ private: ...@@ -754,6 +765,7 @@ private:
OpenCLArray* energyBuffer; OpenCLArray* energyBuffer;
OpenCLArray* energyParamDerivBuffer; OpenCLArray* energyParamDerivBuffer;
OpenCLArray* atomIndexDevice; OpenCLArray* atomIndexDevice;
OpenCLArray* chargeBuffer;
std::vector<std::string> energyParamDerivNames; std::vector<std::string> energyParamDerivNames;
std::map<std::string, double> energyParamDerivWorkspace; std::map<std::string, double> energyParamDerivWorkspace;
std::vector<int> atomIndex; std::vector<int> atomIndex;
......
platforms/opencl/include/OpenCLKernels.h
View file @ 2c559596
...@@ -9,7 +9,7 @@ ...@@ -9,7 +9,7 @@
* Biological Structures at Stanford, funded under the NIH Roadmap for * * Biological Structures at Stanford, funded under the NIH Roadmap for *
* Medical Research, grant U54 GM072970. See https://simtk.org. * * Medical Research, grant U54 GM072970. See https://simtk.org. *
* * * *
* Portions copyright (c) 2008-2016 Stanford University and the Authors. * * Portions copyright (c) 2008-2017 Stanford University and the Authors. *
* Authors: Peter Eastman * * Authors: Peter Eastman *
* Contributors: * * Contributors: *
* * * *
...@@ -270,9 +270,11 @@ public: ...@@ -270,9 +270,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force); void copyParametersToContext(ContextImpl& context, const HarmonicBondForce& force);
private: private:
class ForceInfo;
int numBonds; int numBonds;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLArray* params; OpenCLArray* params;
}; };
...@@ -310,9 +312,11 @@ public: ...@@ -310,9 +312,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomBondForce& force); void copyParametersToContext(ContextImpl& context, const CustomBondForce& force);
private: private:
class ForceInfo;
int numBonds; int numBonds;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
...@@ -353,9 +357,11 @@ public: ...@@ -353,9 +357,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force); void copyParametersToContext(ContextImpl& context, const HarmonicAngleForce& force);
private: private:
class ForceInfo;
int numAngles; int numAngles;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLArray* params; OpenCLArray* params;
}; };
...@@ -393,9 +399,11 @@ public: ...@@ -393,9 +399,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomAngleForce& force); void copyParametersToContext(ContextImpl& context, const CustomAngleForce& force);
private: private:
class ForceInfo;
int numAngles; int numAngles;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
...@@ -436,9 +444,11 @@ public: ...@@ -436,9 +444,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force); void copyParametersToContext(ContextImpl& context, const PeriodicTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLArray* params; OpenCLArray* params;
}; };
...@@ -476,9 +486,11 @@ public: ...@@ -476,9 +486,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const RBTorsionForce& force); void copyParametersToContext(ContextImpl& context, const RBTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLArray* params; OpenCLArray* params;
}; };
...@@ -516,9 +528,11 @@ public: ...@@ -516,9 +528,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force); void copyParametersToContext(ContextImpl& context, const CMAPTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
std::vector<mm_int2> mapPositionsVec; std::vector<mm_int2> mapPositionsVec;
OpenCLArray* coefficients; OpenCLArray* coefficients;
...@@ -559,9 +573,11 @@ public: ...@@ -559,9 +573,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force); void copyParametersToContext(ContextImpl& context, const CustomTorsionForce& force);
private: private:
class ForceInfo;
int numTorsions; int numTorsions;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
...@@ -625,12 +641,14 @@ private: ...@@ -625,12 +641,14 @@ private:
const char* getMaxValue() const {return "(int2) (INT_MAX, INT_MAX)";} const char* getMaxValue() const {return "(int2) (INT_MAX, INT_MAX)";}
const char* getSortKey() const {return "value.y";} const char* getSortKey() const {return "value.y";}
}; };
class ForceInfo;
class PmeIO; class PmeIO;
class PmePreComputation; class PmePreComputation;
class PmePostComputation; class PmePostComputation;
class SyncQueuePreComputation; class SyncQueuePreComputation;
class SyncQueuePostComputation; class SyncQueuePostComputation;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLArray* sigmaEpsilon; OpenCLArray* sigmaEpsilon;
OpenCLArray* exceptionParams; OpenCLArray* exceptionParams;
...@@ -704,8 +722,10 @@ public: ...@@ -704,8 +722,10 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force); void copyParametersToContext(ContextImpl& context, const CustomNonbondedForce& force);
private: private:
class ForceInfo;
void initInteractionGroups(const CustomNonbondedForce& force, const std::string& interactionSource, const std::vector<std::string>& tableTypes); void initInteractionGroups(const CustomNonbondedForce& force, const std::string& interactionSource, const std::vector<std::string>& tableTypes);
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
OpenCLArray* interactionGroupData; OpenCLArray* interactionGroupData;
...@@ -756,10 +776,12 @@ public: ...@@ -756,10 +776,12 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const GBSAOBCForce& force); void copyParametersToContext(ContextImpl& context, const GBSAOBCForce& force);
private: private:
class ForceInfo;
double prefactor, surfaceAreaFactor, cutoff; double prefactor, surfaceAreaFactor, cutoff;
bool hasCreatedKernels; bool hasCreatedKernels;
int maxTiles; int maxTiles;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
OpenCLArray* params; OpenCLArray* params;
OpenCLArray* bornSum; OpenCLArray* bornSum;
OpenCLArray* longBornSum; OpenCLArray* longBornSum;
...@@ -807,10 +829,12 @@ public: ...@@ -807,10 +829,12 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomGBForce& force); void copyParametersToContext(ContextImpl& context, const CustomGBForce& force);
private: private:
class ForceInfo;
double cutoff; double cutoff;
bool hasInitializedKernels, needParameterGradient, needEnergyParamDerivs; bool hasInitializedKernels, needParameterGradient, needEnergyParamDerivs;
int maxTiles, numComputedValues; int maxTiles, numComputedValues;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLParameterSet* computedValues; OpenCLParameterSet* computedValues;
OpenCLParameterSet* energyDerivs; OpenCLParameterSet* energyDerivs;
...@@ -864,9 +888,11 @@ public: ...@@ -864,9 +888,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomExternalForce& force); void copyParametersToContext(ContextImpl& context, const CustomExternalForce& force);
private: private:
class ForceInfo;
int numParticles; int numParticles;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
const System& system; const System& system;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
...@@ -908,9 +934,11 @@ public: ...@@ -908,9 +934,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const CustomHbondForce& force); void copyParametersToContext(ContextImpl& context, const CustomHbondForce& force);
private: private:
class ForceInfo;
int numDonors, numAcceptors; int numDonors, numAcceptors;
bool hasInitializedKernel; bool hasInitializedKernel;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
OpenCLParameterSet* donorParams; OpenCLParameterSet* donorParams;
OpenCLParameterSet* acceptorParams; OpenCLParameterSet* acceptorParams;
OpenCLArray* globals; OpenCLArray* globals;
...@@ -961,9 +989,11 @@ public: ...@@ -961,9 +989,11 @@ public:
void copyParametersToContext(ContextImpl& context, const CustomCentroidBondForce& force); void copyParametersToContext(ContextImpl& context, const CustomCentroidBondForce& force);
private: private:
class ForceInfo;
int numGroups, numBonds; int numGroups, numBonds;
bool needEnergyParamDerivs; bool needEnergyParamDerivs;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
OpenCLArray* groupParticles; OpenCLArray* groupParticles;
...@@ -1013,8 +1043,10 @@ public: ...@@ -1013,8 +1043,10 @@ public:
void copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force); void copyParametersToContext(ContextImpl& context, const CustomCompoundBondForce& force);
private: private:
class ForceInfo;
int numBonds; int numBonds;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
OpenCLParameterSet* params; OpenCLParameterSet* params;
OpenCLArray* globals; OpenCLArray* globals;
std::vector<std::string> globalParamNames; std::vector<std::string> globalParamNames;
...@@ -1059,7 +1091,9 @@ public: ...@@ -1059,7 +1091,9 @@ public:
void copyParametersToContext(ContextImpl& context, const CustomManyParticleForce& force); void copyParametersToContext(ContextImpl& context, const CustomManyParticleForce& force);
private: private:
class ForceInfo;
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
bool hasInitializedKernel; bool hasInitializedKernel;
NonbondedMethod nonbondedMethod; NonbondedMethod nonbondedMethod;
int maxNeighborPairs, forceWorkgroupSize, findNeighborsWorkgroupSize; int maxNeighborPairs, forceWorkgroupSize, findNeighborsWorkgroupSize;
...@@ -1119,9 +1153,11 @@ public: ...@@ -1119,9 +1153,11 @@ public:
*/ */
void copyParametersToContext(ContextImpl& context, const GayBerneForce& force); void copyParametersToContext(ContextImpl& context, const GayBerneForce& force);
private: private:
class ForceInfo;
class ReorderListener; class ReorderListener;
void sortAtoms(); void sortAtoms();
OpenCLContext& cl; OpenCLContext& cl;
ForceInfo* info;
bool hasInitializedKernels; bool hasInitializedKernels;
int numRealParticles, maxNeighborBlocks; int numRealParticles, maxNeighborBlocks;
GayBerneForce::NonbondedMethod nonbondedMethod; GayBerneForce::NonbondedMethod nonbondedMethod;
......
platforms/opencl/src/OpenCLContext.cpp
View file @ 2c559596
...@@ -69,8 +69,8 @@ static void CL_CALLBACK errorCallback(const char* errinfo, const void* private_i ...@@ -69,8 +69,8 @@ static void CL_CALLBACK errorCallback(const char* errinfo, const void* private_i
OpenCLContext::OpenCLContext(const System& system, int platformIndex, int deviceIndex, const string& precision, OpenCLPlatform::PlatformData& platformData) : OpenCLContext::OpenCLContext(const System& system, int platformIndex, int deviceIndex, const string& precision, OpenCLPlatform::PlatformData& platformData) :
system(system), time(0.0), platformData(platformData), stepCount(0), computeForceCount(0), stepsSinceReorder(99999), atomsWereReordered(false), posq(NULL), system(system), time(0.0), platformData(platformData), stepCount(0), computeForceCount(0), stepsSinceReorder(99999), atomsWereReordered(false), posq(NULL),
posqCorrection(NULL), velm(NULL), forceBuffers(NULL), longForceBuffer(NULL), energyBuffer(NULL), energyParamDerivBuffer(NULL), atomIndexDevice(NULL), integration(NULL), posqCorrection(NULL), velm(NULL), forceBuffers(NULL), longForceBuffer(NULL), energyBuffer(NULL), energyParamDerivBuffer(NULL), atomIndexDevice(NULL),
expression(NULL), bonded(NULL), nonbonded(NULL), thread(NULL) { chargeBuffer(NULL), integration(NULL), expression(NULL), bonded(NULL), nonbonded(NULL), thread(NULL) {
if (precision == "single") { if (precision == "single") {
useDoublePrecision = false; useDoublePrecision = false;
useMixedPrecision = false; useMixedPrecision = false;
...@@ -309,6 +309,7 @@ OpenCLContext::OpenCLContext(const System& system, int platformIndex, int device ...@@ -309,6 +309,7 @@ OpenCLContext::OpenCLContext(const System& system, int platformIndex, int device
reduceReal4Kernel = cl::Kernel(utilities, "reduceReal4Buffer"); reduceReal4Kernel = cl::Kernel(utilities, "reduceReal4Buffer");
if (supports64BitGlobalAtomics) if (supports64BitGlobalAtomics)
reduceForcesKernel = cl::Kernel(utilities, "reduceForces"); reduceForcesKernel = cl::Kernel(utilities, "reduceForces");
setChargesKernel = cl::Kernel(utilities, "setCharges");
// Decide whether native_sqrt(), native_rsqrt(), and native_recip() are sufficiently accurate to use. // Decide whether native_sqrt(), native_rsqrt(), and native_recip() are sufficiently accurate to use.
...@@ -439,6 +440,8 @@ OpenCLContext::~OpenCLContext() { ...@@ -439,6 +440,8 @@ OpenCLContext::~OpenCLContext() {
delete energyParamDerivBuffer; delete energyParamDerivBuffer;
if (atomIndexDevice != NULL) if (atomIndexDevice != NULL)
delete atomIndexDevice; delete atomIndexDevice;
if (chargeBuffer != NULL)
delete chargeBuffer;
if (integration != NULL) if (integration != NULL)
delete integration; delete integration;
if (expression != NULL) if (expression != NULL)
...@@ -747,6 +750,28 @@ void OpenCLContext::reduceBuffer(OpenCLArray& array, int numBuffers) { ...@@ -747,6 +750,28 @@ void OpenCLContext::reduceBuffer(OpenCLArray& array, int numBuffers) {
executeKernel(reduceReal4Kernel, bufferSize, 128); executeKernel(reduceReal4Kernel, bufferSize, 128);
} }
void OpenCLContext::setCharges(const vector<double>& charges) {
if (chargeBuffer == NULL)
chargeBuffer = new OpenCLArray(*this, numAtoms, useDoublePrecision ? sizeof(double) : sizeof(float), "chargeBuffer");
if (getUseDoublePrecision()) {
double* c = (double*) getPinnedBuffer();
for (int i = 0; i < charges.size(); i++)
c[i] = charges[i];
chargeBuffer->upload(c);
}
else {
float* c = (float*) getPinnedBuffer();
for (int i = 0; i < charges.size(); i++)
c[i] = (float) charges[i];
chargeBuffer->upload(c);
}
setChargesKernel.setArg<cl::Buffer>(0, chargeBuffer->getDeviceBuffer());
setChargesKernel.setArg<cl::Buffer>(1, posq->getDeviceBuffer());
setChargesKernel.setArg<cl::Buffer>(2, atomIndexDevice->getDeviceBuffer());
setChargesKernel.setArg<cl_int>(3, numAtoms);
executeKernel(setChargesKernel, numAtoms);
}
/** /**
* This class ensures that atom reordering doesn't break virtual sites. * This class ensures that atom reordering doesn't break virtual sites.
*/ */
...@@ -945,9 +970,19 @@ void OpenCLContext::findMoleculeGroups() { ...@@ -945,9 +970,19 @@ void OpenCLContext::findMoleculeGroups() {
} }
void OpenCLContext::invalidateMolecules() { void OpenCLContext::invalidateMolecules() {
for (int i = 0; i < forces.size(); i++)
if (invalidateMolecules(forces[i]))
return;
}
bool OpenCLContext::invalidateMolecules(OpenCLForceInfo* force) {
if (numAtoms == 0 || nonbonded == NULL || !nonbonded->getUseCutoff()) if (numAtoms == 0 || nonbonded == NULL || !nonbonded->getUseCutoff())
return; return false;
bool valid = true; bool valid = true;
int forceIndex = -1;
for (int i = 0; i < forces.size(); i++)
if (forces[i] == force)
forceIndex = i;
for (int group = 0; valid && group < (int) moleculeGroups.size(); group++) { for (int group = 0; valid && group < (int) moleculeGroups.size(); group++) {
MoleculeGroup& mol = moleculeGroups[group]; MoleculeGroup& mol = moleculeGroups[group];
vector<int>& instances = mol.instances; vector<int>& instances = mol.instances;
...@@ -962,22 +997,21 @@ void OpenCLContext::invalidateMolecules() { ...@@ -962,22 +997,21 @@ void OpenCLContext::invalidateMolecules() {
Molecule& m2 = molecules[instances[j]]; Molecule& m2 = molecules[instances[j]];
int offset2 = offsets[j]; int offset2 = offsets[j];
for (int i = 0; i < (int) atoms.size() && valid; i++) { for (int i = 0; i < (int) atoms.size() && valid; i++) {
for (int k = 0; k < (int) forces.size(); k++) if (!force->areParticlesIdentical(atoms[i]+offset1, atoms[i]+offset2))
if (!forces[k]->areParticlesIdentical(atoms[i]+offset1, atoms[i]+offset2)) valid = false;
valid = false;
} }
// See if the force groups are identical. // See if the force groups are identical.
for (int i = 0; i < (int) forces.size() && valid; i++) { if (valid && forceIndex > -1) {
for (int k = 0; k < (int) m1.groups[i].size() && valid; k++) for (int k = 0; k < (int) m1.groups[forceIndex].size() && valid; k++)
if (!forces[i]->areGroupsIdentical(m1.groups[i][k], m2.groups[i][k])) if (!force->areGroupsIdentical(m1.groups[forceIndex][k], m2.groups[forceIndex][k]))
valid = false; valid = false;
} }
} }
} }
if (valid) if (valid)
return; return false;
// The list of which molecules are identical is no longer valid. We need to restore the // The list of which molecules are identical is no longer valid. We need to restore the
// atoms to their original order, rebuild the list of identical molecules, and sort them // atoms to their original order, rebuild the list of identical molecules, and sort them
...@@ -1045,6 +1079,7 @@ void OpenCLContext::invalidateMolecules() { ...@@ -1045,6 +1079,7 @@ void OpenCLContext::invalidateMolecules() {
for (int i = 0; i < (int) reorderListeners.size(); i++) for (int i = 0; i < (int) reorderListeners.size(); i++)
reorderListeners[i]->execute(); reorderListeners[i]->execute();
reorderAtoms(); reorderAtoms();
return true;
} }
void OpenCLContext::reorderAtoms() { void OpenCLContext::reorderAtoms() {
......
platforms/opencl/src/OpenCLKernels.cpp
View file @ 2c559596
This diff is collapsed.
platforms/opencl/src/kernels/utilities.cl
View file @ 2c559596
...@@ -107,3 +107,11 @@ __kernel void determineNativeAccuracy(__global float8* restrict values, int numV ...@@ -107,3 +107,11 @@ __kernel void determineNativeAccuracy(__global float8* restrict values, int numV
values[i] = (float8) (v, native_sqrt(v), native_rsqrt(v), native_recip(v), native_exp(v), native_log(v), 0.0f, 0.0f); values[i] = (float8) (v, native_sqrt(v), native_rsqrt(v), native_recip(v), native_exp(v), native_log(v), 0.0f, 0.0f);
} }
} }
/**
* Record the atomic charges into the posq array.
*/
__kernel void setCharges(__global real* restrict charges, __global real4* restrict posq, __global int* restrict atomOrder, int numAtoms) {
for (int i = get_global_id(0); i < numAtoms; i += get_global_size(0))
posq[i].w = charges[atomOrder[i]];
}
\ No newline at end of file
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