A hydrologist models pollutant diffusion in a lake with the explicit scheme on a grid…
PHYSICS · Multi-part · Fermi University Physics · Computational Physics · High-performance and data-driven computational physics · High-performance and data-driven computational physics
Scientific Computing for Physicists, Ch Graphics processors: kernels, memory and the roofline
Problem
A hydrologist models pollutant diffusion in a lake with the explicit scheme on a 3072×3072 grid of binary64 values, stored in C order, and tracks the concentration ui,jn at the injection cell. One kernel maps threadIdx.x to the fast (column) index, so a warp reads 32 consecutive entries of a row; a careless port instead maps threadIdx.x to the slow (row) index, so a warp reads 32 consecutive rows at fixed column, each thread's load a single binary64 value.
((a)) Predict how many 32 byte sectors one warp's load fetches in the mis-mapped version, and the factor by which this exceeds the coalesced version.
((b)) The device has FP64 peak 9.7TFLOPs−1 and data-sheet peak bandwidth 1555GBs−1; the coalesced stencil (7 flops, 16B moved per cell) already runs at its bandwidth-limited cap. Assuming the mis-mapped version is limited only by the extra traffic found in (a), predict its attained rate in TFLOPs−1.
((c)) The injection-cell value is stored in float16 once it falls into [0.25,0.5), which happens near step 600. A binary64 reference run shows the decrement per step near this range fits Δun≈6.0×104/n3. Using the float16 rounding-unit argument, predict the step at which the stored value stops changing.
((d)) State what you would run on your own device to check parts (a)-(c), and what result in each case would show a prediction wrong.