/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 dict_hw/src/normxcorr_hw.cu.h

  • Committer: Suren A. Chilingaryan
  • Date: 2010-08-08 20:10:46 UTC
  • Revision ID: csa@dside.dyndns.org-20100808201046-ad223bdx3ahc9j3a
Fixes, support for batched FFT transforms

Show diffs side-by-side

added added

removed removed

Lines of Context:
dict_hw/src/normxcorr_hw.cu.h
50
50
    
51
51
    int ncp_cache;
52
52
    int cache_memory;
 
53
    
 
54
#ifdef CUDA_BATCH_FFT
 
55
    int dims[2] = { ps->fft_real_size, ps->fft_real_size };
53
56
 
 
57
    cufft_err = cufftPlanMany(&ps->cufft_r2c_plan, 2, dims, NULL, 1, 0, NULL, 1, 0, CUFFT_R2C, CP_BLOCK);
 
58
#else /* CUDA_BATCH_FFT */
54
59
    cufft_err = cufftPlan2d(&ps->cufft_r2c_plan, ps->fft_real_size, ps->fft_real_size, CUFFT_R2C);
 
60
#endif /* CUDA_BATCH_FFT */
55
61
    if (cufft_err) {
56
62
        reportError("Problem initializing c2r plan, cufft code: %i", cufft_err);
57
63
        return DICT_ERROR_CUFFT;
58
64
    }   
59
65
    
 
66
#ifdef CUDA_BATCH_FFT
 
67
    cufft_err = cufftPlanMany(&ps->cufft_c2r_plan, 2, dims, NULL, 1, 0, NULL, 1, 0, CUFFT_C2R, CP_BLOCK);
 
68
#else /* CUDA_BATCH_FFT */
60
69
    cufft_err = cufftPlan2d(&ps->cufft_c2r_plan, ps->fft_real_size, ps->fft_real_size, CUFFT_C2R);
 
70
#endif /* CUDA_BATCH_FFT */
61
71
    if (cufft_err) {
62
72
        reportError("Problem initializing r2c plan, cufft code: %i", cufft_err);
63
73
        cufftDestroy(ps->cufft_r2c_plan);
493
503
            lsum_temp + (2 * lsum_step), lsum_temp + (3 * lsum_step),
494
504
            lsum_temp, lsum_temp + lsum_step);
495
505
 
 
506
#ifndef CUDA_BATCH_FFT
496
507
        cufftExecR2C(ps->cufft_r2c_plan, cuda_base_buffer + i * alloc_size, cache + i * alloc_size);
 
508
#endif /* ! CUDA_BATCH_FFT */
 
509
 
497
510
    }
498
 
    
 
511
 
 
512
#ifdef CUDA_BATCH_FFT
 
513
    cufftExecR2C(ps->cufft_r2c_plan, cuda_base_buffer, cache);
 
514
#endif /* CUDA_BATCH_FFT */
499
515
 
500
516
    return 0;
501
517
}
667
683
 
668
684
    }
669
685
 
 
686
 
670
687
    if (!image_mode) {
671
688
        cudaMemcpy3DParms copy_params = { 0 };
672
689
 
682
699
        cudaMemcpy3D(&copy_params);
683
700
    }
684
701
 
 
702
 
685
703
    return 0;
686
704
}
687
705
 
817
835
        // Performing FFT's
818
836
    cufftComplex *cuda_fft_buffer = ((cufftComplex*)ps->cuda_temp_buffer) + alloc_size;
819
837
 
 
838
#ifdef CUDA_BATCH_FFT
 
839
    cufftExecR2C(ps->cufft_r2c_plan, cuda_data_buffer, cuda_fft_buffer);
 
840
#else /* CUDA_BATCH_FFT */
820
841
    for (int i = 0;i < ncp;i++) {
821
842
        if (banlist[i]) continue;
822
843
        cufftExecR2C(ps->cufft_r2c_plan, cuda_data_buffer + i * alloc_size, cuda_fft_buffer + i * alloc_size);
823
844
    }
 
845
#endif /* CUDA_BATCH_FFT */
824
846
 
 
847
#ifdef CUDA_BATCH_FFT
 
848
        // No interleave in current version
 
849
    int complex_blocks = calc_blocks(fft_real_size * fft_real_size, SIDE_BLOCK_SIZE);
 
850
#else /* CUDA_BATCH_FFT */
825
851
    int complex_blocks = calc_blocks(fft_real_size * (fft_real_size / 2 + 1), SIDE_BLOCK_SIZE);
 
852
#endif /* CUDA_BATCH_FFT */
 
853
 
826
854
    dim3 complex_grid_dim(complex_blocks, cp_blocks, 1);
827
855
    vecMul<<<complex_grid_dim,block_side_cp>>>(cuda_fft_buffer, ps->cuda_fft_cache + cache_icp * alloc_size, alloc_size, fft_real_size/2+1);
828
856
 
829
857
        // First in-place transform for some reason is failing, therefore we
830
858
        // have one alloc_size spacing between starts (see cuda_fft_buffer set above)
831
859
    cufftReal *cuda_result_buffer = (cufftReal*)ps->cuda_temp_buffer;
 
860
 
 
861
#ifdef CUDA_BATCH_FFT
 
862
    cufftExecC2R(ps->cufft_c2r_plan, cuda_fft_buffer,  cuda_result_buffer);
 
863
#else /* CUDA_BATCH_FFT */
832
864
    for (int i = 0;i < ncp;i++) {
833
865
        if (banlist[i]) continue;
834
866
        cufftExecC2R(ps->cufft_c2r_plan, cuda_fft_buffer + i * alloc_size,  cuda_result_buffer + i * alloc_size);
835
867
    }
 
868
#endif /* CUDA_BATCH_FFT */
 
869
 
836
870
 
837
871
    return 0;
838
872
}