Our sample program uses CUDA events to report the performance of some number of consecutive kernel launches (default 100), and reports the rates of both output coefficients (which varies with the template size) and the “template-pixel” rate, or the number of inner loop iterations per unit time.
The raw performance of GPUs at this computation is astonishing. A mid-range GeForce RTX 3060 (Ampere) runs the shared-memory kernel at more than 320 billion template-pixel calculations per second (Gtpix/s) for a 16×16 template; the DP4A kernel developed in Section 15.5 pushes that past 1900 Gtpix/s—well over an order of magnitude beyond the flagship GPUs of a decade ago.
The default parameters of the program are not ideal for performance
measurement; they are set to detect the dime in the lower right corner,
and optionally write out the image of Figure 15-3. In particular, the
image is too small to keep the GPU fully busy. The image is only 300×246
pixels (74K in size), so only 310 blocks are needed by the shared memory
implementation to perform the computation. The --padWidth and
--padHeight command-line options can be used in the sample program to
increase the size of the image and hence the number of correlation
coefficients computed (there are no data dependencies in the code, so
the padding can be filled with arbitrary data); a 1024×1024 image is
both more realistic and keeps a modern GPU well utilized.
The following table summarizes the relative performance of our four implementations, measured on a GeForce RTX 3060 with the 1024×1024 padded image and a 16×16 template:
| Implementation | Rate (Gtpix/s) |
|---|---|
corrTexTex() |
104 |
corrTexConstant() |
203 |
corrShared() |
324 |
| corrShared4Dp4a | 333 |
| corrShared4Dp4aIlp | 1932 |
Moving the template into constant memory is the single biggest structural win, nearly doubling throughput over the texture-texture baseline (104 to 203 Gtpix/s); staging the image through shared memory adds another 60% (203 to 324 Gtpix/s). From there, the two-stage rewrite of Section 15.5 is transformative. DP4A on its own leaves the latency-bound kernel essentially unchanged (333 Gtpix/s); but once the two-column instruction-level parallelism makes the kernel arithmetic-bound, DP4A reaches 1932 Gtpix/s—about 6× the plain shared-memory version, and nearly 19× the texture-texture baseline.
The size of the template also bears on efficiency: the larger the
template, the more efficient the computation on a per-template-pixel
basis, since the fixed per-thread setup is amortized across more
inner-loop work. The corrShared4Dp4aIlp rate climbs with
template size on the RTX 3060:
| Template | Rate (Gtpix/s) |
|---|---|
| 8×8 | 1192 |
| 12×12 | 1768 |
| 16×16 | 1933 |
| 20×20 | 2029 |
| 24×24 | 2193 |
| 28×28 | 2106 |
Gonzalez and Woods includes both a discussion of normalized correlation (pp. 583-586) and the logarithmic transform used to compute the output pixels in our sample program (pp. 168-169). Gonzalez, Rafael C. and Richard E. Woods. Digital Image Processing. Addison-Wesley, 1992.
J.P. Lewis gives the reference treatment of the FFT numerator and the running-sum (summed-area-table) denominator used in Sections 15.7 and 15.8. Lewis, J.P. “Fast Normalized Cross-Correlation.” Vision Interface, pp. 120-123, 1995.
Kai Briechle and Uwe D. Hanebeck accelerate the numerator a different way, approximating the template as a sum of rectangular basis functions evaluated through summed-area tables—an approach that extends naturally to matching template variants such as rotations and scales. Briechle, K. and Hanebeck, U.D. “Template Matching using Fast Normalized Cross Correlation.” Proc. SPIE 4387, Optical Pattern Recognition XII, pp. 95-102, 2001.
Summed-area tables are due to Franklin Crow, who introduced them for texture antialiasing. Crow, Franklin C. “Summed-Area Tables for Texture Mapping.” Computer Graphics (Proceedings of SIGGRAPH ’84) 18(3), pp. 207-212, 1984.
The single-pass scan behind the fast summed-area-table build is Merrill
and Garland’s decoupled look-back, the algorithm inside
cub::DeviceScan (Chapter 13). Merrill, D. and Garland, M.
“Single-pass Parallel Prefix Scan with Decoupled Look-back.” NVIDIA
Technical Report NVR-2016-002, 2016.
The instruction-level-parallelism argument of Section 15.5—more independent work per thread hides latency better than more threads do—is Vasily Volkov’s “Better Performance at Lower Occupancy,” GPU Technology Conference, 2010.