For rectangles of correlation values such as the ones computed by our sample program, the CUDA kernel exhibits a tremendous amount of reuse of the image data as the template matches are swept across the image. So far, our code has relied on the texture caches to service these redundant reads without going to external memory. For smaller templates, however, shared memory can be used to further increase performance by making the image data available with lower latency.
The kernels of Listings 15-1 and 15-4 implicitly divided the input
image into tiles that were the same size as the threadblocks. For our
shared memory implementation, shown in Listing 15-5, we’ll use the
height of the threadblock (blockDim.y), but specify an explicit tile
width of wTile. In our sample program, wTile is 32. Figure 15-5 shows
how the kernel “overfetches” a rectangle of
wTemplate×hTemplate pixels outside the tile;
boundary conditions are handled by the texture addressing mode. Once the
shared memory has been populated with image data, the kernel does
__syncthreads() and computes and writes out the tile’s
correlation coefficients.
Figure 15-5. Image in shared memory
__global__ void corrShared_kernel( float *pCorr, size_t CorrPitch, int wTile, int wTemplate, int hTemplate, float xOffset, float yOffset, float cPixels, float fDenomExp, int SharedPitch, float xUL, float yUL, int w, int h ){ int uTile = blockIdx.x*wTile; int vTile = blockIdx.y*blockDim.y; int v = vTile + threadIdx.y; float *pOut = (float *) (((char *) pCorr)+v*CorrPitch); for ( int row = threadIdx.y; row < blockDim.y+hTemplate; row += blockDim.y ) { int SharedIdx = row * SharedPitch; for ( int col = threadIdx.x; col < wTile+wTemplate; col += blockDim.x ) { LocalBlock[SharedIdx+col] = tex2D( texImage, (float) (uTile+col+xUL+xOffset), (float) (vTile+row+yUL+yOffset) ); } } __syncthreads(); for ( int col = threadIdx.x; col < wTile; col += blockDim.x ) { int SumI = 0; int SumISq = 0; int SumIT = 0; int idx = 0; int SharedIdx = threadIdx.y * SharedPitch + col; for ( int j = 0; j < hTemplate; j++ ) { for ( int i = 0; i < wTemplate; i++) { unsigned char I = LocalBlock[SharedIdx+i]; unsigned char T = g_Tpix[idx++]; SumI += I; SumISq += I*I; SumIT += I*T; } SharedIdx += SharedPitch; } if ( uTile+col < w && v < h ) { pOut[uTile+col] = CorrelationValue( SumI, SumISq, SumIT, g_SumT, cPixels, fDenomExp ); } } __syncthreads();}corrShared_kernel() (source on GitHub)To ensure that shared memory references have the same characteristics from one row to the next, the amount of shared memory per row is padded to the next multiple of 64:
sharedPitch = ~63&(((wTile+wTemplate)+63));
The total amount of shared memory needed per block is then the pitch multiplied by the number of rows (block height plus template height):
sharedMem = sharedPitch*(threads.y+hTemplate);
The host code to launch corrShared_kernel(), shown in Listing 15-6,
detects whether the kernel launch will require more shared memory than
is available. If that is the case, it calls corrTexTex(), which will
work for any template size.
voidcorrShared( 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 ){ int device; cudaDeviceProp props; cudaError_t status; cuda(GetDevice( &device ) ); cuda(GetDeviceProperties( &props, device ) ); if ( sharedMem > props.sharedMemPerBlock ) { dim3 threads88(8, 8, 1); dim3 blocks88; blocks88.x = INTCEIL(w,8); blocks88.y = INTCEIL(h,8); blocks88.z = 1; return corrTexTex2D( dCorr, CorrPitch, wTile, wTemplate, hTemplate, cPixels, fDenomExp, sharedPitch, xOffset, yOffset, xTemplate, yTemplate, xUL, yUL, w, h, threads88, blocks88, sharedMem ); } corrShared_kernel<<<blocks, threads, sharedMem>>>( dCorr, CorrPitch, wTile, wTemplate, hTemplate, (float) xOffset, (float) yOffset, cPixels, fDenomExp, sharedPitch, (float) xUL, (float) yUL, w, h );Error: return;}corrShared() (host code) (source on GitHub)