/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.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.cu
315
315
        banlist[i] = 0;
316
316
    }
317
317
 
318
 
    dim3 input_block_dim(size, 1, 1);
319
 
    dim3 input_grid_dim(size, 1, 1);
320
318
 
321
319
    cufftReal *cuda_base_buffer = ps->cuda_base_buffer;
322
320
    cufftComplex *cache = ps->cuda_fft_cache +  icp * alloc_size;
323
321
    float *lsum_cache = ps->cuda_lsum_cache + icp * alloc_size;
324
322
    float *denom_cache = ps->cuda_denom_cache + icp * alloc_size;
325
323
 
 
324
    int blocks = calc_blocks(size, BLOCK_SIZE_1D);
 
325
    int base_blocks = blocks * blocks * BLOCK_SIZE_1D;
 
326
    
 
327
    char blocks_power;
 
328
    if ((blocks&(blocks-1))) {
 
329
        blocks_power = -1;
 
330
    } else {
 
331
        blocks_power = debruijn[((uint32_t)blocks * 0x077CB531) >> 27];
 
332
    }
 
333
 
 
334
    int lsum_size = ps->lsum_size;
 
335
    int lsum_alloc = ps->lsum_alloc_size;
 
336
 
326
337
    for (int i = 0;i < ncp;i++) {
327
338
        if (banlist[i]) continue;
328
339
        
329
 
        vecPackBase<<<input_grid_dim, input_block_dim>>>(
330
 
            cuda_input_buffer + i * side_alloc2, side_alloc, 
331
 
            cuda_base_buffer, fft_size, 
332
 
            lsum_temp, lsum_temp + lsum_step, ps->lsum_alloc_size, ps->lsum_size
333
 
        );
 
340
        if (blocks_power < 0) {
 
341
            vecBasePack<<<base_blocks, BLOCK_SIZE_1D>>>(
 
342
                cuda_input_buffer + i * side_alloc2, side_alloc, 
 
343
                cuda_base_buffer, fft_size, 
 
344
                lsum_temp + lsum_size * (lsum_alloc + 1), 
 
345
                lsum_temp + lsum_step + lsum_size * (lsum_alloc + 1), 
 
346
                lsum_alloc,
 
347
                size, blocks
 
348
            );
 
349
        } else {
 
350
            vecBasePackFast<<<base_blocks, BLOCK_SIZE_1D>>>(
 
351
                cuda_input_buffer + i * side_alloc2, side_alloc, 
 
352
                cuda_base_buffer, fft_size, 
 
353
                lsum_temp + lsum_size * (lsum_alloc + 1), 
 
354
                lsum_temp + lsum_step + lsum_size * (lsum_alloc + 1), 
 
355
                lsum_alloc,
 
356
                size, blocks_power
 
357
            );
 
358
        }
334
359
 
335
360
        // In general we should expect non-zero denominals, therefore the Nonzero array is not computed
336
361
        local_sum(ps,