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

2021
Views:

356

Pages:

20
Author:

Clark Serrano

Input Binning 1 Objective • To understand how data scalability problems in gather parallel execution motivate input binning • To learn basic input binning techniques • To understand common tradeoffs in input binning 2 Scatter to Gather Transformation in Thread 1 Thread 2 … in out Thread 1 out Thread 2 … 3 However • Input tend to be much less regular than output – It may be difficult for each thread to efficiently locate all inputs relevant to its output – Or, to efficiently exclude all inputs irrelevant to its output • In a naïve arrangement, all threads may have to process all inputs to decide if each input is relevant to its output – This makes execution time scale poorly with data set size – data scalability problem – Especially a problem for many-cores designed to process large data sets 4 DCS Algorithm for Electrostatic Potentials Revisited • At each grid point, sum the electrostatic potential from all atoms – All threads read all inputs • Highly data-parallel • But has quadratic complexity – Number of grid points × number of atoms – Both proportional to volume – Poor data scalability in terms of volume 5 Direct Summation is accurate but has poor data scalability Same scalability among all cutoff implementations Scalability and Performance of different algorithms for calculating electrostatic potential map. 6 Algorithm for Electrostatic Potentials With a Cutoff • Ignore atoms beyond a cutoff distance, rc – Typically 8Å–12Å – Long-range potential may be computed separately • Number of atoms within cutoff distance is roughly constant (uniform atom density) – 200 to 700 atoms within 8Å–12Å cutoff sphere for typical biomolecular structures 7 Cut-off Summation • With fixed partial charge qi, electrostatic potential V at position r over all N atoms: , where 8 Implementation Challenge • For each tile of grid points, we need to identify the set of atoms that need to be examined – One could naively examine all atoms and only use the ones whose distance is within the given range – But this examination still takes time, and brings the time complexity right back to • number of atoms * number of grid points – Each thread needs to avoid examining the atoms outside the range of its grid point(s) 9 Binning • A process that groups data to form a chunk called bin • Each bin collectively represents a property for data in the bin • Helps problem solving due to data coarsening • Uniform bin arrays, Variable bins, KD Trees, … 10 Binning for Cut-Off Potential • Divide the simulation volume with nonoverlapping uniform cubes • Every atom in the simulation volume falls into a cube based on its spatial location – Bins represent location property of atoms • After binning, each cube has a unique index in the simulation space for easy parallel access 11 Spatial Sorting Using Binning • Presort atoms into bins by location in space • Each bin holds several atoms • Cutoff potential only uses bins within rc – Yields a linear complexity cutoff potential algorithm – Some atoms will be examined by a thread but not used 12 Terminology • Bin Size – The size of the bin cubes that partition the simulation volume – The bigger the bin size, the more the atoms that will fall into each bin • Bin Capacity – The number of atoms that can be accommodated by each bin in the array implementation 13 Bin Capacity Considerations • Capacity of atom bins needs to be balanced – Too large – many dummy atoms in bins – Too small – some atoms will not fit into bins – Target bin capacity to cover more than 95% or atoms • Place all atoms that do not fit into bins into an overflow bin – Use a CPU sequential algorithm to calculate their contributions to the energy grid lattice points. – CPU and GPU can do potential calculations in parallel 14 Bin Design • Uniform sized/capacity bins allow array implementation – And the relative offset list approach • Bin capacity should be big enough to contain all the atoms that fall into a bin – Cut-off will screen away atoms that weren’t processed – Performance penalty if too many are screened away 15 Going from DCS Kernel to Large Bin Cut-off Kernel • Adaptation of techniques from the direct Coulomb summation kernel for a cutoff kernel • Atoms are stored in constant memory as with DCS kernel • CPU loops over potential map regions that are (24Å)3 in volume (cube containing cutoff sphere) – Each map region requires only one large bin of atoms – For each map region, atoms in the corresponding large bin are appended to the constant memory atom buffer until full, then GPU kernel is launched – Host continue to reload constant memory and launch GPU kernels until all atoms in the corresponding large bin are consumed 16 Large Bin Design Concept • Map regions are (24Å)3 in volume • Regions are sized large enough to provide the GPU enough work in a single kernel launch – (48 lattice points)3 for lattice with 0.5Å spacing – (20 atoms)3 on average for each large bin • Bin size and capacity are designed to allow each kernel launch to cover enough lattice points to justify the kernel launch overhead and fully utilize the GPU hardware 17 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); } … 18 Large-bin Cutoff Kernel Evaluation • 6× speedup relative to fast CPU version • Work-inefficient – Coarse spatial hashing into (24Å)3 bins – Only 6.5% of the atoms a thread tests are within the cutoff distance • Small-bin designs improve work efficiency but requires more sophisticated kernel code 19 Small-bin Kernels – 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 M N A • 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) 20

Input Binning Lecture

Related Documents

Recommended Documents

Description

Related Documents

Free up your schedule!

Take 5 seconds to unlock