/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-06 01:52:56 UTC
  • Revision ID: csa@dside.dyndns.org-20091206015256-evn0sne8d18ovm8o
A little more computations are moved to CUDA

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw_kernel.cu
11
11
}
12
12
 
13
13
 
14
 
static __global__ void vecPack(cufftReal *a, int asize, uint8_t *b, int bsize) {
15
 
    int i = threadIdx.x + blockIdx.x*bsize;
16
 
    
 
14
static __global__ void vecPack(cufftReal *a, int asize, uint8_t *b, int bsize, int size) {
 
15
//    int i = threadIdx.x + blockIdx.x*bsize;
 
16
//    int i = bsize - threadIdx.x - 1 + (bsize - blockIdx.x - 1)*bsize;
 
17
 
 
18
        // Invalid memory access (possibly)
 
19
//    int i = (bsize - blockIdx.x)*bsize - threadIdx.x - 1;
 
20
    int i = size - threadIdx.x - 1 + (size - blockIdx.x - 1)*bsize;
17
21
    a[threadIdx.x + asize*blockIdx.x] = b[i];
18
22
}
19
23
 
34
38
}
35
39
 
36
40
 
 
41
static __global__ void stat1(int *buf1, int *buf2, uint8_t *img, int image_pitch, int row_pitch, int size) {
 
42
    int i;
 
43
    int end = size * row_pitch;
 
44
 
 
45
    int side_idx =  blockIdx.x * blockDim.x + threadIdx.x;
 
46
    int img_idx = blockIdx.y * blockDim.y + threadIdx.y;
 
47
 
 
48
    int sum = 0;
 
49
    int sum2 = 0;
 
50
 
 
51
    uint8_t *vec = img + img_idx * image_pitch + side_idx;
 
52
 
 
53
    for (i = 0; i < end; i+=row_pitch) {
 
54
        int val = vec[i];
 
55
        sum += val;
 
56
        sum2 += val*val;
 
57
    }
 
58
 
 
59
    buf1[side_idx * CP_BLOCK + img_idx] = sum;
 
60
    buf2[side_idx * CP_BLOCK + img_idx] = sum2;
 
61
}
 
62
 
 
63
static __global__ void stat2(float *res1, float *res2, int *buf1, int *buf2, int size) {
 
64
    int i;
 
65
    int end = size * CP_BLOCK;
 
66
    int img_idx =  blockIdx.x * blockDim.x + threadIdx.x;
 
67
 
 
68
    int sum = 0;
 
69
    int sum2 = 0;
 
70
 
 
71
    int *vec1 = buf1 + img_idx;
 
72
    int *vec2 = buf2 + img_idx;
 
73
 
 
74
    for (i = 0; i < end; i+=CP_BLOCK) {
 
75
        sum += vec1[i];
 
76
        sum2 += vec2[i];
 
77
    }
 
78
 
 
79
    res1[img_idx] = sum;
 
80
 
 
81
    float cnt = size * size;
 
82
    float mean = ((float)sum) / cnt;
 
83
    
 
84
    res2[img_idx] = sqrtf(fmaxf(((float)sum2) / cnt - mean*mean,0));
 
85
}
 
86
 
 
87
 
 
88
 
 
89
 
 
90
 
37
91
/*
38
92
static __global__ void vecCompute(
39
93
    uint16_t *items, float *res,
46
100
}
47
101
*/
48
102
 
 
103
/*
49
104
static __global__ void vecCompute(
50
105
    float *res,
51
106
    cufftReal *corr, float corr_scale, 
59
114
        res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);
60
115
    }
61
116
}
62
 
 
 
117
*/
 
118
 
 
119
static __global__ void vecCompute(
 
120
    float *res,
 
121
    cufftReal *corr, float corr_scale, 
 
122
    float *lsum, float *lsum_scale_ptr, float lsum_mult,
 
123
    float *denom, float *denom_scale_ptr,
 
124
    int size
 
125
) {
 
126
    int pos = threadIdx.x + blockIdx.x*size;
 
127
 
 
128
    float lsum_scale = (*lsum_scale_ptr) * lsum_mult;
 
129
    float denom_scale = (*denom_scale_ptr);
 
130
 
 
131
    if (denom[pos]&&denom_scale) {
 
132
        res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);
 
133
    }
 
134
}