A group models heat conduction in a copper plate on a grid stored in C order as…
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 group models heat conduction in a copper plate on a 5120×5120 grid stored in C order as binary64 (8 byte) values. Their diffusion-stencil kernel computes the row i=blockIdx.y⋅blockDim.y+threadIdx.y and the column j=blockIdx.x⋅blockDim.x+threadIdx.x of the cell ui,j, with blocks of 32×8 threads, 32 along x. A warp is formed from 32 consecutive threadIdx.x at fixed threadIdx.y, so it reads 32 consecutive entries of row i. A collaborator, porting habits from a column-major solver, swaps the two formulas, so that threadIdx.x, which still runs fastest within a warp, now selects the row: each thread of the warp reads a different row at the same column.
((a)) Predict, in 32 byte sectors and in bytes, what one warp's load of ui,j fetches in the second (mis-mapped) version, and the factor by which this exceeds the first (coalesced) version.
((b)) The card's measured triad bandwidth is 1400GBs−1 and the coalesced stencil moves the perfect-reuse 16B per cell update. Assuming each version is limited only by the memory traffic found in part (a), predict the attained update rate, in cells per second, for the coalesced and for the mis-mapped version.
((c)) If the whole run were instead carried out in float32 (4 byte values), predict the new ratio between the mis-mapped and the coalesced version's bytes fetched per warp, using the same sector argument, and say briefly why it differs from the binary64 ratio.
((d)) State what you would time on your own device to check parts (a)-(c), and what outcome in each case would show the prediction wrong.