Skip to content
Snippets Groups Projects
Commit 3e2ac53a authored by cmaffeo2's avatar cmaffeo2
Browse files

Began reorganizing project

parent 0544da93
No related branches found
No related tags found
No related merge requests found
Showing with 0 additions and 174 deletions
* TODO active
** Fix wrapping routines (trouble with machine precision
** make units and kT more uniform (references to 295 K are hardcoded into RB code and Temperature code)
*** I would prefer to specify temperature in K, rather than specifying a scaling factor to be applied to all energies to make them kT
** handle periodic boundaries for RBs
** RB read restart file
** RB throw exception if an RB type uses a grid key multiple times
** RB each RB gets temperature from grid
* TODO eventually
** enable flag to disable physics-checking assertions in kernels (also write these assertions)
** split computation into multiple regions
*** simulate regions, also in overlap
**** atoms at edge of overlap can have DoF frozen
*** periodically synchronize positions
*** optimize to minimize communication & error
**** monitor energy jumps associated with this process
** organize code a bit better
** increase cells/cutoff
** improve pairlist algorithm
http://arxiv.org/pdf/1306.1737.pdf
** RB optionally don't eval forces every ts
** pairlists for rigid bodies
*** maybe for grids, depending on parallel structure of code
* Organization of code
** RgidBodyController (RBC) holds device pointers, manages force evaluation and integration
* other ideas
** interpolate density grid?
BaseGrid.h: // RBTODO Fix?
BaseGrid.h- BaseGrid(); // cmaffeo2 (2015) moved this out of protected, cause I wanted BaseGrid in a struct
BaseGrid.h- // The most obvious of constructors.
BaseGrid.h- BaseGrid(Matrix3 basis0, Vector3 origin0, int nx0, int ny0, int nz0);
BaseGrid.h-
--
BaseGrid.h- float a0, a1, a2, a3;
BaseGrid.h-
BaseGrid.h: // RBTODO parallelize loops?
BaseGrid.h-
BaseGrid.h- // Mix along x, taking the derivative.
BaseGrid.h- float g2[4][4];
BaseGrid.h- for (int iy = 0; iy < 4; iy++) {
--
BaseGrid.h- }
BaseGrid.h-
BaseGrid.h: // RBTODO overload with optimized algorithm
BaseGrid.h- // skip transforms (assume identity basis)
BaseGrid.h- HOST DEVICE inline float interpolatePotential(Vector3 pos) const {
BaseGrid.h- // Find the home node.
BaseGrid.h- Vector3 l = basisInv.transform(pos - origin);
--
BaseGrid.h-
BaseGrid.h- // out of grid? return 0
BaseGrid.h: // RBTODO
BaseGrid.h-
BaseGrid.h- // Get the array jumps.
BaseGrid.h- int jump[3];
BaseGrid.h- jump[0] = nz*ny;
--
BaseGrid.h- // Find the values at the neighbors.
BaseGrid.h- float g1[4][4][4];
BaseGrid.h: //RBTODO parallelize?
BaseGrid.h- for (int ix = 0; ix < 4; ix++) {
BaseGrid.h- for (int iy = 0; iy < 4; iy++) {
BaseGrid.h- for (int iz = 0; iz < 4; iz++) {
BaseGrid.h- // Wrap around the periodic boundaries.
--
ComputeGridGrid.cuh- Matrix3 basis_rho, Vector3 origin_rho,
ComputeGridGrid.cuh- Matrix3 basis_u, Vector3 origin_u) {
ComputeGridGrid.cuh: // RBTODO http://devblogs.nvidia.com/parallelforall/cuda-pro-tip-write-flexible-kernels-grid-stride-loops/
ComputeGridGrid.cuh- const unsigned int r_id = blockIdx.x * blockDim.x + threadIdx.x;
ComputeGridGrid.cuh-
ComputeGridGrid.cuh: // RBTODO parallelize transform
ComputeGridGrid.cuh- if (r_id > rho->size) // skip threads with no data
ComputeGridGrid.cuh- return;
ComputeGridGrid.cuh-
ComputeGridGrid.cuh- // Maybe: Tile grid data into shared memory
ComputeGridGrid.cuh: // RBTODO: think about localizing regions of grid data
ComputeGridGrid.cuh- Vector3 p = rho->getPosition(r_id, basis, origin);
ComputeGridGrid.cuh- float val = rho->val[r_id];
ComputeGridGrid.cuh-
ComputeGridGrid.cuh: // RBTODO reduce forces and torques
ComputeGridGrid.cuh- // http://www.cuvilib.com/Reduction.pdf
ComputeGridGrid.cuh-
ComputeGridGrid.cuh: // RBTODO combine interp methods and reduce repetition!
ComputeGridGrid.cuh- float energy = u->interpolatePotential(p);
ComputeGridGrid.cuh- Vector3 f = u->interpolateForceD(p);
ComputeGridGrid.cuh- Vector3 t = cross(p,f); // test if sign is correct!
ComputeGridGrid.cuh-
ComputeGridGrid.cuh: // RBTODO 3rd-law forces + torques
ComputeGridGrid.cuh-}
--
Configuration.cpp- cudaMemcpyHostToDevice));
Configuration.cpp- }
Configuration.cpp: // RBTODO: moved this out of preceding loop; was that correct?
Configuration.cpp- gpuErrchk(cudaMemcpyAsync(part_d, part_addr, sizeof(BrownianParticleType*) * numParts,
Configuration.cpp- cudaMemcpyHostToDevice));
Configuration.cpp-
Configuration.cpp-
--
Configuration.cpp- sz = sizeof(float) * len;
Configuration.cpp- gpuErrchk(cudaMemcpy( tmpData, g->val, sz, cudaMemcpyHostToDevice));
Configuration.cpp: // RBTODO: why can't this be deleted?
Configuration.cpp- // delete[] tmpData;
Configuration.cpp- }
Configuration.cpp- }
Configuration.cpp-
--
Configuration.cpp- sz = sizeof(float) * len;
Configuration.cpp- gpuErrchk(cudaMemcpy( tmpData, g->val, sz, cudaMemcpyHostToDevice));
Configuration.cpp: // RBTODO: why can't this be deleted?
Configuration.cpp- // delete[] tmpData;
Configuration.cpp- }
Configuration.cpp-
Configuration.cpp- }
--
RigidBodyGrid.h- \===============================*/
RigidBodyGrid.h-
RigidBodyGrid.h: // RBTODO Fix?
RigidBodyGrid.h- RigidBodyGrid(); // cmaffeo2 (2015) moved this out of protected, cause I wanted RigidBodyGrid in a struct
RigidBodyGrid.h- // The most obvious of constructors.
RigidBodyGrid.h- RigidBodyGrid(int nx0, int ny0, int nz0);
RigidBodyGrid.h-
--
RigidBodyGrid.h- float a0, a1, a2, a3;
RigidBodyGrid.h-
RigidBodyGrid.h: // RBTODO further parallelize loops? unlikely?
RigidBodyGrid.h-
RigidBodyGrid.h- // Mix along x, taking the derivative.
RigidBodyGrid.h- float g2[4][4];
RigidBodyGrid.h- for (int iy = 0; iy < 4; iy++) {
--
RigidBodyGrid.h-
RigidBodyGrid.h- // out of grid? return 0
RigidBodyGrid.h: // RBTODO
RigidBodyGrid.h-
RigidBodyGrid.h- // Get the array jumps.
RigidBodyGrid.h- int jump[3];
RigidBodyGrid.h- jump[0] = nz*ny;
--
RigidBodyGrid.h- w[2] = l.z - homeZ;
RigidBodyGrid.h- // Find the values at the neighbors.
RigidBodyGrid.h: float g1[4][4][4]; /* RBTODO: inefficient for my algorithm? */
RigidBodyGrid.h- for (int ix = 0; ix < 4; ix++) {
RigidBodyGrid.h- int jx = ix-1 + home[0];
RigidBodyGrid.h- for (int iy = 0; iy < 4; iy++) {
RigidBodyGrid.h- int jy = iy-1 + home[1];
--
RigidBodyGrid.h- // Assume zero value at edges
RigidBodyGrid.h- int jz = iz-1 + home[2];
RigidBodyGrid.h: // RBTODO: possible branch divergence in warp?
RigidBodyGrid.h- if (jx < 0 || jy < 0 || jz < 0 ||
RigidBodyGrid.h- jx >= nx || jz >= nz || jz >= nz) {
RigidBodyGrid.h- g1[ix][iy][iz] = 0;
RigidBodyGrid.h- } else {
--
RigidBodyGrid.h- // Find the values at the neighbors.
RigidBodyGrid.h- float g1[4][4][4];
RigidBodyGrid.h: //RBTODO parallelize?
RigidBodyGrid.h- for (int ix = 0; ix < 4; ix++) {
RigidBodyGrid.h- for (int iy = 0; iy < 4; iy++) {
RigidBodyGrid.h- for (int iz = 0; iz < 4; iz++) {
RigidBodyGrid.h- // Wrap around the periodic boundaries.
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
File moved
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment