Most template-matching applications perform many correlation computations with the same template, at different offsets of the input image. In that case, the template statistics (SumT and fDenomExp) can be precomputed, and the template data can be moved to special memory or otherwise pre-massaged. For CUDA, the obvious place to put the template data is in constant memory, so each template pixel can be broadcast to the threads computing correlation values for different image locations.
Figure 15-4. Template in __constant__ memory
The CopyToTemplate() function given in Listing 15-3 pulls a rectangular
area of pixels out of the input image, computes the statistics, and
copies the data and statistics to __constant__ memory:
const int maxTemplatePixels = 3072;
__constant__ int g_xOffset[maxTemplatePixels];
__constant__ int g_yOffset[maxTemplatePixels];
__constant__ unsigned char g_Tpix[maxTemplatePixels];
__constant__ float g_cPixels, g_SumT, g_fDenomExp;
unsigned int g_cpuSumT, g_cpuSumTSq;
cudaError_tCopyToTemplate( unsigned char *img, size_t imgPitch, int xTemplate, int yTemplate, int wTemplate, int hTemplate, int OffsetX, int OffsetY){ cudaError_t status_cudart; unsigned char pixels[maxTemplatePixels]; int inx = 0; int SumT = 0; int SumTSq = 0; int cPixels = wTemplate*hTemplate; size_t sizeOffsets = cPixels*sizeof(int); float fSumT, fDenomExp, fcPixels; cuda(Memcpy2D( pixels, wTemplate, img+yTemplate*imgPitch+xTemplate, imgPitch, wTemplate, hTemplate, cudaMemcpyDeviceToHost ) ); cuda(MemcpyToSymbol( g_Tpix, pixels, cPixels ) ); for ( int i = OffsetY; i < OffsetY+hTemplate; i++ ) { for ( int j = OffsetX; j < OffsetX+wTemplate; j++) { SumT += pixels[inx]; SumTSq += pixels[inx]*pixels[inx]; poffsetx[inx] = j; poffsety[inx] = i; inx += 1; } } g_cpuSumT = SumT; g_cpuSumTSq = SumTSq; cuda(MemcpyToSymbol(g_xOffset, poffsetx, sizeOffsets) ); cuda(MemcpyToSymbol(g_yOffset, poffsety, sizeOffsets) ); fSumT = (float) SumT; cuda(MemcpyToSymbol(g_SumT, &fSumT, sizeof(float)) ); fDenomExp = float( (double)cPixels*SumTSq - (double) SumT*SumT); cuda(MemcpyToSymbol(g_fDenomExp, &fDenomExp, sizeof(float)) ); fcPixels = (float) cPixels; cuda(MemcpyToSymbol(g_cPixels, &fcPixels, sizeof(float)) );Error_cudart: return status_cudart;}The corrTexConstant() kernel given in Listing 15-4 then can read the
template values from g_Tpix[], which resides in constant memory.
__global__ void corrTemplate2D_kernel( float *pCorr, size_t CorrPitch, float cPixels, float fDenomExp, float xUL, float yUL, int w, int h, int xOffset, int yOffset, int wTemplate, int hTemplate ){ size_t row = blockIdx.y*blockDim.y + threadIdx.y; size_t col = blockIdx.x*blockDim.x + threadIdx.x; // adjust pointers 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 SumISq = 0; int SumIT = 0; int inx = 0; for ( int j = 0; j < hTemplate; j++ ) { for ( int i = 0; i < wTemplate; i++ ) { unsigned char I = tex2D( texImage, (float) col+xUL+xOffset+i, (float) row+yUL+yOffset+j ); unsigned char T = g_Tpix[inx++]; SumI += I; SumISq += I*I; SumIT += I*T; } } pCorr[col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp );}corrTexConstant() kernel. (source on GitHub)corrTexConstant() is even simpler and shorter than corrTexTex(),
since it does not have to compute the template statistics.