/normxcorr/trunk

To get this branch, use:
bzr branch http://suren.me/webbzr/normxcorr/trunk

« back to all changes in this revision

Viewing changes to cuda/normxcorr_hw_kernel.cu

  • Committer: Suren A. Chilingaryan
  • Date: 2009-12-10 03:16:21 UTC
  • Revision ID: csa@dside.dyndns.org-20091210031621-2a15m2tdumdz3s39
Block computational kernels

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw_kernel.cu
1
 
static __global__ void vecMul(cuComplex *a, cuComplex *b, int size) {
 
1
static __global__ void vecMul(cuComplex *a, cuComplex *b, int pitch, int size) {
2
2
    float tmp;
3
 
    
4
 
    int i = threadIdx.x + blockIdx.x*size;
 
3
 
 
4
    int point = blockIdx.y * blockDim.y + threadIdx.y;
 
5
    int i = threadIdx.x + blockIdx.x * blockDim.x + point*pitch;
5
6
    
6
7
    tmp = a[i].x * b[i].x - a[i].y * b[i].y;
7
8
    a[i].y = a[i].x * b[i].y + a[i].y * b[i].x; 
8
9
    a[i].x = tmp;
9
10
}
10
11
 
11
 
 
12
 
static __global__ void vecPack(uint8_t *b, int bsize, cufftReal *a, int asize, int size) {
13
 
//    int i = threadIdx.x + blockIdx.x*bsize;
14
 
//    int i = bsize - threadIdx.x - 1 + (bsize - blockIdx.x - 1)*bsize;
15
 
 
16
 
        // Invalid memory access (possibly)
17
 
//    int i = (bsize - blockIdx.x)*bsize - threadIdx.x - 1;
18
 
    int i = size - threadIdx.x - 1 + (size - blockIdx.x - 1)*bsize;
19
 
    a[threadIdx.x + asize*blockIdx.x] = b[i];
20
 
}
 
12
static __global__ void vecPack(uint8_t *b, int bpitch, int bsize, cufftReal *a, int apitch, int asize, int size, int blocks_size) {
 
13
        // Includes rotation on 180 grad
 
14
    
 
15
    int point = blockIdx.y * blockDim.y + threadIdx.y;
 
16
 
 
17
    //__fdiv_rz(pos / bsize)
 
18
    int y = __float2int_rz(__fdividef(blockIdx.x, blocks_size));
 
19
    int x = (blockIdx.x - y * blocks_size) * blockDim.x + threadIdx.x ;
 
20
    
 
21
/*
 
22
    int pos = blockIdx.x * blockDim.x + threadIdx.x;
 
23
    int y = pos / (blocks_size * blockDim.x);
 
24
    int x = pos - (y * blocks_size * blockDim.x);
 
25
*/
 
26
 
 
27
    if ((x < size)&&(y < size)) {
 
28
        int i = (size - y - 1)*bsize + size - x - 1;
 
29
        a[point * apitch + y * asize + x] = b[point * bpitch + i];
 
30
    }
 
31
 
 
32
}
 
33
 
 
34
static __global__ void vecPackFast(uint8_t *b, int bpitch, int bsize, cufftReal *a, int apitch, int asize, int size, int blocks_shift) {
 
35
    __shared__ float data[CP_BLOCK_SIZE][SIDE_BLOCK_SIZE + 1];
 
36
 
 
37
    int point = blockIdx.y * blockDim.y + threadIdx.y;
 
38
 
 
39
    int y = blockIdx.x>>blocks_shift;
 
40
    int bx = (blockIdx.x - (y<<blocks_shift)) * blockDim.x ;
 
41
    int x = bx + threadIdx.x;
 
42
 
 
43
//    int x = (blockIdx.x - (y<<blocks_shift)) * blockDim.x + threadIdx.x ;
 
44
 
 
45
        // threadIdx.x depends only on x
 
46
    data[threadIdx.y][threadIdx.x] = b[point * bpitch + y * bsize + x];
 
47
    
 
48
    __syncthreads();
 
49
 
 
50
    int pos  = size - bx - blockDim.x + threadIdx.x;
 
51
 
 
52
    if ((pos>=0)&&(y < size)) {
 
53
        a[point * apitch + (size - y - 1) * asize + pos] = 
 
54
            data[threadIdx.y][blockDim.x - threadIdx.x - 1];
 
55
    }
 
56
 
 
57
//    if ((x < size)&&(y < size)) {
 
58
//        int i = (size - y - 1)*bsize + size - x - 1;
 
59
//      a[point * apitch + y * asize + x] = b[point * bpitch + i];
 
60
//    }
 
61
 
 
62
}
 
63
 
21
64
 
22
65
static __global__ void vecPackBase(
23
66
    uint8_t *b, int bsize,
87
130
    cufftReal *corr, float corr_scale, 
88
131
    float *lsum, float *lsum_scale_ptr, float lsum_mult,
89
132
    float *denom, float *denom_scale_ptr,
90
 
    int size
 
133
    int pitch, int size
91
134
) {
92
 
    int pos = threadIdx.x + blockIdx.x*size;
93
 
 
94
 
    float lsum_scale = (*lsum_scale_ptr) * lsum_mult;
95
 
    float denom_scale = (*denom_scale_ptr);
 
135
//    int pos = threadIdx.x + blockIdx.x*size;
 
136
 
 
137
    int point = blockIdx.y * blockDim.y + threadIdx.y;
 
138
    int pos = threadIdx.x + blockIdx.x * blockDim.x + point * pitch;
 
139
 
 
140
 
 
141
    float lsum_scale = lsum_scale_ptr[point] * lsum_mult;
 
142
    float denom_scale = denom_scale_ptr[point];
96
143
 
97
144
    if (denom[pos]&&denom_scale) {
98
145
        res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);