/normxcorr/trunk

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

    __shared__ float data[BLOCK_SIZE_2D][BLOCK_SIZE_2D];

 

    int dx = blockIdx.x * blockDim.x;

    int dy = blockIdx.y * blockDim.y;

 

    int x = threadIdx.x + dx;

    int y = threadIdx.y + dy;

    

    data[threadIdx.x][threadIdx.y] = 

        src[x + step + y*width] - src[x + y*width];

    

    __syncthreads();

 

    x = threadIdx.x + dy;

    y = threadIdx.y + dx;

    

    dst[x + y*height] = data[threadIdx.y][threadIdx.x];

}

 

 

__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];

 

    int dx = blockIdx.x * blockDim.x;

    int dy = blockIdx.y * blockDim.y;

 

    int x = threadIdx.x + dx;

    int y = threadIdx.y + dy;

    

    data1[threadIdx.x][threadIdx.y] = 

        src1[x + step + y*width] - src1[x + y*width];

 

    data2[threadIdx.x][threadIdx.y] = 

        src2[x + step + y*width] - src2[x + y*width];

    

    __syncthreads();

 

    x = threadIdx.x + dy;

    y = threadIdx.y + dx;

    

    float val1 = data1[threadIdx.y][threadIdx.x];

    float val2 = data2[threadIdx.y][threadIdx.x];

    float coef = step * step;

    

    if (x < height) {

        lsum[x + y*height] = val1;

        denom[x + y*height] = sqrtf(fmaxf((val2 - val1*val1/coef)/(coef - 1) ,0));

    }

}