/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 04:26:05 UTC
  • Revision ID: csa@dside.dyndns.org-20091210042605-6wmxmhj3mj6happs
Eleminate last kernel with NxN geometr

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw_kernel.cu
62
62
}
63
63
 
64
64
 
65
 
static __global__ void vecPackBase(
66
 
    uint8_t *b, int bsize,
67
 
    cufftReal *a, int asize, 
68
 
    float *c, float *c2, int csize, int coffset) {
69
 
    
70
 
    float v;
71
 
    int i = threadIdx.x + (blockIdx.x+coffset)*csize + coffset;
72
 
    
73
 
    v = b[threadIdx.x + blockIdx.x*bsize];
74
 
 
75
 
    a[threadIdx.x + asize*blockIdx.x] = v;
76
 
    
77
 
    c[i] = v;
78
 
    c2[i] = v * v;
 
65
static __global__ void vecBasePack(
 
66
    uint8_t *b, int bsize,
 
67
    cufftReal *a, int asize, 
 
68
    float *c, float *c2, int csize,
 
69
    int size, int blocks_size) {
 
70
 
 
71
    int y = __float2int_rz(__fdividef(blockIdx.x, blocks_size));
 
72
    int x = (blockIdx.x - y * blocks_size) * blockDim.x + threadIdx.x ;
 
73
 
 
74
    if ((x<size)&&(y<size)) {
 
75
        float v = b[x + y*bsize];
 
76
 
 
77
        a[x + y*asize] = v;
 
78
        
 
79
        int i = x + y*csize;
 
80
        c[i] = v;
 
81
        c2[i] = v * v;
 
82
    }
 
83
}
 
84
 
 
85
static __global__ void vecBasePackFast(
 
86
    uint8_t *b, int bsize,
 
87
    cufftReal *a, int asize, 
 
88
    float *c, float *c2, int csize,
 
89
    int size, int blocks_shift) {
 
90
 
 
91
    int y = blockIdx.x>>blocks_shift;
 
92
    int x = (blockIdx.x - (y<<blocks_shift)) * blockDim.x + threadIdx.x ;
 
93
    
 
94
    if ((x<size)&&(y<size)) {
 
95
        float v = b[x + y*bsize];
 
96
        a[x + y*asize] = v;
 
97
 
 
98
        int i = x + y*csize;
 
99
        c[i] = v;
 
100
        c2[i] = v * v;
 
101
    }
79
102
}
80
103
 
81
104