/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 16:04:14 UTC
  • Revision ID: csa@dside.dyndns.org-20091210160414-jxqhujlctfp0k8dv
Precompute if side and base blocks amount is power of 2

Show diffs side-by-side

added added

removed removed

Lines of Context:
cuda/normxcorr_hw.cu
292
292
    int blocks = calc_blocks(size, BLOCK_SIZE_1D);
293
293
    int base_blocks = blocks * blocks * BLOCK_SIZE_1D;
294
294
    
295
 
    char blocks_power;
296
 
    if ((blocks&(blocks-1))) {
297
 
        blocks_power = -1;
298
 
    } else {
299
 
        blocks_power = debruijn[((uint32_t)blocks * 0x077CB531) >> 27];
300
 
    }
301
 
 
302
295
    int lsum_size = ps->lsum_size;
303
296
    int lsum_alloc = ps->lsum_alloc_size;
304
297
 
354
347
      if (i > 0) {
355
348
        int j = i - 1;
356
349
        
357
 
        if (blocks_power < 0) {
 
350
        if (ps->base_blocks_power < 0) {
358
351
            vecBasePack<<<base_blocks, BLOCK_SIZE_1D, 0, stream[j%2]>>>(
359
352
                cuda_input_buffer + j * side_alloc2, side_alloc, 
360
353
                cuda_base_buffer + j*alloc_size, fft_real_size, 
370
363
                lsum_temp + lsum_size * (lsum_alloc + 1), 
371
364
                lsum_temp + lsum_step + lsum_size * (lsum_alloc + 1), 
372
365
                lsum_alloc,
373
 
                size, blocks_power
 
366
                size, ps->base_blocks_power
374
367
            );
375
368
        }
376
369
 
508
501
        // Packing input data for FFT
509
502
    dim3 input_grid_dim(input_blocks, cp_blocks, 1);
510
503
 
511
 
    char side_blocks_power;
512
 
    if ((side_blocks&(side_blocks-1))) {
513
 
        side_blocks_power = -1;
514
 
    } else {
515
 
        side_blocks_power = debruijn[((uint32_t)side_blocks * 0x077CB531) >> 27];
516
 
    }
517
 
//    printf("power %i\n", side_blocks_power);
518
 
 
519
 
//    cudaMemset((void*)ps->cuda_data_buffer, 0, CP_BLOCK * ps->fft_alloc_size * sizeof(cufftReal));
520
 
    if (side_blocks_power < 0) {
 
504
    if (ps->side_blocks_power < 0) {
521
505
        vecPack<<<input_grid_dim, block_side_cp>>>(
522
506
            cuda_input_buffer, side_alloc2, side_alloc, 
523
507
            cuda_data_buffer, alloc_size, fft_real_size, 
527
511
        vecPackFast<<<input_grid_dim, block_side_cp>>>(
528
512
            cuda_input_buffer, side_alloc2, side_alloc, 
529
513
            cuda_data_buffer, alloc_size, fft_real_size, 
530
 
            size, side_blocks_power
 
514
            size, ps->side_blocks_power
531
515
        );
532
516
    }
533
517
 
795
779
    int width, height;
796
780
    int size, size2;
797
781
    int base_size, base_size2;
 
782
    int base_blocks, side_blocks;
798
783
 
799
784
    if (!nrhs) {
800
785
        reportMessage("Initializing normxcorr_hw instance");
1175
1160
                return;
1176
1161
            }
1177
1162
        }
 
1163
 
 
1164
        side_blocks = calc_blocks(2 * ps->corr_size  + 1, SIDE_BLOCK_SIZE);
 
1165
        if ((side_blocks&(side_blocks-1))) {
 
1166
            ps->side_blocks_power = -1;
 
1167
        } else {
 
1168
            ps->side_blocks_power = debruijn[((uint32_t)side_blocks * 0x077CB531) >> 27];
 
1169
        }
 
1170
 
 
1171
        base_blocks = calc_blocks(4 * ps->corr_size  + 1, BLOCK_SIZE_1D);
 
1172
        if ((base_blocks&(base_blocks-1))) {
 
1173
            ps->base_blocks_power = -1;
 
1174
        } else {
 
1175
            ps->base_blocks_power = debruijn[((uint32_t)base_blocks * 0x077CB531) >> 27];
 
1176
        }
1178
1177
     break;
1179
1178
     case ACTION_PREPARE:
1180
1179
        fftPrepare(ps);