15.2 简单的纹理实现

我们利用纹理单元读取图像和模板的像素值来实现第一个归一化互相关。虽然它并不是一个最优实现，甚至没有预先计算模板的统计量。但这种方法简单易懂，并且是更进一步优化（当然也更复杂）的一个良好的基础。

代码清单15-1给出了执行此计算的内核程序，它计算出了5个数据的加和信息，然后利用前面给出的函数CorrelationValue()将浮点型的相关系数写入到输出数组中。但要注意的是，计算fDenomExp的表达式可能会发出警告，这是因为SM1.3架构之前不支持双精度浮点型的计算，但只要模板像素数量在允许的范围之内，内核仍可以正常工作。

(xUL, yUL) 表示图像左上角，w和h表示搜索窗口的宽和高，也就是表示系数的输出数组。如果该模板是在纹理内存中，则其纹理图像的左上角可由 (xTemplate, yTemplate) 表示。

最后，偏移量(xOffset, yOffset)描述了模板是如何与待比较的图像进行叠加的。当读取图像像素时，这个偏移量被添加到了左上角为(xUL, yUL)的搜索矩形框所在的坐标中。

从模板所在位置向周边行进时，我们看到其相关函数递减，这对我们很有启发。示例程序normalizedCrossCorrelation.cu记录下的模

板附近的相关系数如下所示：

模板周围：

图15-1中包含了很多硬币，右下角的被高亮显示的 $52 \times 52$ 大小的镍币图像便是默认模板。默认的程序可以有选择的输出一个PGM文件，并转换相关像素值的范围为 $0 \sim 255$ 。图15-1中的高亮部分的相关图像如图15-2所示。其他镍币所在的位置的亮度较强，表明匹配的程度较好，而那些硬币的响应更弱。

图15-1 Coins.pgm（高亮部分为默认的模板）

图15-2 默认模板的相关图像

代码清单15-1 corrTexTex2D_kernel

global void
corrTexTex2D_kernel()
float *pCorr, size_t CorrPitch,
float cPixels,
int xOffset, int yOffset,
int xTemplate, int yTemplate,
int wTemplate, int hTemplate,
float xUL, float yUL, int w, int h)
{
size_t row = blockIdx.y*blockDim.y + threadIdx.y;
size_t col = blockIdx.x*blockDim.x + threadIdx.x;
// adjust pCorr to point to row
pCorr = (float *) ((char *) pCorr+row*CorrPitch);
// No syncthreads in this kernel, so we can early-out
// without worrying about the effects of divergence.
if (col >= w || row >= h)
return;
int SumI = 0;
int SumT = 0;
int SumISq = 0;
int SumTSq = 0;
int SumIT = 0;
for (int y = 0; y < hTemplate; y++) {
    for (int x = 0; x < wTemplate; x++) {
        unsigned char I = tex2D(texImage,
            (float) col+xUL+xOffset+x, (float) row+yUL+yOffset+y);
        unsigned char T = tex2D(texTemplate,
            (float) xTemplate+x, (float) yTemplate+y);
        SumI += I;
        SumT += T;
        SumISq += I*T;
        SumTSq += T*T;
        SumIT += I*T;
    }
    float fDenomExp = (float) ((double) cPixels*SumTSq -
            (double) SumT*SumT);
pCorr[col] = CorrelationValue(
SumI, SumISq, SumIT, SumT, cPixels, fDenomExp);
}

代码清单15-2给出了调用corrTexTex2D_kernel()的主机端代码。它是用来测试源文件normalizedCrossCorrelation.cu的性能表现的，所以有很多参数。此主机端的函数此时只是用来传递内核函数所需要

的参数，但它在后面的实现中，还会检查它所发现的设备的属性，并启动不同的内核。对于一个大小适中的图像来说，做这样的检查的成本相比于GPU运行时间是微不足道的。

代码清单15-2 corrTexTex2D()（主机端代码）

void  
corrTexTex2D( float \*dCorr, int CorrPitch, int wTile, int wTemplate, int hTemplate, float cPixels, float fDenomExp, int sharedPitch, int xOffset, int yOffset, int xTemplate, int yTemplate, int xUL, int yUL, int w, int h, dim3 threads, dim3 blocks, int sharedMem)  
{ corrTexTex2D_kernel<<blocks, threads>>>( dCorr, CorrPitch, cPixels, xOffset, yOffset, xTemplate+xOffset, yTemplate+yOffset, wTemplate, hTemplate, (float) xUL, (float) yUL, w, h); }

在搜索过程中，如果应用程序选择不同的模板和不同的图像时，均采用纹理实现是非常适宜的一例如，在模板和图像比对时，需要变换模板的数据时。但是在大多数应用中，模板仅被选择一次，并且用来和图像内很多不同偏移的内容作计算和比较。本章的剩余部分将会针对这一情况继续进行优化。