/normxcorr/trunk : revision 9

To get this branch, use:

bzr branch
http://suren.me/webbzr/normxcorr/trunk

Avoid bank conflicts in local sum kernels

added added

removed removed

__global__ void transpose1(float *dst, float *src, int width, int height, int step) {

__shared__ float data[BLOCK_SIZE_2D][BLOCK_SIZE_2D];

// + 1 is to resolve bank conflict: Sequential threads within the same

// warp are accessing locations that are exactly BLOCK_SIZE apart, which

// will cause bank conflicts if BLOCK_SIZE is a multiple of 16

__shared__ float data[BLOCK_SIZE_2D][BLOCK_SIZE_2D + 1];

int dx = blockIdx.x * blockDim.x;

int dy = blockIdx.y * blockDim.y;

__global__ void transpose2(float *lsum, float *denom, float *src1, float *src2, int width, int height, int step) {

__shared__ float data1[BLOCK_SIZE_2D][BLOCK_SIZE_2D];

__shared__ float data2[BLOCK_SIZE_2D][BLOCK_SIZE_2D];

__shared__ float data1[BLOCK_SIZE_2D][BLOCK_SIZE_2D + 1];

__shared__ float data2[BLOCK_SIZE_2D][BLOCK_SIZE_2D + 1];

int dx = blockIdx.x * blockDim.x;

int dy = blockIdx.y * blockDim.y;