/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-01-15 13:50:29 UTC
  • Revision ID: csa@dside.dyndns.org-20090115135029-wleapicg9a4593tp
Initial import

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw_kernel.cu
 
1
#define DOUBLE(x) ((double)(x))
 
2
 
 
3
static __global__ void vecMul(cuComplex *a, cuComplex *b, int size) {
 
4
    float tmp;
 
5
    
 
6
    int i = threadIdx.x + blockIdx.x*size;
 
7
    
 
8
    tmp = a[i].x * b[i].x - a[i].y * b[i].y;
 
9
    a[i].y = a[i].x * b[i].y + a[i].y * b[i].x; 
 
10
    a[i].x = tmp;
 
11
}
 
12
 
 
13
 
 
14
static __global__ void vecPack(cufftReal *a, int asize, uint8_t *b, int bsize) {
 
15
    int i = threadIdx.x + blockIdx.x*bsize;
 
16
    
 
17
    a[threadIdx.x + asize*blockIdx.x] = b[i];
 
18
}
 
19
 
 
20
 
 
21
static __global__ void vecCompute(
 
22
    uint16_t *items, float *res,
 
23
    cufftReal *corr, float corr_scale, 
 
24
    float *lsum, float lsum_scale,
 
25
    float *denom, float denom_scale
 
26
) {
 
27
    int pos = items[threadIdx.x + blockIdx.x*BLOCK_SIZE];//correct (row/column)?
 
28
    res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);
 
29
}
 
30