/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 15:47:44 UTC
  • Revision ID: csa@dside.dyndns.org-20091210154744-min3x71y3tgrkvpu
Optimize FFT size

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw_kernel.cu
148
148
    res2[img_idx] = sqrtf(fmaxf(((float)sum2) / cnt - mean*mean,0));
149
149
}
150
150
 
 
151
/*
 
152
Use real size everthere
151
153
static __global__ void vecCompute(
152
154
    float *res,
153
155
    cufftReal *corr, float corr_scale, 
154
156
    float *lsum, float *lsum_scale_ptr, float lsum_mult,
155
157
    float *denom, float *denom_scale_ptr,
156
 
    int pitch, int size
 
158
    int pitch
157
159
) {
158
 
//    int pos = threadIdx.x + blockIdx.x*size;
159
 
 
160
160
    int point = blockIdx.y * blockDim.y + threadIdx.y;
161
161
    int pos = threadIdx.x + blockIdx.x * blockDim.x + point * pitch;
162
162
 
168
168
        res[pos] = (corr[pos] * corr_scale - lsum[pos]*lsum_scale) / (denom[pos] * denom_scale);
169
169
    }
170
170
}
 
171
*/
 
172
 
 
173
static __global__ void vecCompute(
 
174
    float *res, int row_pitch,
 
175
    cufftReal *corr, int corr_row_pitch, float corr_scale, 
 
176
    float *lsum, float *lsum_scale_ptr, float lsum_mult,
 
177
    float *denom, float *denom_scale_ptr,
 
178
    int pitch, int blocks_size
 
179
) {
 
180
    int point = blockIdx.y * blockDim.y + threadIdx.y;
 
181
 
 
182
    int y = __float2int_rz(__fdividef(blockIdx.x, blocks_size));
 
183
    int x = (blockIdx.x - y * blocks_size) * blockDim.x + threadIdx.x ;
 
184
 
 
185
    int i = x + y * corr_row_pitch + point * pitch;
 
186
    int pos = x + y * row_pitch + point * pitch;
 
187
 
 
188
    float lsum_scale = lsum_scale_ptr[point] * lsum_mult;
 
189
    float denom_scale = denom_scale_ptr[point];
 
190
 
 
191
    if (denom[pos]&&denom_scale) {
 
192
        res[pos] = (corr[i] * corr_scale - lsum[pos] * lsum_scale) / (denom[pos] * denom_scale);
 
193
    }
 
194
}
 
195
 
171
196
 
172
197
static __global__ void find_max1(float *buf1, int32_t *buf2, float *corr, int image_pitch, int row_pitch, int size) {
173
198
    int i;