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
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.