diff --git a/src/ComputeForce.cu b/src/ComputeForce.cu index 54a6ca222d01c5604e26fc69b23c1953acdd41b4..95b5b23c57bf164ad552462caf666bb9ec915609 100644 --- a/src/ComputeForce.cu +++ b/src/ComputeForce.cu @@ -4,6 +4,7 @@ #include "ComputeForce.h" #include "ComputeForce.cuh" +#include "Configuration.h" #include <cuda_profiler_api.h> @@ -21,55 +22,65 @@ void runSort(int2 *d1, int *d2, float *key, int2 *scratch1, int *scratch2, float *scratchKey, unsigned int count); -ComputeForce::ComputeForce(int num, const BrownianParticleType part[], - int numParts, const BaseGrid* g, float switchStart, - float switchLen, float electricConst, - int fullLongRange, int numBonds, int numTabBondFiles, - int numExcludes, int numAngles, int numTabAngleFiles, - int numDihedrals, int numTabDihedralFiles, - float pairlistDistance, int numReplicas) : - num(num), numParts(numParts), sys(g), switchStart(switchStart), - switchLen(switchLen), electricConst(electricConst), - cutoff2((switchLen + switchStart) * (switchLen + switchStart)), - decomp(g->getBox(), g->getOrigin(), switchStart + switchLen, numReplicas), - numBonds(numBonds), numTabBondFiles(numTabBondFiles), - numExcludes(numExcludes), numAngles(numAngles), - numTabAngleFiles(numTabAngleFiles), numDihedrals(numDihedrals), - numTabDihedralFiles(numTabDihedralFiles), numReplicas(numReplicas) { +ComputeForce::ComputeForce(const Configuration& c, const int numReplicas) : + num(c.num), numParts(c.numParts), sys(c.sys), switchStart(c.switchStart), + switchLen(c.switchLen), electricConst(c.coulombConst), + cutoff2((c.switchLen + c.switchStart) * (c.switchLen + c.switchStart)), + decomp(c.sys->getBox(), c.sys->getOrigin(), c.switchStart + c.switchLen, numReplicas), + numBonds(c.numBonds), numTabBondFiles(c.numTabBondFiles), + numExcludes(c.numExcludes), numAngles(c.numAngles), + numTabAngleFiles(c.numTabAngleFiles), numDihedrals(c.numDihedrals), + numTabDihedralFiles(c.numTabDihedralFiles), numReplicas(numReplicas) { + +// int num, const BrownianParticleType part[], +// int numParts, const BaseGrid* g, float switchStart, +// float switchLen, float electricConst, +// int fullLongRange, int numBonds, int numTabBondFiles, +// int numExcludes, int numAngles, int numTabAngleFiles, +// int numDihedrals, int numTabDihedralFiles, +// float pairlistDistance, int numReplicas) : +// num(num), numParts(numParts), sys(g), switchStart(switchStart), +// switchLen(switchLen), electricConst(electricConst), +// cutoff2((switchLen + switchStart) * (switchLen + switchStart)), +// decomp(g->getBox(), g->getOrigin(), switchStart + switchLen, numReplicas), +// numBonds(numBonds), numTabBondFiles(numTabBondFiles), +// numExcludes(numExcludes), numAngles(numAngles), +// numTabAngleFiles(numTabAngleFiles), numDihedrals(numDihedrals), +// numTabDihedralFiles(numTabDihedralFiles), numReplicas(numReplicas) { // Allocate the parameter tables. + decomp_d = NULL; - pairlistdist2 = (sqrt(cutoff2) + pairlistDistance); + pairlistdist2 = (sqrt(cutoff2) + c.pairlistDistance); pairlistdist2 *= pairlistdist2; - tableEps = new float[numParts * numParts]; - tableRad6 = new float[numParts * numParts]; - tableAlpha = new float[numParts * numParts]; + int np2 = numParts*numParts; + tableEps = new float[np2]; + tableRad6 = new float[np2]; + tableAlpha = new float[np2]; - const size_t tableSize = sizeof(float) * numParts * numParts; + const size_t tableSize = sizeof(float) * np2; gpuErrchk(cudaMalloc(&tableEps_d, tableSize)); gpuErrchk(cudaMalloc(&tableRad6_d, tableSize)); gpuErrchk(cudaMalloc(&tableAlpha_d, tableSize)); gpuErrchk(cudaMalloc(&sys_d, sizeof(BaseGrid))); gpuErrchk(cudaMemcpyAsync(sys_d, sys, sizeof(BaseGrid), cudaMemcpyHostToDevice)); - // Form the parameter tables. - makeTables(part); - gpuErrchk(cudaMemcpyAsync(tableAlpha_d, tableAlpha, tableSize, - cudaMemcpyHostToDevice)); - gpuErrchk(cudaMemcpyAsync(tableEps_d, tableEps, tableSize, - cudaMemcpyHostToDevice)); - gpuErrchk(cudaMemcpyAsync(tableRad6_d, tableRad6, tableSize, - cudaMemcpyHostToDevice)); + // Build the parameter tables. + makeTables(c.part); + + gpuErrchk(cudaMemcpyAsync(tableAlpha_d, tableAlpha, tableSize, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpyAsync(tableEps_d, tableEps, tableSize, cudaMemcpyHostToDevice)); + gpuErrchk(cudaMemcpyAsync(tableRad6_d, tableRad6, tableSize, cudaMemcpyHostToDevice)); // Create the potential table - tablePot = new TabulatedPotential*[numParts * numParts]; - tablePot_addr = new TabulatedPotential*[numParts * numParts]; - for (int i = 0; i < numParts*numParts; i++) { + tablePot = new TabulatedPotential*[np2]; + tablePot_addr = new TabulatedPotential*[np2]; + for (int i = 0; i < np2; ++i) { tablePot_addr[i] = NULL; tablePot[i] = NULL; } - gpuErrchk(cudaMalloc(&tablePot_d, sizeof(TabulatedPotential*) * numParts * numParts)); + gpuErrchk(cudaMalloc(&tablePot_d, sizeof(TabulatedPotential*) * np2)); // Create the bond table tableBond = new TabulatedPotential*[numTabBondFiles]; diff --git a/src/ComputeForce.h b/src/ComputeForce.h index dc20c4e20637425e2263594bb231f754091fe6ad..fb103462e16c3f6c72661968944e56ee4574c37a 100644 --- a/src/ComputeForce.h +++ b/src/ComputeForce.h @@ -30,15 +30,13 @@ const unsigned int NUM_THREADS = 256; +class Configuration; + class ComputeForce { public: - ComputeForce(int num, const BrownianParticleType* part, int numParts, - const BaseGrid* g, float switchStart, float switchLen, - float electricConst, int fullLongRange, int numBonds, - int numTabBondFiles, int numExcludes, int numAngles, int numTabAngleFiles, - int numDihedrals, int numTabDihedralFiles, float pairlistDistance, int numReplicas = 1); - ~ComputeForce(); - + ComputeForce(const Configuration &c, const int numReplicas); + ~ComputeForce(); + void updateNumber(int newNum); void makeTables(const BrownianParticleType* part); @@ -140,6 +138,7 @@ public: static EnergyForce softcoreForce(Vector3 r, float eps, float rad6); private: + // Configuration* c; int numReplicas; int num; int numParts; diff --git a/src/GrandBrownTown.cu b/src/GrandBrownTown.cu index 9584117fd8a7ec1b244f5102e3997831e5702d10..41ff55b1317c4c9f90695d69a0df58dee4a4b6c0 100644 --- a/src/GrandBrownTown.cu +++ b/src/GrandBrownTown.cu @@ -178,9 +178,10 @@ GrandBrownTown::GrandBrownTown(const Configuration& c, const char* outArg, } // Prepare internal force computation - internal = new ComputeForce(num, part, numParts, sys, switchStart, switchLen, coulombConst, - fullLongRange, numBonds, numTabBondFiles, numExcludes, numAngles, numTabAngleFiles, - numDihedrals, numTabDihedralFiles, c.pairlistDistance, numReplicas); + // internal = new ComputeForce(num, part, numParts, sys, switchStart, switchLen, coulombConst, + // fullLongRange, numBonds, numTabBondFiles, numExcludes, numAngles, numTabAngleFiles, + // numDihedrals, numTabDihedralFiles, c.pairlistDistance, numReplicas); + internal = new ComputeForce(c, numReplicas); //MLog: I did the other halve of the copyToCUDA function from the Configuration class here, keep an eye on any mistakes that may occur due to the location. internal -> copyToCUDA(c.simNum, c.type, c.bonds, c.bondMap, c.excludes, c.excludeMap, c.angles, c.dihedrals);