/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-02 05:08:22 UTC
  • Revision ID: csa@dside.dyndns.org-20091202050822-n6ouznm1zp2n2i5l
Instead of transfer compute local sums and denormals on board

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw_kernel.cu
17
17
    a[threadIdx.x + asize*blockIdx.x] = b[i];
18
18
}
19
19
 
20
 
 
 
20
static __global__ void vecPackBase(
 
21
    uint8_t *b, int bsize,
 
22
    cufftReal *a, int asize, 
 
23
    float *c, float *c2, int csize, int coffset) {
 
24
    
 
25
    float v;
 
26
    int i = threadIdx.x + (blockIdx.x+coffset)*csize + coffset;
 
27
    
 
28
    v = b[threadIdx.x + blockIdx.x*bsize];
 
29
 
 
30
    a[threadIdx.x + asize*blockIdx.x] = v;
 
31
    
 
32
    c[i] = v;
 
33
    c2[i] = v * v;
 
34
}
 
35
 
 
36
 
 
37
/*
21
38
static __global__ void vecCompute(
22
39
    uint16_t *items, float *res,
23
40
    cufftReal *corr, float corr_scale, 
24
41
    float *lsum, float lsum_scale,
25
42
    float *denom, float denom_scale
26
43
) {
27
 
    int pos = items[threadIdx.x + blockIdx.x*BLOCK_SIZE];//correct (row/column)?
 
44
    int pos = items[threadIdx.x + blockIdx.x*BLOCK_SIZE_1D];//correct (row/column)?
28
45
    res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);
29
46
}
 
47
*/
 
48
 
 
49
static __global__ void vecCompute(
 
50
    float *res,
 
51
    cufftReal *corr, float corr_scale, 
 
52
    float *lsum, float lsum_scale,
 
53
    float *denom, float denom_scale,
 
54
    int size
 
55
) {
 
56
    int pos = threadIdx.x + blockIdx.x*size;
 
57
 
 
58
    if (denom[pos]) {
 
59
        res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);
 
60
    }
 
61
}
30
62