From 2832cfdfe581c8a18c695d6737d8ab92f7f6f846 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Fri, 27 Jun 2014 13:38:21 -0700 Subject: [PATCH] No longer use constant memory for image step. --- modules/cudastereo/src/cuda/stereocsbp.cu | 66 ++++++++++++++---------------- modules/cudastereo/src/cuda/stereocsbp.hpp | 7 ++-- modules/cudastereo/src/stereocsbp.cpp | 10 ++--- 3 files changed, 39 insertions(+), 44 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index fa17742..d0097f3 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -67,14 +67,12 @@ namespace cv { namespace cuda { namespace device __constant__ int cth; - __constant__ size_t cimg_step; __constant__ size_t cmsg_step; __constant__ size_t cdisp_step1; __constant__ size_t cdisp_step2; - void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th, - size_t leftstep) + void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th) { cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); @@ -84,8 +82,6 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) ); - - cudaSafeCall( cudaMemcpyToSymbol(cimg_step, &leftstep, sizeof(size_t)) ); } /////////////////////////////////////////////////////////////// @@ -217,7 +213,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, int h, int w, int level) + __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int h, int w, int level) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -257,7 +253,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost_reduce(const uchar *cleft, const uchar *cright, uchar *ctemp, int level, int rows, int cols, int h) + __global__ void init_data_cost_reduce(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int level, int rows, int cols, int h) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -305,7 +301,7 @@ namespace cv { namespace cuda { namespace device template - void init_data_cost_caller_(const uchar *cleft, const uchar *cright, uchar *ctemp, int /*rows*/, int /*cols*/, int h, int w, int level, int /*ndisp*/, int channels, cudaStream_t stream) + void init_data_cost_caller_(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int /*rows*/, int /*cols*/, int h, int w, int level, int /*ndisp*/, int channels, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -315,15 +311,15 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost<<>>(cleft, cright, ctemp, h, w, level); break; - case 3: init_data_cost<<>>(cleft, cright, ctemp, h, w, level); break; - case 4: init_data_cost<<>>(cleft, cright, ctemp, h, w, level); break; + case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level); break; + case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level); break; + case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void init_data_cost_reduce_caller_(const uchar *cleft, const uchar *cright, uchar *ctemp, int rows, int cols, int h, int w, int level, int ndisp, int channels, cudaStream_t stream) + void init_data_cost_reduce_caller_(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int rows, int cols, int h, int w, int level, int ndisp, int channels, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -334,19 +330,19 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, level, rows, cols, h); break; - case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, level, rows, cols, h); break; - case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, level, rows, cols, h); break; + case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h); break; + case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h); break; + case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, + void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, bool use_local_init_data_cost, cudaStream_t stream) { - typedef void (*InitDataCostCaller)(const uchar *cleft, const uchar *cright, uchar *ctemp, int cols, int rows, int w, int h, int level, int ndisp, int channels, cudaStream_t stream); + typedef void (*InitDataCostCaller)(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int cols, int rows, int w, int h, int level, int ndisp, int channels, cudaStream_t stream); static const InitDataCostCaller init_data_cost_callers[] = { @@ -359,7 +355,7 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) ); cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, sizeof(size_t)) ); - init_data_cost_callers[level](cleft, cright, ctemp, rows, cols, h, w, level, ndisp, channels, stream); + init_data_cost_callers[level](cleft, cright, ctemp, cimg_step, rows, cols, h, w, level, ndisp, channels, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -382,10 +378,10 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } - template void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, int rows, int cols, short* disp_selected_pyr, short* data_cost_selected, size_t msg_step, + template void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int rows, int cols, short* disp_selected_pyr, short* data_cost_selected, size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, bool use_local_init_data_cost, cudaStream_t stream); - template void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, int rows, int cols, float* disp_selected_pyr, float* data_cost_selected, size_t msg_step, + template void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int rows, int cols, float* disp_selected_pyr, float* data_cost_selected, size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, bool use_local_init_data_cost, cudaStream_t stream); /////////////////////////////////////////////////////////////// @@ -393,7 +389,7 @@ namespace cv { namespace cuda { namespace device /////////////////////////////////////////////////////////////// template - __global__ void compute_data_cost(const uchar *cleft, const uchar *cright, const T* selected_disp_pyr, T* data_cost_, int h, int w, int level, int nr_plane) + __global__ void compute_data_cost(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* selected_disp_pyr, T* data_cost_, int h, int w, int level, int nr_plane) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -436,7 +432,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void compute_data_cost_reduce(const uchar *cleft, const uchar *cright, const T* selected_disp_pyr, T* data_cost_, int level, int rows, int cols, int h, int nr_plane) + __global__ void compute_data_cost_reduce(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* selected_disp_pyr, T* data_cost_, int level, int rows, int cols, int h, int nr_plane) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -486,7 +482,7 @@ namespace cv { namespace cuda { namespace device } template - void compute_data_cost_caller_(const uchar *cleft, const uchar *cright, const T* disp_selected_pyr, T* data_cost, int /*rows*/, int /*cols*/, + void compute_data_cost_caller_(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* disp_selected_pyr, T* data_cost, int /*rows*/, int /*cols*/, int h, int w, int level, int nr_plane, int channels, cudaStream_t stream) { dim3 threads(32, 8, 1); @@ -497,15 +493,15 @@ namespace cv { namespace cuda { namespace device switch(channels) { - case 1: compute_data_cost<<>>(cleft, cright, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; - case 3: compute_data_cost<<>>(cleft, cright, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; - case 4: compute_data_cost<<>>(cleft, cright, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; + case 1: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; + case 3: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; + case 4: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void compute_data_cost_reduce_caller_(const uchar *cleft, const uchar *cright, const T* disp_selected_pyr, T* data_cost, int rows, int cols, + void compute_data_cost_reduce_caller_(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* disp_selected_pyr, T* data_cost, int rows, int cols, int h, int w, int level, int nr_plane, int channels, cudaStream_t stream) { const int threadsNum = 256; @@ -517,18 +513,18 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: compute_data_cost_reduce<<>>(cleft, cright, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; - case 3: compute_data_cost_reduce<<>>(cleft, cright, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; - case 4: compute_data_cost_reduce<<>>(cleft, cright, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; + case 1: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; + case 3: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; + case 4: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void compute_data_cost(const uchar *cleft, const uchar *cright, const T* disp_selected_pyr, T* data_cost, size_t msg_step, + void compute_data_cost(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* disp_selected_pyr, T* data_cost, size_t msg_step, int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream) { - typedef void (*ComputeDataCostCaller)(const uchar *cleft, const uchar *cright, const T* disp_selected_pyr, T* data_cost, int rows, int cols, + typedef void (*ComputeDataCostCaller)(const uchar *cleft, const uchar *cright, size_t cimg_step, const T* disp_selected_pyr, T* data_cost, int rows, int cols, int h, int w, int level, int nr_plane, int channels, cudaStream_t stream); static const ComputeDataCostCaller callers[] = @@ -544,17 +540,17 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaMemcpyToSymbol(cdisp_step2, &disp_step2, sizeof(size_t)) ); cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, sizeof(size_t)) ); - callers[level](cleft, cright, disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, stream); + callers[level](cleft, cright, cimg_step, disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); } - template void compute_data_cost(const uchar *cleft, const uchar *cright, const short* disp_selected_pyr, short* data_cost, size_t msg_step, + template void compute_data_cost(const uchar *cleft, const uchar *cright, size_t cimg_step, const short* disp_selected_pyr, short* data_cost, size_t msg_step, int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream); - template void compute_data_cost(const uchar *cleft, const uchar *cright, const float* disp_selected_pyr, float* data_cost, size_t msg_step, + template void compute_data_cost(const uchar *cleft, const uchar *cright, size_t cimg_step, const float* disp_selected_pyr, float* data_cost, size_t msg_step, int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream); diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index 72b5ab0..0854a92 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -2,15 +2,14 @@ namespace cv { namespace cuda { namespace device { namespace stereocsbp { - void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th, - size_t leftstep); + void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th); template - void init_data_cost(const uchar *left, const uchar *right, uchar *ctemp, int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, + void init_data_cost(const uchar *left, const uchar *right, uchar *ctemp, size_t cimg_step, int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, bool use_local_init_data_cost, cudaStream_t stream); template - void compute_data_cost(const uchar *left, const uchar *right, const T* disp_selected_pyr, T* data_cost, size_t msg_step, + void compute_data_cost(const uchar *left, const uchar *right, size_t cimg_step, const T* disp_selected_pyr, T* data_cost, size_t msg_step, int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream); template diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 55afd8c..be4d8f1 100644 --- a/modules/cudastereo/src/stereocsbp.cpp +++ b/modules/cudastereo/src/stereocsbp.cpp @@ -222,7 +222,7 @@ namespace //////////////////////////////////////////////////////////////////////////// // Compute - load_constants(ndisp_, max_data_term_, data_weight_, max_disc_term_, disc_single_jump_, min_disp_th_, left.step); + load_constants(ndisp_, max_data_term_, data_weight_, max_disc_term_, disc_single_jump_, min_disp_th_); l[0].setTo(0, _stream); d[0].setTo(0, _stream); @@ -245,12 +245,12 @@ namespace { if (i == levels_ - 1) { - init_data_cost(left.ptr(), right.ptr(), temp_.ptr(), left.rows, left.cols, disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), + init_data_cost(left.ptr(), right.ptr(), temp_.ptr(), left.step, left.rows, left.cols, disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), elem_step, rows_pyr[i], cols_pyr[i], i, nr_plane_pyr[i], ndisp_, left.channels(), use_local_init_data_cost_, stream); } else { - compute_data_cost(left.ptr(), right.ptr(), disp_selected_pyr[cur_idx].ptr(), data_cost.ptr(), elem_step, + compute_data_cost(left.ptr(), right.ptr(), left.step, disp_selected_pyr[cur_idx].ptr(), data_cost.ptr(), elem_step, left.rows, left.cols, rows_pyr[i], cols_pyr[i], rows_pyr[i+1], i, nr_plane_pyr[i+1], left.channels(), stream); int new_idx = (cur_idx + 1) & 1; @@ -276,12 +276,12 @@ namespace { if (i == levels_ - 1) { - init_data_cost(left.ptr(), right.ptr(), temp_.ptr(), left.rows, left.cols, disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), + init_data_cost(left.ptr(), right.ptr(), temp_.ptr(), left.step, left.rows, left.cols, disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), elem_step, rows_pyr[i], cols_pyr[i], i, nr_plane_pyr[i], ndisp_, left.channels(), use_local_init_data_cost_, stream); } else { - compute_data_cost(left.ptr(), right.ptr(), disp_selected_pyr[cur_idx].ptr(), data_cost.ptr(), elem_step, + compute_data_cost(left.ptr(), right.ptr(), left.step, disp_selected_pyr[cur_idx].ptr(), data_cost.ptr(), elem_step, left.rows, left.cols, rows_pyr[i], cols_pyr[i], rows_pyr[i+1], i, nr_plane_pyr[i+1], left.channels(), stream); int new_idx = (cur_idx + 1) & 1; -- 2.7.4