University:
Boston University
Course:
CAS PY 105 | Physics 1
Academic year:

2022
Views:

47

Pages:

19
Author:

Jayson Mcconnell

The Neighborhood is a volume • Calculating and specifying all bin indexes of the sphere can be quite complex – Rough approximations reduce efficiency 21 Neighborhood Offset List (Pre-calculated) • A list of relative offsets enumerating the bins that are located within the cutoff distance for a given location in the simulation volume • Detection of surrounding atoms becomes realistic for output grid points – By visiting bins in the neighborhood offset list and iterating atoms they contain a bin in the (1, 2) neighborhood list (-1, -1) cutoff distance not included center (0, 0) 22 Large Sized Bin Design • Each bin covers a simulation volume for all grid points processed by a thread block – All input atoms potentially needed – It also need to have sufficient capacity • The efficiency is quite low, about 6.5% in CP 23 Large Bin Cut-off Kernel Code static __constant__ float4 atominfo[MAXATOMS]; __global__ static void mgpot_shortrng_energy(…) { […] for (n = 0; n < natoms; n++) { float dx = coorx - atominfo[n].x; float dy = coory - atominfo[n].y; float dz = coorz - atominfo[n].z; float q = atominfo[n].w; float dxdy2 = dx*dx + dy*dy; float r2 = dxdy2 + dz*dz; if (r2 < CUTOFF2) { float gr2 = GC0 + r2*(GC1 + r2*GC2); float r_1 = 1.f/sqrtf(r2); accum_energy_z0 += q * (r_1 - gr2); } … 24 Improving Work Efficiency • Thread block examines atom bins up to the cutoff distance – Use a sphere of bins – All threads in a block scan the same bins and atoms • No hardware penalty for multiple simultaneous reads of the same address • Simplifies fetching of data – The sphere has to be big enough to cover all grid point at corners – There will be a small level of divergence • Not all grid points processed by a thread block relate to all atoms in a bin the same way • (A within cut-off distance of N but outside cut-off of M) M N A 25 The Neighborhood is a volume • Calculating and specifying all bin indexes of the sphere can be quite complex – Rough approximations reduce efficiency 26 Neighborhood Offset List (Pre-calculated) • A list of relative offsets enumerating the bins that are located within the cutoff distance for a given location in the simulation volume • Detection of surrounding atoms becomes realistic for output grid points – By visiting bins in the neighborhood offset list and iterating atoms they contain a bin in the (1, 2) neighborhood list (-1, -1) cutoff distance not included center (0, 0) 27 Pseudo Code of an Implementation // 1. binning for each atom in the simulation volume, index_of_bin := atom.addr / BIN_SIZE bin[index_of_bin] += atom // 2. generate the neighborhood offset list for each c from -cutoff to cutoff, if distance(0, c) < cutoff, nlist += c CPU // 3. do the computation GPU for each point in the output grid, index_of_bin := point.addr / BIN_SIZE for each offset in nlist, for each atom in bin[index_of_bin + offset], point.potential += atom.charge / (distance from point to atom) 28 Performance • O(MN’) where M and N’ are the number of output grid points and atoms in the neighborhood offset list, respectively – In general, N’ is small compared to the number of all atoms • Works well if the distribution of atoms is uniform 29 Details on Small Bin Design • For 0.5Å lattice spacing, a (4Å)3 cube of the potential map is computed by each thread block – 8×8×8 potential map points – 128 threads per block (4 points/thread) – 34% of examined atoms are within cutoff distance 30 More Design Considerations for the Cutoff Kernel • High memory throughput to atom data essential – Group threads together for locality – Fetch bins of data into shared memory – Structure atom data to allow fetching • After taking care of memory demand, optimize to reduce instruction count – Loop and instruction-level optimization 31 Tiling Atom Data • Shared memory used to reduce Global Memory bandwidth consumption – Threads in a thread block collectively load one bin at a time into shared memory – Once loaded, threads scan atoms in shared memory – Reuse: Loaded bins used 128 times Time Data returned from global memory Ready Threads individually Collectively compute potentials load next using bin in shared mem bin Suspend Execution cycle of a thread block Write bin to shared memory Another thread block runs while this one waits 32 Coalesced Global Memory Access to Atom Data • Full global memory bandwidth only with 64byte, 64-byte-aligned memory accesses – Each bin is exactly 128 bytes – Bins stored in a 3D array – 32 threads in each block load one bin into shared memory, then processed by all threads in the block • 128 bytes = 8 atoms (x,y,z,q) – Nearly uniform density of atoms in typical systems • 1 atom per 10 Å3 – Bins hold atoms from (4Å)3 of space (example) – Number of atoms in a bin varies • For water test systems, 5.35 atoms in a bin on average • Some bins overfull 33 Handling Overfull Bins • In typical use, 2.6% of atoms exceed bin capacity • Spatial sorting puts these into a list of extra atoms • Extra atoms processed by the CPU – Computed with CPU-optimized algorithm – Takes about 66% as long as GPU computation – Overlapping GPU and CPU computation yields in additional speedup – CPU performs final integration of grid data 34 CPU Grid Data Integration • Effect of overflow atoms are added to the CPU master energygrid array • Slice of grid point values calculated by GPU are added into the master energygrid array while removing the padded elements 0,0 0,1 … 1,0 1,1 … … … … 35 GPU Thread Coarsening • Each thread computes potentials at four potential map points – Reuse x and z components of distance calculation – Check x and z components against cutoff distance (cylinder test) • Exit inner loop early upon encountering the first empty slot in a bin 36 GPU Thread Inner Loop Exit when an empty atom bin entry is encountered aq = AtomBinCache[i].w; if (aq == 0) break; dx = AtomBinCache[i].x - x; dz = AtomBinCache[i].z - z; dxdz2 = dx*dx + dz*dz; if (dxdz2 < cutoff2) continue; Cylinder test dy = AtomBinCache[i].y - y; r2 = dy*dy + dxdz2; if (r2 < cutoff2) poten0 += aq * rsqrtf(r2); // Simplified example Cutoff test and potential value calculation dy = dy - 2 * grid_spacing; /* Repeat three more times */ } 37 Cutoff Summation Runtime GPU cutoff with CPU overlap: 12x-21x faster than CPU core 50k–1M atom structure size 38 Summary • Large bins allow re-use of all-input kernels with little code change – But work efficiency can be very low • Use of small-sized bins require more sophisticated kernel code to traverse list of small bins – Much higher work efficiency – Small bins also serve as tiles for locality • CPU process overflow atoms from fixed capacity bins 39

Input Binning Lecture Chapter 2

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock